100 DAYS OF GPU Course preview
← Glossary
CUDA glossaryMemory
CC 7.5

What is memory coalescing in CUDA?

The hardware merging a warp's 32 addresses into the fewest memory transactions that cover them: cheap when the addresses are consecutive, expensive when they are not.

A warp issues one load instruction for all 32 of its lanes, and the memory system has to turn those 32 addresses into fetches. Global memory is served in 32-byte sectors, and each sector fetched is one transaction. When lane 0 reads data[0] and lane 31 reads data[31], those 32 floats sit in 128 contiguous bytes, so four sectors cover the warp and every byte fetched gets used. Coalescing is not a flag you pass or a pass the compiler runs. It is what falls out of 32 lanes sharing one memory pipe.

Space the addresses out and the same instruction costs more. A 32-byte sector holds eight floats, so at stride s a warp uses at most 1 / min(s, 8) of every byte it pulls. By stride 8 each lane owns a sector of its own: the warp needs 32 transactions instead of 4, and seven eighths of the bandwidth carries data nobody asked for. Past stride 8 the transaction count cannot grow, because one sector per lane is the ceiling.

That is where most explanations stop, and stopping there is the mistake. Transactions saturate at stride 8; the slowdown does not. Day 11 measured 20.1 GB/s at stride 8, 14.2 at stride 16 and 9.6 at stride 32, while Nsight Compute reported the same 67,108,864 sectors for all three. What keeps costing you past that point is locality, not transactions: the lanes spread across more 128-byte cache lines, then across more DRAM pages, and a sector count sees neither. Treat 1 / min(s, 8) as a ceiling on efficiency rather than an estimate of it, because even inside its own regime the T4 does worse than the model predicts. At stride 2 the model says half and the card delivered a third.

The pattern that catches good CPU programmers is giving each thread its own contiguous chunk. With a private cache per core that is the right answer. On a GPU the 32 lanes of a warp are on the same loop iteration at the same instant, so their addresses sit one chunk apart, and a 32-element chunk rebuilds the stride-32 case by accident. Day 11 measured it at 32.6 GB/s: three times better than stride 32, because a warp's contiguous working set stays cache-resident, and still about seven times worse than doing it right. The fix is the grid-stride loop from day 8, where consecutive lanes hold consecutive addresses on every iteration.

Measured

On a Tesla T4, driver 595.84, CUDA 12.6 (V12.6.85), built with nvcc -O3 -arch=sm_75, day 11 copied the same buffer seven ways. Every variant moved 536,870,912 bytes, asserted by the harness, so no row can win by doing less work.

Pattern GB/s Load requests Load sectors Sectors per request
stride 1 232.9 2,097,152 8,388,608 4
stride 2 87.2 2,097,152 16,777,216 8
stride 4 41.8 2,097,152 33,554,432 16
stride 8 20.1 2,097,152 67,108,864 32
stride 16 14.2 2,097,152 67,108,864 32
stride 32 9.6 2,097,152 67,108,864 32
chunked, 32 per thread 32.6 not profiled not profiled not profiled

Bandwidth comes from the timed run, the counters from an Nsight Compute pass on the same card, in the same session, on the same kernel. Two things there matter more than the headline ratio. The request count never moves, because a pattern changes what an instruction costs and not how many issue. And the last three rows share one sector count while bandwidth halves twice, which is the measurement that separates the transaction model from what you actually pay.

Getting the counters needed root. Plain ncu returned ERR_NVGPUCTRPERM, sudo ncu worked, and the cause is the stock driver default RmProfilingAdminOnly: 1 in /proc/driver/nvidia/params. On a hosted tier such as Colab you do not get either, which is why day 11 ships its own report.

Diagram

coalescing-visualizer: 32 lane boxes above a strip of 32-byte sectors, with a stride control. Stride 1 lands all 32 arrows in four adjacent sectors, every byte shaded as used. Stride 8 fans them out to 32 sectors with 28 of each 32 bytes greyed. Stride 32 keeps the same 32 sectors and spreads them over 32 different 128-byte lines.

Alt text: "One warp at three strides. Contiguous addresses need four sectors and use every byte. Stride eight and stride thirty-two both need thirty-two sectors and use an eighth of the bytes, but stride thirty-two spreads them over four times as many cache lines."

Code

From code/day11-coalescing/coalescing.cu. One kernel serves every stride, so only the address pattern changes across rows. The 2D launch is what makes it a permutation rather than a sampling, which is the invariant the benchmark rests on.

__global__ void copyPermuted(const float* __restrict__ in,
                             float* __restrict__ out, size_t m, int stride) {
    size_t col = blockIdx.x * static_cast<size_t>(blockDim.x) + threadIdx.x;
    if (col < m) {
        size_t j = col * static_cast<size_t>(stride) + blockIdx.y;
        out[j] = in[j];
    }
}

Consecutive lanes get consecutive col, so their addresses sit stride elements apart. There is no modulo: an earlier version used % n on a runtime size_t, which nvcc cannot reduce to a mask, so the strided rows paid for a division the coalesced row did not.

Related terms

Where you meet this

Sources

Byline

Written by: pending. Reviewed by: pending. Written on: pending. Last checked on: pending. This entry stays a draft until a named author and a different named reviewer sign it.