Memory Coalescing

Memory does not arrive one value at a time. It arrives in fixed-size chunks, and the hardware serves a whole thread group's requests together. When the lanes in a group ask for addresses that fall inside the same few chunks, those requests merge into a handful of transactions. When they ask for scattered addresses, each one drags in its own chunk and almost all of the bytes fetched are thrown away.

Same data, same instruction count, an order of magnitude difference in time. This is usually the largest single lever in a memory-bound kernel.

A load instruction is issued for the whole group at once. The memory system looks at the addresses every lane wants, works out which chunks of memory cover them, and issues one transaction per distinct chunk.

So the cost of a load is not the number of lanes, it is the number of distinct chunks touched. Thirty-two lanes reading thirty-two consecutive floats touch a small number of chunks and cost a small number of transactions. Thirty-two lanes reading floats a kilobyte apart touch thirty-two different chunks and cost thirty-two transactions, of which you use four bytes each.

The useful mental model is bytes requested against bytes delivered. Coalesced access uses nearly all of what the memory system fetched. Scattered access can use a few percent of it, and the wasted fraction is bandwidth you paid for and threw away.

Consecutive threadIdx.x should touch consecutive addresses. That is the whole rule, and it is why the global index formula puts threadIdx.x in the least significant position:

int i = blockIdx.x * blockDim.x + threadIdx.x;
float v = data[i];      // coalesced: lane n reads element n
            

Striding. Each lane reads every Nth element. The classic case is a column-major walk through a row-major array, or a matrix transpose where one of the two accesses is necessarily strided.

// coalesced read, strided write
out[x * height + y] = in[y * width + x];
            

Array of structs. Storing particles as {x, y, z, vx, vy, vz} and having each lane read one particle's x means the lanes touch addresses 24 bytes apart. Flipping to a struct of arrays, one array per field, makes every field access contiguous. This is often the single highest-value change available in a simulation kernel, and it is a data layout decision rather than a kernel one.

Indirection. Gather and scatter through an index array is inherently scattered and there is no trick that makes it contiguous. Here the goal shifts to improving locality, by sorting or partitioning the indices so that lanes in a group at least land near each other.

Beyond ordering, the base address matters. An allocation returned by the device allocator is generously aligned, so arrays start on a chunk boundary and a contiguous group access falls inside the minimum number of chunks. Slicing into the middle of an array at an arbitrary offset can push a group across one extra boundary and cost an additional transaction per access.

The effect is far smaller than the ordering problem and is usually not worth restructuring for, but it is why padding a row stride sometimes produces an unexplained few percent.

This is one of the easier things to confirm, because the profilers report the wasted fraction directly.