GPU Glossary
/
Performance·3 min read

Memory Coalescing

A hardware technique that serves several logical global memory reads with one physical access, so memory bandwidth is used fully.

Memory coalescing is a hardware technique that improves use of memory bandwidth by servicing multiple logical memory reads with a single physical memory access. It applies to global memory; for efficient shared memory access, see bank conflict.

Why it's needed#

In CUDA GPUs, global memory is backed by GPU RAM, built from DRAM technologies such as GDDR or HBM. These offer high bandwidth but long access latency, even compared with DDR5, the peer technology in CPU RAM.

DRAM latency is set by how fast tiny capacitors can charge their access lines, which is fundamentally limited by thermal, power and size constraints. With latency that high, serving each logical access as its own physical access would leave much of the GPU's bandwidth unused.

How it works: DRAM bursts#

Each time a DRAM address is accessed, several consecutive addresses are fetched together in parallel in a single clock. This transfer is called a DRAM burst. When multiple concurrent logical accesses are served by one burst, the access is coalesced. (A physical access is part of a memory transaction, a term you'll also see in discussions of coalescing.)

For more detail, see Section 6.1 of the 4th edition of Programming Massively Parallel Processors; for comprehensive detail, see Ulrich Drepper's What Every Programmer Should Know About Memory.

CPUs get a similar benefit by mapping bursts onto cache lines. As often in GPU programming, what CPUs handle automatically in cache is here up to the programmer.

Why it fits the SIMT model#

That's less painful than it sounds, because bursts line up neatly with the single-instruction, multiple-thread (SIMT) model of CUDA PTX. Normally all threads in a warp run the same instruction at the same time. That makes it easy for a CUDA programmer to write coalesced accesses, and easy for the memory hardware to detect them.

Key idea A single burst typically serves 128 bytes: not by coincidence, exactly one 32-bit float for each of the 32 threads in a warp. Have consecutive threads read consecutive addresses.

Measuring the impact#

This kernel reads an array with a variable stride between accessed elements. As stride grows, each warp's read needs more DRAM bursts, so there are more physical accesses per logical access and throughput drops.

__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 * (size_t)stride; j < N; j += (size_t)T * (size_t)stride) {
        // across a warp, addresses differ by (stride * sizeof(float))
        float v = in[j]; // perfectly coalesced for stride == 1
        acc = acc * 1.000000119f + v;  // force compiler to keep the load
    }

    // do one write per thread (negligible vs reads)
    if (t < N) out[t] = acc;
}

A micro-benchmark on Godbolt (reproduce it here) shows the expected relationship:

# Device: Tesla T4 (SM 75)
# N = 67108864 floats (256.0 MB), iters = 10
stride        GB/s
    1       206.0
    2       130.5
    4        68.8
    8        33.8
   16        16.8
   32        15.2
   64        13.6
  128        11.2

A stride of 2 roughly halves throughput, because each warp's request needs twice as many bursts. Stride 4 halves it again. By stride 16, throughput is down 16×, and after that it degrades differently, presumably because other parts of the memory subsystem start to show, with worse performance from reduced locality (e.g. on-device TLB misses).

For more best practices, see How to Access Global Memory Efficiently in CUDA C/C++ Kernels on the NVIDIA Developer blog.