Coalescing is what happens between the moment a warp issues one load instruction and the moment DRAM hands back data: the memory subsystem takes the 32 addresses the warp’s lanes produced and folds them into as few off-chip transactions as it can. When the addresses are neighbours, that folding is nearly free and the warp pays for exactly the bytes it asked for. When they are scattered, the hardware still moves data in fixed-size chunks — so the warp pays for a whole chunk per lane and throws away most of it. Nothing in the instruction stream changes; only the addresses do. So coalescing is less a tuning knob than a property of your data layout. This piece covers global memory: the granularity, the arithmetic, the layouts, and the number to look at.
One instruction, thirty-two addresses
A warp executes in lockstep: one instruction, 32 lanes. When that instruction is a global load, each lane computes its own address, so the load-store unit is handed 32 addresses at once and must decide how many memory requests they imply. This is the coalescing step. It is done in hardware, per warp, per instruction, and it has no visibility into your intent — only into the numeric addresses.
The canonical good case is a[blockIdx.x * blockDim.x + threadIdx.x] on a 4-byte type. Lane 0 wants bytes 0–3, lane 1 wants 4–7, and so on up to lane 31 at bytes 124–127. Those 32 addresses cover one contiguous 128-byte run, and the hardware satisfies the whole warp with a single line’s worth of traffic. The canonical bad case is the same instruction — only the index expression differs — producing 32 addresses in 32 unrelated places: same code shape, an order of magnitude more traffic. Coalescing is invisible in the source and enormous in the profile, which is why it gets missed.
Sectors and cache lines — the granularity floor
The reason scattered access is expensive is that memory is not byte-addressable all the way down to DRAM. Global loads are serviced through the L1/texture path and L2 in fixed blocks: a 128-byte cache line made of four 32-byte sectors. Traffic between L1 and L2 is counted in sectors, and a sector is the smallest thing that can move. Ask for one byte and 32 bytes travel.
That floor sets the whole cost model. The ideal 32-lane, 4-byte load touches exactly four sectors — the minimum possible for 128 bytes of payload, so every byte moved is a byte wanted. Now let the same warp read one 4-byte element from 32 widely separated places. No two addresses share a sector, so the warp generates 32 sector requests: 1024 bytes moved to deliver 128 useful bytes. Same single instruction, eight times the work. Coalescing, stated precisely, is the art of getting a warp’s addresses to share sectors.
Sector efficiency — requested bytes over transferred bytes
The useful metric follows directly: sector efficiency = requested bytes ÷ transferred bytes, where requested is what your lanes actually consume and transferred is what the sectors carried. It is a single ratio, it has an obvious ceiling of 100%, and unlike wall-clock timing it tells you why a kernel is slow rather than just that it is.
Work the arithmetic for a warp reading floats. Stride 1: 128 bytes requested, 4 sectors × 32 = 128 bytes transferred, 100%. Stride 2: the same 128 bytes requested now spread across 256 bytes of address space, 8 sectors, 256 bytes transferred — 50%. Stride 8: 32 sectors, 1024 bytes, 12.5%. Beyond stride 8 for a 4-byte type nothing more can be lost: each lane already owns a private sector. These figures are arithmetic, not benchmarks — they fall out of the sector size and the element size — and a kernel at 12.5% is doing 8× the DRAM traffic it needs to.
Stride and the slide toward one transaction per lane
Stride is the variable that moves you along that curve, and the degradation is not gradual. Going from stride 1 to stride 2 halves efficiency in one step; once the stride in bytes reaches the sector size, every lane is isolated and you are paying the maximum. There is no long tail of mild penalties.
Where do large strides come from? Almost always from indexing a 2-D array along the wrong axis. For a row-major matrix, A[row][col] with col = threadIdx.x walks consecutive lanes across consecutive elements — stride 1, perfect. Swap the roles so that row = threadIdx.x and each lane jumps a full row: the stride is the row length, which for any realistic matrix is far past the sector size. The fix is usually not to change the storage but to change which index the fastest-varying thread dimension maps to. threadIdx.x should always drive the contiguous dimension. That one habit prevents most stride problems before they exist.
Array-of-structs vs struct-of-arrays
The layout decision that determines coalescing is made before a single line of kernel code is written, and it usually looks like a data-modelling choice rather than a performance one. An array of structs — struct Particle { float x, y, z; }; Particle p[N]; — is the natural CPU-side representation: one object, its fields together, cache-friendly for a single thread that touches all three.
On a GPU it is close to the worst case. A warp evaluating p[i].x has lanes 12 bytes apart, so a 128-byte payload of x-values is smeared across roughly 384 bytes and 12 sectors; two thirds of everything you move is y and z that this instruction does not want. The struct of arrays form — separate float x[N], y[N], z[N] — restores stride 1 for each field, and each of the three loads is individually perfect. The cost is ergonomic: you lose the tidy object and pass three pointers instead of one. That trade is almost always worth taking, and it is why SoA is the default layout in GPU-facing code.
Alignment — why a shifted base pointer costs extra
Contiguity is necessary but not sufficient: the run of bytes also has to sit where the sectors sit. A warp reading 128 contiguous bytes starting at a 128-byte boundary occupies exactly four sectors. Start the same 128 bytes four bytes later and the run straddles a boundary, spilling into a fifth sector — 160 bytes moved for 128 wanted, an unforced 25% overhead on an access that looks perfectly sequential in the source.
Under the coarser 128-byte transaction granularity of older hardware the same misalignment cost a full second transaction, doubling the traffic; sectored access softened that penalty but did not remove it. In practice the base pointers you get from cudaMalloc are already generously aligned, so misalignment is nearly always self-inflicted: an offset added to a pointer before passing it to a kernel, a row pitch that is not a multiple of the sector size, or a header packed in front of a payload. Padding rows to a friendly multiple — what a pitched allocation does for you — costs a little memory and buys back alignment on every row.
The transpose problem and the staging buffer
Some access patterns cannot be fixed by layout because the kernel genuinely needs both orientations. Transpose is the archetype: read in[row][col], write out[col][row]. Whichever way you assign threadIdx.x, one of those two accesses is stride 1 and the other has a stride of a full row. You can move the problem from the read to the write, but you cannot make it disappear.
The standard escape is to stage through on-chip memory. Have the block read a tile from global memory with a fully coalesced, stride 1 pattern and deposit it in a shared-memory buffer. Synchronize. Then have the block read that tile back out of shared memory in transposed order and write it to global memory — again stride 1, again coalesced. The transposition happens entirely on chip, where scattered access is cheap, and both off-chip accesses stay perfect. (The on-chip half has its own hazard, bank conflicts, which belongs to the shared-memory article.) The lesson generalizes: when a permutation is unavoidable, perform it where the granularity floor does not apply.
Measure it, do not eyeball it
Reading an index expression and pronouncing it coalesced is unreliable, because the address a lane produces depends on the block geometry, the launch configuration, and the pitch of the allocation, none of which are visible in the line you are staring at. Measure instead.
The number that settles the question is sectors per request: the average count of sectors a single warp-wide global access generated. For a 32-lane, 4-byte load the perfect value is 4 and the floor is 32, so the ratio you see maps straight onto the efficiency arithmetic above — 8 means half your traffic is waste, 32 means seven eighths of it is. The equivalent framing, global load and store efficiency, expresses the same thing as a percentage. Both come out of the memory workload analysis in a GPU profiler, which reports requested versus transferred bytes per access instruction; the profiling article covers driving that tool. The discipline is to make sectors-per-request a number you check, not a theory you argue about.
When you cannot get to stride 1
Some access patterns are irreducibly irregular — sparse matrices, gather and scatter through an index array, graph traversal. Coalescing is not a binary you win or lose outright. Sorting or bucketing indices so lanes within a warp land in the same neighbourhood recovers sector sharing without making the access truly sequential, and reordering the data so elements accessed together are stored together does the same thing permanently.
Two other levers are worth knowing. Vectorized loads — having each lane fetch a float4 instead of a float — keep the access fully coalesced while moving four times the payload per instruction, which cuts instruction count and improves the memory pipeline’s ability to keep requests in flight. And broadcast is the degenerate best case: when all 32 lanes read the same address, the hardware fetches one sector and hands it to everyone. If a value is uniform across a warp, reading it costs almost nothing, so do not contort a layout to avoid it.
threadIdx.x on the contiguous dimension, rows padded so each base stays aligned. Where a permutation is genuinely required, stage it through shared memory so both global accesses stay stride 1. And check sectors-per-request rather than reading the index expression and hoping.