The memory hierarchy: coalescing, shared memory, banks

1. The ladder

Fast memory is small, and big memory is slow. The reason is physical: a bit you can read in one cycle has to sit close to the arithmetic units, and there is only so much room close to them. DRAM stores a bit as a charge in a tiny capacitor that must drive a long wire before anyone can sense it, which takes tens of nanoseconds per access no matter how fast the rest of the chip is. So every processor builds a ladder of memories, and the GPU’s has five rungs:

Two numbers on this ladder matter more than the rest. The first is the ratio between compute and DRAM that chapter 1 turned into the ridge point: your SMs can do about 46 FP32 operations in the time DRAM delivers one byte. The second is how much faster the on-chip memories are. Shared memory has 32 banks per SM, each delivering 4 bytes per cycle: 128 bytes per cycle per SM, which across 48 SMs at 2.5 GHz is roughly 15 TB/s, over 20× DRAM. Registers are faster still. Every optimization in this chapter is a way to serve more of your loads from the upper rungs.

2. Memory spaces in CUDA

CUDA C++ names memory by who can see it and how long it lives, which is not the same thing as where it physically sits. The mapping is worth knowing exactly, because one row of it is a trap:

“Local” memory is not local

The name is the trap. Local memory is private to a thread, but it lives in DRAM, cached in L1 and L2 like any global access. Two things put data there. A per-thread array indexed with a value the compiler cannot resolve at compile time (a[i] where i comes from a loop it cannot unroll) goes to local memory, because registers cannot be indexed dynamically. And when a kernel needs more registers than it may use, the compiler spills the excess to local memory. That is exactly what lab 1’s experiment D showed: __launch_bounds__(256, 6) capped the tile kernel at 40 registers, and ptxas reported 428 bytes spill stores, 504 bytes spill loads: hundreds of bytes of extra memory traffic per thread, created by the compiler. Whenever -Xptxas -v prints a nonzero stack frame or spill, you are using local memory.

Constant memory and the read-only path

__constant__ data has its own small cache designed for one case: every thread of a warp reading the same address, which is served in one go. When lanes read different addresses, the accesses are serialized, so it suits kernel parameters and small lookup tables indexed uniformly, not arrays indexed by thread. A more generally useful cousin is the read-only data path. Mark a kernel’s input const float* __restrict__, promising the compiler that nothing else writes to it during the kernel, and it may load it with a non-coherent read-only instruction. You can see it in lab 2’s SASS: the copy kernels load with LDG.E.CONSTANT, the pointer-chasing kernel without __restrict__ with plain LDG.E. Get in the habit of writing const T* __restrict__ for inputs.

One pool, two uses

On your card each SM has a single 128 KB array that serves as both L1 cache and shared memory. The split, the carveout, is chosen per kernel from a fixed menu (0, 8, 16, 32, 64 or 100 KB of shared memory on compute capability 12.x), and whatever is not shared memory is L1. A kernel that declares a lot of shared memory therefore shrinks L1 for itself, and a kernel with none gets the whole 128 KB as cache. Two limits to remember: a single block can use at most 99 KB, and anything above 48 KB must be allocated dynamically and explicitly allowed with cudaFuncSetAttribute(kernel, cudaFuncAttributeMaxDynamicSharedMemorySize, bytes), a compatibility rule that trips everyone once.

3. How a warp talks to memory: sectors

Chapter 2 established that a warp executes one instruction for 32 threads at once. For a load, that means one instruction produces 32 addresses at the same moment, and the hardware serves them together. The rule that decides the cost is short, and it is the most important rule in this chapter:

So the question for every memory access in a kernel is: how many sectors do these 32 addresses span? With 4-byte floats, the answers bracket a factor of 8:

• 32 threads reading 32 consecutive floats: 128 bytes, 4 sectors, every byte used. This is a coalesced access.
• 32 threads reading floats 8 or more apart: each lands in its own sector, 32 sectors (1,024 bytes) moved for 128 bytes used. 12.5% efficiency.

Chapter 2’s uniform/varying lens predicts the pattern without drawing anything. In an index expression, find where the varying part (threadIdx.x, or anything derived from it) sits. If it is added at unit stride to a uniform base, consecutive lanes hit consecutive addresses: coalesced. If it is multiplied by a stride, the lanes spread out, and the sector count grows with the stride until every lane has its own sector. If the index is entirely uniform, all 32 lanes read the same address, which is a broadcast: one sector serves everyone, and it is cheap. The campaign’s level 1 walks this lens through a naive matrix multiply, where a two-letter swap of indices is worth 6×.

