Memory Coalescing

Definition

Coalescing is the hardware combining the 32 addresses a warp issues in one memory instruction into the minimum number of sector transactions (32-byte sectors; 128-byte cache lines = 4 sectors). Global memory is never fetched byte-by-byte — how a warp’s threads map onto sectors determines your effective HBM bandwidth. Uncoalesced access is the #1 beginner GPU performance bug.

Key math

Coalesced (ideal): 32 threads read 32 consecutive 4-byte words → 128 contiguous bytes → exactly one 128-byte line = 4 sectors, all used:
$\text{addr}(t)=\text{base}+4t,t=0\dots 31\Rightarrow \text{1 transaction, 128 B, 0 waste}$

Strided: thread $t$ reads base + 4·stride·t. With stride = 32, all 32 threads touch 32 different 128-byte lines — fetch $32\times 128=4096$ B to deliver the $128$ B wanted:
$\text{efficiency}=\frac{128}{4096}=\frac{1}{32}\approx 3.1\%$
A 32× effective-bandwidth loss from address pattern alone: the DRAM pins still run at full 3.35 TB/s, but 97% of the bytes crossing them are discarded. Efficiency roughly halves per doubling of stride until it floors at $1/32$.

Why it matters

Row-major arrays where consecutive threads map to consecutive rows (not columns), Array-of-Structs layouts, and random gather/scatter all destroy coalescing. The fix is a data-layout change: Struct-of-Arrays, transposing which index the thread ID maps to, or staging through shared memory so the global access is coalesced even if the SMEM access isn’t. Rule: make threadIdx.x (fastest-varying thread index) map to the fastest-varying (contiguous) memory dimension. Coalesced tile loads are half the recipe (with conflict-free SMEM) that takes GEMM and FlashAttention to near-peak; for bandwidth-bound LLM ops, coalescing directly sets the achieved throughput.

Taught in

• 02-gpu-memory-hierarchy — §6 global memory and coalescing.

See also

• gpu-memory-hierarchy — the HBM tier being accessed
• warp-simt — the 32 addresses combined
• shared-memory-bank-conflicts — the on-chip analog
• gemm — coalesced tile loads
• arithmetic-intensity — bytes moved per FLOP
• gpu-systems-for-llms