Coalescing and shared-memory tiling

Most "my kernel is slow" bugs are too many trips to global memory. Two levers come first after correctness: coalescing (make each trip wide) and tiling (make each trip useful many times). FlashAttention is the same idea applied to attention.

The big idea

A delivery truck that drops one letter per trip is coalescing gone wrong. Fill the truck with adjacent addresses. Then keep the letters on a side table (shared memory) so twenty calculations reuse one trip.

The picture breaks at bank conflicts. Shared memory is organized into banks (typically 32). If multiple threads in a warp hit the same bank for different addresses, accesses serialize. Padding tiles (for example tile[32][33]) is a common fix for column-strided patterns.

Walking through naive matmul

A naive kernel assigns one thread to one output C[row, col]. That thread loops over K, reading A and B from HBM every time.

__global__ void matmul_naive(const float* A, const float* B, float* C,
                             int M, int N, int K) {
  int row = blockIdx.y * blockDim.y + threadIdx.y;
  int col = blockIdx.x * blockDim.x + threadIdx.x;
  if (row >= M || col >= N) return;
  float sum = 0.f;
  for (int k = 0; k < K; ++k)
    sum += A[row * K + k] * B[k * N + col];
  C[row * N + col] = sum;
}

Each output does O(K) HBM loads with terrible reuse. Real GEMMs tile into shared memory, use vectorized loads, and on modern GPUs hand the inner tile to tensor cores.

The standard tiled pattern:

1. Cooperatively load a tile of A and B from HBM into shared memory.
2. __syncthreads().
3. Compute the partial product from the tile.
4. Repeat across K.
5. Write the result.

That raises arithmetic intensity by reusing each loaded element many times — the roofline lever.

Rules of thumb for coalescing:

• Have threadIdx.x map to the fastest-moving index in memory (often the column in row-major).
• Prefer layouts that keep warp lanes contiguous.
• Misaligned starts are cheaper than they used to be. Strided and sparse gathers remain expensive.

Simon Boehm's blog (linked from the CUDA hub) walks a matmul from this naive form to something near cuBLAS.

Common mistakes

• threadIdx.x on the slow index. The warp gathers scattered addresses and effective bandwidth collapses.
• Tiling without __syncthreads(). Threads read a tile that has not finished loading.
• Ignoring bank conflicts. A pretty tile can still serialize on shared memory.

Check yourself

1. Why is naive matmul slow even though it is correct? Answer: each output does O(K) HBM loads with almost no reuse.
2. What does coalescing require of a warp's loads? Answer: adjacent addresses, so the hardware merges them into few transactions.

Key takeaways

• Coalesced HBM access and shared-memory tiling are the first two performance levers after correctness.
• Naive per-element matmul teaches indexing and why libraries exist.
• Padding tiles avoids bank conflicts on column-strided patterns.

Further learning

Next, hide the remaining HBM latency with occupancy and launch overhead.