Below the API

Coalescing and Access Patterns

Intermediate Advanced 1h 15m Difficulty 4/5

Prerequisites 03, 08, IV.07

1. What is it?#

Coalescing is the hardware combining a warp’s 32 memory accesses into the minimum number of memory transactions.

COALESCED:                          UNCOALESCED (stride 32 floats):
thread 0 → byte 0..3                thread 0 → byte 0
thread 1 → byte 4..7                thread 1 → byte 128
...                                 ...
thread 31 → byte 124..127           thread 31 → byte 3968

→ 128 contiguous bytes              → 32 separate 32-byte sectors
→ 4 sectors, 1 burst                → 32 transactions
→ 100% of fetched bytes used        → 12.5% of fetched bytes used (4 of 32)

Up to 32x bandwidth waste from a memory access pattern, on a machine where bandwidth is your binding constraint.

2. Why does it exist?#

Because DRAM is efficient in bursts, not in individual bytes. HBM delivers data in 32-byte sectors. If a warp needs 32 floats scattered across 32 different sectors, the memory system must fetch 1,024 bytes to deliver 128.

This is the same principle as CPU cache lines (Section II.02), but the penalty is larger because a warp’s 32 lanes issue simultaneously.

3. Simple analogy#

A delivery van serving a street.

Coalesced: 32 packages for houses 1-32 on one street. One trip, van full.

Uncoalesced: 32 packages for house 1 on 32 different streets. Thirty-two trips, each with a van carrying one package. Same packages, 32x the fuel.

The van’s capacity (the 32-byte sector) is fixed. Your job is to fill it.

4. Tiny example#

// COALESCED: consecutive threads → consecutive addresses
__global__ void good(const float* in, float* out, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) out[i] = in[i] * 2.0f;
}

// UNCOALESCED: consecutive threads → addresses 32 floats apart
__global__ void bad(const float* in, float* out, int n, int stride) {
    int i = (blockIdx.x * blockDim.x + threadIdx.x) * stride;
    if (i < n) out[i] = in[i] * 2.0f;
}

Measure:

ncu --metrics l1tex__t_sectors_pipe_lsu_mem_global_op_ld.sum,\
l1tex__t_requests_pipe_lsu_mem_global_op_ld.sum,\
dram__bytes_read.sum ./bench

good:  sectors/request = 4    bandwidth 2,900 GB/s
bad:   sectors/request = 32   bandwidth   340 GB/s      8.5x worse

The sectors/request ratio is the diagnostic. 4 is perfect (128 bytes = 4 sectors per warp request). 32 is worst case.

5. Technical explanation#

The rules#

1. A warp's global memory access is serviced in 32-byte sectors.
2. The hardware coalesces addresses falling in the same sector.
3. Best case: 32 threads × 4 bytes = 128 contiguous bytes = 4 sectors.
4. Alignment matters: 128 bytes starting at a non-128-aligned address
   touches 5 sectors instead of 4.

Patterns, ranked#

PERFECT      out[tid] = in[tid]                   4 sectors/warp
GOOD         out[tid] = in[tid + offset]          4-5 (alignment)
OK           out[tid] = in[tid * 2]               8 (half the bytes used)
BAD          out[tid] = in[tid * 32]              32
WORST        out[tid] = in[random[tid]]           up to 32, unpredictable
BROADCAST    out[tid] = in[0]                     1 (hardware broadcasts — fine!)

Broadcast is worth noting: all threads reading the same address is efficient, not pathological. The hardware detects it and broadcasts. This is why __constant__ memory works well for values all threads need.

Vectorized loads#

// 4 bytes per thread: 128 bytes per warp, 4 sectors
float x = in[i];

// 16 bytes per thread: 512 bytes per warp, 16 sectors — but ONE instruction
float4 x = reinterpret_cast<const float4*>(in)[i];

Vectorized loads:

  • Reduce instruction count 4x.
  • Improve memory pipeline utilization.
  • Require 16-byte alignment (the pointer AND the index).

Typical gain on a bandwidth-bound kernel: 1.3-1.8x. This is one of the easiest optimizations available and is frequently missed in hand-written kernels.

// half8 for FP16 (16 bytes = 8 halves)
using half8 = uint4;   // reinterpret

The KV cache case, concretely#

Consider decode attention reading the KV cache with layout [num_blocks, block_size, n_kv_heads, head_dim]:

Thread t of a warp handles element t of head_dim (=128, so 4 warps per head).
Reading token j, head h:
  address = ((block × block_size + offset) × n_kv_heads + h) × head_dim + t

Consecutive t → consecutive addresses.  ✓ COALESCED

Now layout [num_blocks, n_kv_heads, head_dim, block_size]:

  address = ((block × n_kv_heads + h) × head_dim + t) × block_size + offset

