GPU memory — coalescing, shared memory & occupancy.

Three rules decide whether a kernel is fast: (1) make a warp's 32 reads land in one contiguous chunk so the hardware coalesces them into a single transaction; (2) stage reused data in shared memory, the on-chip scratchpad — while avoiding bank conflicts; (3) keep enough warps resident (occupancy) to hide the memory latency you cannot remove. Article 1.1 gave you the machine; this article gives you the rules that machine rewards.

In article 1.1 we established the core truth: a GPU is bandwidth-starved, not compute-starved. The arithmetic units can almost always go faster than memory can feed them. So the craft of GPU programming is mostly the craft of moving bytes efficiently — getting the right data, to the right place, at the right time, in the widest possible transactions. This article is the practical core of that craft. Everything in Phases 2 and 3 builds on these three ideas.

Rule 1: Coalesced global memory access

When a warp executes a load instruction, all 32 threads issue their addresses at once. The memory subsystem does not service 32 separate requests. Instead it looks at the 32 addresses and groups them into memory transactions of a fixed size — typically 32, 64, or 128 bytes, aligned to those boundaries. The fewer transactions needed to satisfy the warp, the better.

Coalescing is what happens in the best case: if the 32 threads access 32 consecutive 4-byte words (a contiguous 128-byte block, properly aligned), the hardware satisfies the whole warp with a single 128-byte transaction. Maximum efficiency: every byte fetched is a byte used.

The opposite case is a scattered or strided access. If each thread reads a word far from its neighbors, the warp may need up to 32 separate transactions — and because each transaction still moves a minimum of 32 bytes, you might fetch 32×32 = 1024 bytes to use only 32×4 = 128 of them. That is a 8× waste of bandwidth on the metric that matters most.

The practical rule is short: thread i should access element i. When you index global memory as data[blockIdx.x * blockDim.x + threadIdx.x], consecutive threads hit consecutive addresses and the access coalesces automatically. The classic mistake is the opposite — giving each thread a row of a matrix to march down. Then on any given step, the 32 threads are 32 rows apart in memory, and every access is maximally strided.

// BAD: thread marches down a column-major stride
for (int k = 0; k < N; k++)
    sum += A[threadIdx.x * N + k];   // each thread far from its neighbor

// GOOD: consecutive threads, consecutive addresses
for (int k = 0; k < N; k++)
    sum += A[k * N + threadIdx.x];   // warp reads one contiguous row of A

This single transformation — swapping which index the thread ID drives — is often a 5–10× speedup for memory-bound kernels, and it costs nothing. It is the first thing to check on any slow kernel.

Rule 2: Shared memory, the programmer-managed scratchpad

Coalescing makes each trip to HBM efficient. The next idea is to avoid trips to HBM altogether by keeping reused data on-chip. That on-chip store is shared memory: a small (up to ~228 KB per SM on Hopper), fast (~20–30 cycle latency, roughly 100× lower than HBM) pool that is private to a thread block and explicitly managed by you. Unlike a CPU cache, nothing fills it automatically — you decide what to load, when, and how to index it.

The canonical use is tiling: when many threads in a block need overlapping data, load a tile into shared memory once, then let every thread read it from shared memory many times instead of hammering HBM. Matrix multiply is the textbook case. Naively, computing one output element reads a full row and column from global memory; for an N×N output that re-reads each input value N times. Tiled multiply loads a small sub-block of A and B into shared memory, and each loaded value is reused by an entire tile of threads.

Threads cooperating through shared memory must synchronize. After loading a tile, every thread calls __syncthreads() — a barrier that makes all threads in the block wait until the whole tile is loaded before anyone reads it. Forgetting this barrier is one of the most common GPU bugs: threads read shared memory before their neighbors have written it, producing nondeterministic garbage that only shows up under certain timings.

__shared__ float tileA[32][32];
__shared__ float tileB[32][32];

tileA[ty][tx] = A[...];   // each thread loads one element
tileB[ty][tx] = B[...];
__syncthreads();          // <-- wait for the WHOLE tile

for (int k = 0; k < 32; k++)
    acc += tileA[ty][k] * tileB[k][tx];   // fast shared-mem reads
__syncthreads();          // wait before overwriting the tile next iteration

Bank conflicts: shared memory's one trap

Shared memory is not uniformly fast. It is divided into 32 banks (matching the warp size), each serving one 4-byte word per cycle. If the 32 threads of a warp access 32 different banks, all 32 reads happen in parallel — full speed. But if two or more threads hit the same bank (different addresses), the accesses serialize: an N-way bank conflict takes N times as long.