4. Access patterns in real code

Rows and columns

A row-major matrix stores M[r][c] at r × width + c. Walking along a row is unit stride and walking down a column is a stride of width, typically thousands of floats: every lane in its own sector, 12.5% efficiency. Hence the rule you already met in chapter 2’s 2D grids: threadIdx.x must drive the last, fastest-varying index of every large array you touch. When one array must be read by rows and another by columns (a transpose, a matrix multiply with a transposed operand), no single mapping of threads makes both coalesced. That is what shared memory is for (section 8).

Structures: AoS versus SoA

Suppose each particle is struct { float x, y, z, w; } and you store an array of structures (AoS). A kernel that only needs x reads every 16 bytes: stride 4 in floats, 16 sectors per warp for 128 useful bytes, 25% efficiency. Store a structure of arrays (SoA) instead, one array per field, and the same kernel reads x[i] contiguously: 100%. The general rule: lay data out by the order in which neighbouring threads read it, not by the order a human would group it. (If a kernel really does use all four fields, AoS with a single 16-byte load per thread is also perfectly efficient; that is section 5.)

Misalignment costs little now

A warp reading 128 contiguous bytes that start 12 bytes into a sector spans 5 sectors instead of 4: 80% efficiency for that request. Older GPUs punished misalignment much harder; modern ones only charge the extra sector, and neighbouring warps often find it in L2 anyway. Lab 2’s experiment B2 shifts a copy by 1 to 32 floats; expect a few percent, not a cliff.

Gathers: where ML meets the sector rule

Some reads are inherently scattered. The first layer of every language model, the embedding lookup, reads row token_id of a big table for each token: random rows. It is still fast, because the row is contiguous: if a warp (or a whole block) cooperates to read one row, every access is coalesced, and the randomness only decides which rows. The lesson generalizes: when you cannot make the pattern regular, make the unit of randomness large. Sparse matrices, KV-cache pages (Part IV) and mixture-of-experts routing are all designed around this.

5. Bytes in flight: wide loads and ILP

Coalescing decides how many bytes each request moves. It does not decide how many requests are outstanding, and chapter 1 showed that this matters just as much. Little’s law asked for about 336 KB in flight to saturate 672 GB/s at ~500 ns of latency, about 7 KB per SM, and noticed that 1,536 threads per SM, each waiting on one 4-byte load, only put 6 KB in flight. One float per thread at a time is not enough, even at full occupancy. Chapter 2 named the two ways out. Here they are in code.

Wider loads

A thread can load 8 or 16 bytes with one instruction: float2 or float4. The warp then moves 512 bytes per request instead of 128, still perfectly coalesced (16 sectors, all used), and each thread has four times the bytes in flight.

// n4 = n / 4: each thread copies 16 bytes with a single load and a single store
__global__ void copy_f4(const float4* __restrict__ in, float4* __restrict__ out, size_t n4)
{
    size_t i = (size_t)blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n4) out[i] = in[i];          // SASS: LDG.E.128 / STG.E.128
}

The price is alignment: a float4 access must be 16-byte aligned, so the pointer must be (cudaMalloc returns 256-byte-aligned memory) and the array length must be a multiple of 4, or you handle the tail separately. Vectorized loads are everywhere in production kernels for this reason.

More independent loads per thread

The other way keeps 4-byte loads but issues several before using any of them. Each is independent, so the hardware can have all of them in flight at once:

float x0 = in[i], x1 = in[i + s], x2 = in[i + 2*s], x3 = in[i + 3*s];  // 4 loads in flight
out[i] = x0;  out[i + s] = x1;  out[i + 2*s] = x2;  out[i + 3*s] = x3;

Each of the four loads is still coalesced (consecutive threads, consecutive i), and the thread now waits once for four loads instead of four times for one. This is instruction-level parallelism (ILP), and it is what compilers produce when they unroll a loop. You rarely write it by hand like this, but you will write code whose structure allows it: loops with independent iterations and enough registers to hold the results.

Here is the same arithmetic as a machine. Set the two knobs to the cases of experiment C and watch how full the loop gets.

6. Shared memory: the cache you manage

L1 and L2 cache data automatically, and they help. But thousands of threads share each SM’s L1, so the data one warp is about to reuse is often evicted by others before it gets there: a GPU cache is a poor place to count on reuse. Shared memory is the alternative: a scratchpad in the same 128 KB pool that a block allocates for itself and fills explicitly. Nothing is evicted unless you overwrite it. Its latency is close to L1’s, its bandwidth is over 20× DRAM, and every thread of the block can read what any other thread wrote there, after a barrier.

