What is Memory Coalescing?

Memory coalescing is a hardware technique that improves memory bandwidth utilization by folding several logical memory reads into a single physical memory access. It applies to accesses of global memory; the equivalent concern for shared memory is a bank conflict, which is a different mechanism with a different fix.

On CUDA GPUs, global memory is backed by GPU RAM: DRAM built with technologies like GDDR or HBM. DRAM offers very high bandwidth but comparatively long access latency, higher even than the DDR5 used for CPU RAM, because that latency is set by how fast the tiny capacitors that store each bit can charge, a limit set by thermal, power, and physical size constraints rather than anything a chip designer can trivially fix. If every logical memory access were serviced as its own separate physical access, that latency would leave most of the GPU's memory bandwidth unused.

Coalescing gets around this by exploiting how DRAM actually works: each time an address is accessed, a whole run of consecutive addresses is fetched together in a single clock, a unit of transfer usually called a DRAM burst. When several threads' logical accesses fall inside the same burst, the memory controller services them with one physical access, or memory transaction, instead of several. CPUs get a similar win for free from cache lines; on a GPU this behavior is programmer-managed rather than automatic, which is exactly why access-pattern choices matter so much for kernel performance.

Coalescing lines up naturally with CUDA's single-instruction, multiple-thread execution model, because every thread in a warp executes the same instruction at the same moment. If those 32 threads also touch consecutive addresses, the pattern is trivial for the hardware to detect and coalesce. A typical burst services 128 bytes, which not coincidentally is exactly enough for each of a warp's 32 threads to load one 4-byte float.

The performance gap between coalesced and uncoalesced access shows up clearly in a strided read kernel, which reads array elements spaced stride positions apart instead of consecutively:

__global__ void strided_read_kernel(const float* __restrict__ in,
                                     float* __restrict__ out,
                                     size_t N, int stride) {
  const size_t t = blockIdx.x * blockDim.x + threadIdx.x;
  const size_t T = gridDim.x * (size_t)blockDim.x;

  float acc = 0.f;
  for (size_t j = (size_t)t * stride; j < N; j += (size_t)T * stride) {
    // consecutive threads read addresses stride * 4 bytes apart
    float v = in[j];
    acc = acc * 1.000000119f + v;
  }
  if (t < N) out[t] = acc;
}

As stride grows, each warp's 32 logical reads spread across more DRAM bursts, so more physical accesses are needed to service the same amount of useful work. A microbenchmark on a Tesla T4 makes the effect concrete:

Doubling the stride from 1 up through 8 roughly halves throughput at each step, exactly what you'd expect from doubling the number of bursts needed per warp. Past a stride of 16 the relationship flattens out, likely because other parts of the memory subsystem, such as TLB misses from reduced locality, start to dominate instead of burst count alone.

In practice, getting coalesced access is mostly about data layout rather than clever indexing. A structure-of-arrays layout, where each field lives in its own contiguous array, lets a warp read one field for 32 consecutive elements in a single coalesced transaction; an array-of-structures layout scatters that same field 32 struct-widths apart, which is exactly the strided pattern above. Aligning the base address of an allocation to the burst size, and reading multiple elements per thread with vector types like float4 instead of four separate float reads, are the other two levers that show up most often in tuned kernels.

Coalescing is one of the first things worth checking when a kernel underperforms what its arithmetic intensity would predict: a memory-bound kernel with a bad access pattern can be leaving most of its rented bandwidth unused even though the math looks right. Aquanode's GPU metrics surface achieved memory bandwidth per job next to the theoretical peak for whatever card you rented, so a coalescing problem shows up as a number on a dashboard instead of a guess from staring at kernel code.