The classic trigger is a column access of a 32-wide shared array: tile[threadIdx.x][k] with a 32-column tile makes every thread land in the same bank, a 32-way conflict that erases shared memory's advantage. The standard fix is padding — declare the array one column wider ([32][33]) so consecutive rows shift across banks and the column access spreads out.

Rule 3: Occupancy — keeping enough warps resident

Coalescing and shared memory reduce how much you wait on memory. Occupancy determines how well you hide the waiting that remains. Recall from article 1.1: the SM hides memory latency by switching to a ready warp whenever one stalls. That only works if there are ready warps to switch to. Occupancy is the formal measure of how many.

On Ampere/Hopper an SM can hold up to 64 resident warps (2048 threads). If your kernel keeps 32 warps resident, that is 50% occupancy. The maximum is bounded by whichever resource runs out first across all the blocks you want resident:

• Registers. Each SM has a fixed register file (65,536 32-bit registers on Ampere). If each thread uses 64 registers, the SM can host 65,536 / 64 = 1024 threads = 16 warps = 25% occupancy. Register-hungry kernels self-limit.
• Shared memory. Each block claims its shared memory for its whole life. If a block uses 48 KB and the SM has 100 KB available, only 2 blocks fit — regardless of how many warps that is.
• Block size. The SM has a hard cap on resident blocks (e.g. 32) and threads (2048). Tiny blocks can hit the block cap before the thread cap.

NVIDIA ships an Occupancy Calculator (now built into Nsight Compute and the CUDA API cudaOccupancyMaxActiveBlocksPerMultiprocessor) that takes your block size, register count, and shared-memory usage and tells you the theoretical occupancy and what is limiting it. We use it hands-on in article 2.1.

Putting the three rules together

The three rules interact, and the interaction is the whole art. Tiling (shared memory) requires coalesced loads to fill the tile efficiently. Large tiles improve data reuse but consume more shared memory and registers, lowering occupancy. The sweet spot is a kernel that: loads global memory coalesced, stages reused data in conflict-free shared memory, uses enough registers for good per-thread work without spilling, and keeps just enough warps resident to hide the residual latency.

Revisiting divergence with data layout in mind

Article 1.1 introduced warp divergence; data layout gives you the lever to control it. Since divergence only costs you when lanes within a warp disagree, you can often eliminate it by sorting or grouping data so that a warp's 32 elements take the same branch. A ray tracer that sorts rays by which object they hit, or an inference batch that groups sequences of similar length, converts random per-lane branching into per-warp uniform branching — same logic, no serialization. This "data reordering to align with warps" theme returns in the optimization-techniques literature in article 2.1.

A note on registers and spilling

Registers are the fastest memory — one cycle, per-thread-private. The compiler allocates them automatically, but the count per thread is the single biggest driver of occupancy. When a kernel needs more registers than are available at your target occupancy, the compiler spills excess values to "local memory," which despite the name lives in slow global HBM (cached in L1/L2). Spilling can quietly tank performance. You can see per-thread register counts with nvcc -Xptxas -v or in Nsight Compute, and cap them with __launch_bounds__ or -maxrregcount — trading register pressure against occupancy deliberately rather than letting the compiler guess.

FAQ

Takeaways

• Coalesce: map thread i to element i so a warp's 32 reads become one transaction. Prefer Struct-of-Arrays.
• Tile: stage reused data in shared memory, synchronize with __syncthreads(), and collapse HBM traffic by the tile width.
• Avoid bank conflicts with padding ([N][N+1]).
• Occupancy = active/max warps; it is a min() over registers, shared memory, and block caps. Aim for enough, not maximum.
• Registers drive occupancy; spilling to local memory (slow HBM) is a silent killer — watch the per-thread count.
• Use data layout to align branches with warps and kill divergence.

Phase 1 is done — you have the machine and the rules it rewards. Phase 2 opens with article 2.1: turning these rules into a repeatable optimization method using the Roofline model, the optimization-technique taxonomy from the research literature, and the occupancy calculator in anger.

References & further reading

• CUDA C++ Best Practices Guide — Memory Optimizations — the authoritative source on coalescing, shared memory, and bank conflicts.
• Kirk & Hwu — Programming Massively Parallel Processors — the tiled matrix-multiply derivation in full.
• NVIDIA Dev Blog — Using Shared Memory in CUDA — worked examples of tiling and the padding trick.
• NVIDIA Dev Blog — Efficient Global Memory Access — coalescing transactions visualized.
• Nsight Compute Occupancy section — the modern home of the Occupancy Calculator.