Consecutive t → addresses block_size apart.  ✗ UNCOALESCED by 16x

Same data, same kernel logic, 16x difference. This is why KV layout is a real engineering decision (Section IV.07) and why vLLM’s layouts look strange — they’re chosen for the kernel’s access pattern.

Shared memory bank conflicts (the shared-memory analogue)#

32 banks × 4 bytes. A warp's 32 accesses:
  different banks     → 1 cycle
  same bank, same address → broadcast, 1 cycle
  same bank, different addresses → N-way conflict, N cycles
__shared__ float t[32][32];
t[tid][0]           // all 32 threads → bank 0 → 32-way conflict
t[0][tid]           // 32 consecutive → 32 banks → conflict-free

__shared__ float t[32][33];
t[tid][0]           // addresses differ by 33 → banks (33·tid) mod 32 = tid → conflict-free

The padding trick works because 33 is coprime with 32.

6. Under the hood#

Nsight Compute’s memory chart shows the whole path:

   L1/TEX Cache
   ├─ Requests:      how many warp-level instructions
   ├─ Sectors:       how many 32-byte sectors touched
   ├─ Hit rate:      L1 hits
   ↓
   L2 Cache
   ├─ Requests / Sectors / Hit rate
   ↓
   Device Memory
   └─ Bytes read/written

The key derived numbers:

sectors / request           → coalescing quality (4 = perfect for 32-bit loads)
dram_bytes / useful_bytes   → amplification factor

If dram_bytes is 8x what your algorithm needs, you have an access pattern problem.

7. Performance implications#

Issue                              Bandwidth achieved (of peak)
Perfectly coalesced, vectorized    85-95%
Coalesced, scalar loads            60-75%
Stride 2                           40-50%
Stride 8                           15-20%
Stride 32+                         3-6%
Random access                      2-10%

A memory-bound kernel with a bad access pattern can be 20x slower than necessary, and the profiler will show it as “memory bound” either way — which is why the sectors/request metric matters more than the bound classification.

8. Production implications#

  • Layout is chosen for the kernel, not for readability. When you see an odd tensor layout in vLLM or FlashAttention, it’s coalescing.
  • Check sectors/request when evaluating a custom kernel. It’s the fastest quality signal.
  • Gathers are inherently uncoalesced. Embedding lookups, MoE token routing, and sparse operations all pay this. Accept it or restructure (e.g. sort tokens by expert before the MoE GEMM — Section XIII.03).
  • Alignment matters. Slicing a tensor by an odd offset can misalign it and disable vectorized loads.
  • Padding to alignment is usually worth it. A few wasted bytes for 1.5x bandwidth.

9. Common mistakes#

Column-major traversal in a row-major tensor. The classic.

Ignoring alignment when slicing. x[1:] may be misaligned.

Not using vectorized loads. Free 1.3-1.8x.

Bank conflicts from unpadded shared arrays.

Assuming “memory bound” means “optimal.” It might mean “wasting 8x on bad access.”

Choosing a KV layout by copying a paper without checking your kernel.

10. Hands-on exercise#

A. Measure the penalty. Implement the good and bad kernels. Sweep stride from 1 to 64. Plot achieved bandwidth and sectors/request vs stride. Confirm the 32x worst case.

B. Vectorize. Take a bandwidth-bound kernel and convert it to float4 loads. Measure the improvement. Then deliberately misalign the pointer and observe the fallback.

C. KV layout. Implement a simplified decode-attention read for both layouts in section 5. Measure sectors/request and bandwidth for each. Confirm the 16x prediction.

D. Bank conflicts. Implement a shared-memory transpose with and without padding. Measure with l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_ld.sum.

E. Gather. Benchmark an embedding-lookup-style gather from a 1 GB table. What bandwidth do you achieve? Compare to a sequential read of the same byte count. Explain the gap.

11. Interview questions#

  1. What is memory coalescing and what’s the worst-case penalty?
  2. What metric tells you whether a kernel’s accesses are coalesced, and what value is ideal?
  3. Why do vectorized loads help, and what do they require?
  4. Design a KV cache layout for a decode attention kernel. Justify it in terms of coalescing.
  5. What is a shared-memory bank conflict and why does padding by 1 fix it?
  6. Is a broadcast (all threads reading the same address) a coalescing problem?
  7. Why are embedding lookups and MoE routing inherently uncoalesced, and what can you do?

12. Further reading#

  • [REFERENCE] CUDA Best Practices Guide, “Coalesced Access to Global Memory”
  • [REFERENCE] Nsight Compute Memory Workload Analysis
  • [REFERENCE] vLLM csrc/attention/attention_kernels.cu — read the layout comments
  • Next: 10 — Occupancy and warp divergence
esc

↑↓ navigate ↵ open