Why reuse is the whole game

A matrix multiply is where the argument is cleanest (Part II builds the real kernel; the campaign’s level 2 plays it out). In the naive kernel, each thread computing one output element runs a loop that loads one element of A and one of B (8 bytes) for one multiply-add (2 FLOPs). That is 0.25 FLOPs per byte from DRAM, a 184th of your card’s ridge point: at 672 GB/s it is capped at about 170 GFLOPS, half a percent of the 30.9 TFLOPS peak. Yet every element of A is used by a whole row of outputs, and every element of B by a whole column. The data is reused thousands of times; the naive kernel just fetches it again every time.

Tiling fixes that. A block computing a 32×32 tile of outputs loads a 32×32 tile of A and one of B into shared memory (each thread loads one element of each), then every thread reads 32 values from each shared tile. Each byte fetched from DRAM is now used 32 times, so the intensity rises 32-fold, to 8 FLOPs per byte. That is still below the ridge of 46; Part II gets the rest by tiling a second time, in registers.

Watch it happen on a block small enough to follow: 16 threads, a 4×4 tile of outputs, dot products of length 8. Run the naive kernel first, then the tiled one, and compare the shades in DRAM.

The two barriers

Here is the skeleton of every tiled kernel. Click through it: both __syncthreads() are necessary, for two different reasons.

What shared memory costs

Shared memory counts against occupancy (chapter 2): 100 KB per SM, plus 1 KB reserved per block. Spread over the 1,536 threads of a full SM, that is only about 66 bytes per thread on average. The 32×32 tiles above use 8 KB per block, which is harmless; a kernel that wants 64 KB per block runs one block per SM. The trade is the same as with registers: on-chip storage buys reuse and costs parallelism, and the right amount is found by measurement. From Ampere on, GPUs can also copy global memory into shared memory asynchronously, without passing through registers, so that the next tile loads while the current one is used; on your card both the Ampere-style cp.async and the newer bulk-copy engine compile for sm_120 (we checked the SASS). Part II uses them.

7. Bank conflicts

Shared memory is fast because it is divided into 32 banks that work in parallel, each able to deliver one 4-byte word per cycle. Consecutive 4-byte words go to consecutive banks, so the bank of a word is its index modulo 32: s[0] in bank 0, s[31] in bank 31, s[32] in bank 0 again. A warp’s access is served in one pass if its 32 lanes hit 32 different banks. If several lanes want different words in the same bank, those words are fetched one after the other: an N-way bank conflict costs N passes. Lanes reading the same word do not conflict: that word is broadcast to all of them.

So the analysis is the coalescing analysis again, with a different modulus. Take the word index as a function of the lane:

• s[lane]: banks 0…31, one pass.
• s[2*lane]: even banks only, two lanes per bank, 2-way.
• s[3*lane], or any odd stride: still 32 distinct banks, since an odd number has no factor in common with 32. One pass.
• s[32*lane]: every lane in bank 0, 32-way: the access runs at 1/32 speed.

The last case is not exotic. It is what happens when a warp reads a column of a 32×32 tile, tile[lane][c]: consecutive rows are 32 words apart, so the whole column lives in one bank. And reading columns of a tile is the entire point of the transpose, and of half the loads in a matrix multiply.

Two notes for later. First, padding is not the only fix. Production kernels (CUTLASS, FlashAttention) usually swizzle instead: they permute the column index with an XOR, tile[r][c ^ r], which spreads columns over all banks without wasting memory and keeps rows 16-byte aligned for vector loads. It is the last option in the widget; Part II uses it. Second, accesses wider than 4 bytes are processed in several passes anyway (a warp’s float4 read moves 512 bytes, four times what the 32 banks deliver per cycle), and conflicts are counted within each pass, so the analysis changes slightly for vector types. Nsight Compute counts bank conflicts per instruction, and chapter 5 shows you where.

8. Capstone: the matrix transpose

Transposing a matrix, out[c][r] = in[r][c], does no arithmetic at all. It is pure data movement, which makes it the perfect exam for this chapter: the right yardstick is a plain copy of the same bytes, and anything slower than the copy is wasted memory bandwidth. Lab 2’s experiment D runs three versions on an 8192×8192 matrix (256 MB in, 256 MB out), with a copy as the ceiling. All use blocks of 32×8 threads, each block handling a 32×32 tile (each thread moves four elements), so that a warp is always one row of 32 threads.

