Memory Coalescing — infra.bea.ai

The single most impactful optimization for any GPU kernel is ensuring coalesced memory access — where threads in a warp read consecutive addresses, allowing the hardware to combine 32 reads into a single transaction. Getting this wrong can cost 10–32× performance.

1. Why Memory Access Dominates Performance

Most CUDA kernels are memory-bound — the GPU finishes computation faster than it can fetch data from HBM. The roofline model tells us:

H100 FP16 compute: 989 TFLOPS
H100 HBM bandwidth: 3,350 GB/s
Ops:byte ratio: 295

Any operation with arithmetic intensity below 295 is memory-bound. That includes: element-wise operations (~1), LayerNorm (~5), Softmax (~2), attention score computation (~50-200), and single-token decode GEMMs (~1).

Conclusion: for most kernels, optimizing memory access patterns matters more than optimizing compute.

2. How Coalescing Works

The CUDA Memory Model article introduced coalesced access. Here we go deeper.

Coalesced Access

Memory Transactions

When a warp of 32 threads accesses global memory, the hardware groups the 32 addresses into memory transactions of 32, 64, or 128 bytes. Fewer transactions = faster.

Ideal case (coalesced): 32 threads access 32 consecutive 4-byte floats → one 128-byte transaction.

Worst case (scattered): 32 threads access 32 random addresses → up to 32 separate transactions.

Alignment Requirements

Transactions are aligned to their size. A 128-byte transaction starts at an address that’s a multiple of 128. If a warp’s access spans two 128-byte segments, it generates two transactions instead of one.

// ALIGNED: base address is 128-byte aligned, one transaction
float val = data[blockIdx.x * 256 + threadIdx.x];  // if data is cudaMalloc'd (aligned)

// MISALIGNED: shifted by 1 element, may span two segments
float val = data[blockIdx.x * 256 + threadIdx.x + 1];  // possibly 2 transactions

cudaMalloc always returns 256-byte aligned memory. Misalignment issues arise from pointer arithmetic within kernels.

3. Access Patterns and Their Cost

Pattern 1: Consecutive (Coalesced) ✅

// Thread i reads element i — consecutive addresses
float val = data[threadIdx.x];

All 32 threads in a warp read addresses [0, 4, 8, ..., 124] → 1 transaction (128 bytes).

Pattern 2: Strided ❌

// Thread i reads element i*stride
float val = data[threadIdx.x * stride];

With stride = 2: threads read addresses [0, 8, 16, ..., 248] → 2 transactions (spans 256 bytes). With stride = 32: each thread’s address is in a different 128-byte segment → up to 32 transactions.

Rule: Every doubling of stride roughly doubles the number of transactions.

Pattern 3: Random (Gather) ❌❌

// Thread i reads a random address
float val = data[indices[threadIdx.x]];

Worst case: 32 threads, 32 different cache lines → 32 transactions. This is 32× slower than coalesced.

When you must scatter/gather: Load into shared memory with coalesced reads, then shuffle in shared memory.

4. Structure of Arrays (SoA) vs Array of Structures (AoS)

This is where coalescing has the biggest practical impact:

AoS: Bad for GPUs

struct Particle {
    float x, y, z, w;  // 16 bytes per particle
};
Particle particles[N];

// Thread i accesses particles[i].x
// Memory layout: [x0,y0,z0,w0, x1,y1,z1,w1, x2,y2,z2,w2, ...]
// Thread 0 reads offset 0, thread 1 reads offset 16 → stride of 4 floats
// Result: 4 transactions instead of 1

SoA: Good for GPUs

float x[N], y[N], z[N], w[N];

// Thread i accesses x[i]
// Memory layout: [x0,x1,x2,...,xN, y0,y1,y2,...,yN, ...]
// Thread 0 reads x[0], thread 1 reads x[1] → consecutive!
// Result: 1 transaction

Performance difference: For the particle example, SoA is 4× faster than AoS just from coalescing. For structures with more fields, the gap is even larger.

2D Array Access: Row-Major Matters

// Row-major array: data[row][col]
// Stored as: data[0][0], data[0][1], ..., data[0][N-1], data[1][0], ...

// COALESCED: threads read along a row (consecutive in memory)
float val = data[row][threadIdx.x];  // ✅

// UNCOALESCED: threads read along a column (stride = row_width)
float val = data[threadIdx.x][col];  // ❌ stride = N

This is why matrix transpose is a classic shared memory optimization — load a tile with coalesced row reads, transpose in shared memory, write with coalesced row writes.

5. Vectorized Loads

Instead of loading one float at a time, load 2 or 4 at once:

// Standard: 1 float per thread per load (4 bytes)
float val = data[idx];

// Vectorized: 4 floats per thread per load (16 bytes)
float4 vec = reinterpret_cast<float4*>(data)[idx];
// vec.x, vec.y, vec.z, vec.w — four consecutive floats

Why it helps: One float4 load generates a single 128-bit transaction per thread. With 32 threads, that’s 32 × 16 = 512 bytes in one warp — vs 128 bytes with scalar loads. More data per transaction = better bandwidth utilization.

Requirements: Data must be 16-byte aligned for float4, 8-byte for float2.

// Common pattern: process 4 elements per thread
int idx = (blockIdx.x * blockDim.x + threadIdx.x) * 4;
float4 a = reinterpret_cast<float4*>(A)[idx/4];
float4 b = reinterpret_cast<float4*>(B)[idx/4];
float4 c;
c.x = a.x + b.x;
c.y = a.y + b.y;
c.z = a.z + b.z;
c.w = a.w + b.w;
reinterpret_cast<float4*>(C)[idx/4] = c;

6. How to Check Coalescing

Nsight Compute Metrics

ncu --metrics l1tex__t_sectors_pipe_lsu_mem_global_op_ld.sum kernel

Key metrics:

Global load efficiency: gld_efficiency — percentage of bytes actually used vs bytes transferred. 100% = perfect coalescing. 25% = 4× waste.
Sectors per request: ideal is 4 (one 128-byte cache line = 4 sectors of 32 bytes). Higher means uncoalesced.

Quick Check: Bandwidth Utilization

Measured bandwidth = bytes_read / kernel_time
Achievable bandwidth = HBM peak × efficiency_factor (~80%)
If measured << achievable → likely uncoalesced access

Self-Check

Explain how 32 threads’ addresses map to memory transactions
Calculate the number of transactions for stride-2 vs stride-1 access
Convert an AoS data structure to SoA for GPU
Explain why column-major access in a row-major array is slow
Use float4 vectorized loads and explain the alignment requirement
Check global load efficiency using Nsight Compute