Version 1: naive

out[(size_t)c * R + r] = in[(size_t)r * C + c];     // c varies along the warp

The read in[r*C + c] is coalesced: c is the varying index, in the last position. The write out[c*R + r] has the varying c multiplied by R = 8,192: every lane writes to its own sector, 8,192 floats from its neighbour. One coalesced side, one side at 12.5% efficiency. No mapping of threads can fix this, because whichever index you make vary along the warp, one of the two matrices is walked by columns.

Version 2: through shared memory

The fix is to do the transposition inside the SM, where column access is cheap. Read a 32×32 tile with coalesced row reads into shared memory, synchronize, then write the transposed tile with coalesced row writes of the output, reading the tile by columns:

template <int PAD>
__global__ void transpose_shared(const float* __restrict__ in, float* __restrict__ out, int R, int C)
{
    __shared__ float tile[32][32 + PAD];

    int c  = blockIdx.x * 32 + threadIdx.x;                 // 1. load a tile, row by row
    int r0 = blockIdx.y * 32 + threadIdx.y;
    for (int k = 0; k < 32; k += 8)
        if (r0 + k < R && c < C)
            tile[threadIdx.y + k][threadIdx.x] = in[(size_t)(r0 + k) * C + c];

    __syncthreads();                                        // the whole tile is loaded

    int oc  = blockIdx.y * 32 + threadIdx.x;                // 2. write the mirrored tile
    int or0 = blockIdx.x * 32 + threadIdx.y;
    for (int k = 0; k < 32; k += 8)
        if (or0 + k < C && oc < R)
            out[(size_t)(or0 + k) * R + oc] = tile[threadIdx.x][threadIdx.y + k];   // column read
}

Only one barrier this time: each block handles a single tile, so nobody overwrites it (no write-after-read hazard). Both global accesses now have threadIdx.x in the last index. But look at the shared-memory read in the last line: tile[threadIdx.x][…]. The varying index is the row, so with PAD = 0 a warp reads a column of a 32-word-wide tile: section 7’s 32-way bank conflict, on every read out of the tile.

Version 3: padded

Set PAD = 1. The build output shows the cost (4224 bytes smem instead of 4096), and the column read now spreads over all 32 banks. Nothing else changes.

9. Lab 2: measure the hierarchy

cd ~/unibo-lessons/mlsys/labs/lab02-memory
make          # ptxas: the two transpose_shared kernels, 4096 vs 4224 bytes smem
make test     # every kernel against a CPU reference, awkward sizes included
make bench    # A latency · B1 strides · B2 offsets · C bytes in flight · D transpose

Check your understanding

Sources and further reading

• Hwu, Kirk, El Hajj, Programming Massively Parallel Processors, 4th ed., ch. 5 (memory types, tiling, the two barriers, boundary checks, shared memory and occupancy) and ch. 6 (DRAM bursts, coalescing, corner turning, latency hiding through memory parallelism, thread coarsening). All numbers here are recomputed for the RTX 5070; the book does not cover shared-memory banks.
• NVIDIA, CUDA C++ Best Practices Guide, “Coalesced Access to Global Memory” (32-byte transactions, misaligned and strided examples) and “Shared Memory” (32 banks, 32 bits per bank per cycle, broadcast, serialization).
• NVIDIA, CUDA Programming Guide, Compute Capabilities: shared-memory carveouts for 12.x (0, 8, 16, 32, 64, 100 KB), 99 KB per block, 32 banks, 64 KB constant memory.
• NVIDIA, Blackwell Tuning Guide: the 48 KB static shared-memory limit, the dynamic opt-in, 1 KB reserved per block.
• NVIDIA, RTX Blackwell GPU Architecture whitepaper, Appendix C: 128 KB L1/shared and 256 KB of registers per SM, 48 MB L2 on GB205.
• Mark Harris, An Efficient Matrix Transpose in CUDA C/C++ (NVIDIA blog, 2013): the classic 32×8-thread, padded-tile transpose that section 8 follows. His companion post How to Access Global Memory Efficiently has the strided and offset experiments.
• Z. Jia, M. Maggioni, B. Staiger, D. Scarpazza, Dissecting the NVIDIA Volta GPU Architecture via Microbenchmarking (2018): pointer-chasing methodology for measuring each level’s latency, the model for experiment A.