GPU Memory Hierarchy

1. What is it?#

Five levels of storage, each roughly 10x faster and 10x smaller than the one below.

                        size (H100)     bandwidth      latency
Registers               256 KB/SM       ~500 TB/s      ~1 cycle
Shared memory / L1      256 KB/SM       ~19 TB/s       ~30 cycles
L2 cache                50 MB           ~7 TB/s        ~200 cycles
HBM3 (global)           80 GB           3.35 TB/s      ~500 cycles
Host DRAM (over PCIe)   TBs             ~50 GB/s       ~10,000 cycles

The entire art of GPU optimization is keeping your working set one level higher than the naive implementation does.

Diagram — Fast and tiny at the top, vast and slow at the bottom#

flowchart TB
  R["Registers<br/>per thread - fastest, tiny"] --> S["Shared memory / L1<br/>per SM - cooperative scratchpad for a block"]
  S --> L2["L2 cache<br/>whole GPU - tens of MB"]
  L2 --> H[("HBM global memory<br/>the capacity tier - bandwidth is the hard limit")]
  H -->|"PCIe: orders of magnitude slower"| HOST["Host RAM"]
  HOST --> DISK["NVMe / network storage"]

  class R,S compute
  class L2,H memory
  class HOST,DISK io

2. Why does it exist?#

Same reason as the CPU hierarchy (Section II.02): fast memory is expensive and small. But the GPU’s version is more extreme and more programmer-managed: shared memory is not a cache you hope hits, it’s a scratchpad you explicitly fill.

That explicitness is why GPU kernels can be so much better than automatic caching would achieve — and why writing them requires understanding the hierarchy.

3. Simple analogy#

A workshop.

• Registers — what’s in your hands. Instant, tiny.
• Shared memory — your workbench. You choose what to put there. Shared with your team (the block).
• L2 — the shelf at the end of the room, shared by everyone.
• HBM — the stockroom down the corridor.
• Host memory — the warehouse across town.

The CPU’s caches are an automatic assistant who guesses what you’ll need. The GPU’s shared memory is a workbench you load deliberately. More work, more control, better results.

4. Tiny example#

Measure the hierarchy:

// Latency probe: dependent loads defeat prefetching and ILP
__global__ void chase(int* data, int n, int iters, int* out) {
    int idx = 0;
    for (int i = 0; i < iters; i++) idx = data[idx];   // each load depends on the last
    *out = idx;
}

Fill data with a random permutation of increasing sizes and time it. You’ll see clear plateaus at L1 (~30 cycles), L2 (~200), and HBM (~500).

Bandwidth per level, roughly:

# Shared memory bandwidth (via a kernel that only touches shared)
# HBM bandwidth (large copy)
import torch, time
n = 500_000_000
a = torch.empty(n, device='cuda', dtype=torch.float32)
b = torch.empty_like(a)
for _ in range(3): b.copy_(a)
torch.cuda.synchronize(); t0 = time.perf_counter()
for _ in range(10): b.copy_(a)
torch.cuda.synchronize()
print(f"HBM: {2*n*4/((time.perf_counter()-t0)/10)/1e12:.2f} TB/s")

5. Technical explanation#

Registers#

256 KB per SM (65,536 × 32-bit registers), statically partitioned.
Fastest storage. Private per thread.
Compiler-allocated; you influence it with __launch_bounds__ and code structure.

Register spilling happens when a thread needs more registers than allocated. Spilled values go to “local memory” — which is actually a private region of global memory, ~100x slower. Watch for lmem in nvcc --ptxas-options=-v output; nonzero local memory in a hot kernel is usually a bug.

Shared memory#

Up to 228 KB per SM on Hopper (shared with L1; configurable split).
Per-block scope. Explicitly managed.
Organized in 32 banks of 4 bytes.

Bank conflicts: if threads in a warp access different addresses in the same bank, the accesses serialize.

__shared__ float tile[32][32];
tile[threadIdx.x][0]        // BAD: all 32 threads hit bank 0 → 32-way conflict
tile[0][threadIdx.x]        // GOOD: 32 consecutive floats → 32 different banks

__shared__ float tile[32][33];   // padding: now column access is conflict-free
tile[threadIdx.x][0]        // addresses differ by 33 floats → different banks

The [32][33] padding trick is ubiquitous in CUDA code. Now you know why.

Modern alternative: swizzled layouts (XOR the column index with a function of the row), which avoid conflicts without wasting memory. CUTLASS uses these.

L2 cache#

50 MB on H100, shared by all SMs.
Automatically managed, but with hints:

// Persist frequently-reused data in L2 (Ampere+)
cudaStreamSetAttribute(stream, cudaStreamAttributeAccessPolicyWindow, &attr);

For inference, L2 matters most when a small tensor is read by many blocks — e.g. the same weight tile read by many attention heads, or a small model that partly fits in L2.

Global memory (HBM)#

80 GB, 3.35 TB/s, ~500 cycle latency.
Accessed in 32-byte sectors; a warp's 32 accesses are coalesced into as few
sectors as possible (Section 09).

Latency is hidden by having many warps resident. Bandwidth cannot be hidden — it’s the hard limit.

The Hopper additions#

TMA (Tensor Memory Accelerator)
  A hardware unit that copies multi-dimensional tiles between HBM and shared
  memory asynchronously, with one instruction, computing addresses in hardware.
  Frees threads and registers that previously did address arithmetic.
  Used by FlashAttention-3 and modern CUTLASS GEMMs.

DSMEM (Distributed Shared Memory)
  Blocks in the same cluster can read each other's shared memory directly.
  Effectively a larger cooperative scratchpad.

Async copy (cp.async, Ampere+)
  Copy HBM→shared without going through registers, overlapping with compute.
  The basis of software pipelining in modern GEMM kernels.

These three features are why Hopper kernels look different from Ampere kernels, and why FlashAttention-3 is Hopper-specific.

The canonical use: tiled GEMM#

for each output tile:
    accumulator in REGISTERS
    for each k-tile:
        cp.async / TMA: HBM → SHARED (tile of A and B)
        __syncthreads()
        for each warp sub-tile:
            SHARED → REGISTERS (fragments)
            tensor core MMA: registers → registers
    REGISTERS → HBM (once)

Each level does what it’s good at: HBM for capacity, shared for cooperative reuse within a block, registers for the actual arithmetic.

FlashAttention is the same pattern applied to attention: keep the score tile in shared memory/registers, never write it to HBM.

6. Under the hood#

Measure your hierarchy with Nsight Compute:

ncu --set full --section MemoryWorkloadAnalysis ./your_program

Key metrics:

dram__bytes.sum                          bytes from HBM
lts__t_bytes.sum                         bytes from L2
l1tex__t_bytes.sum                       bytes from L1/shared
smsp__inst_executed_op_shared_ld.sum     shared memory loads
l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_ld.sum    ← bank conflicts

The ratio l1tex bytes / dram bytes tells you your cache reuse factor. For a well-tiled GEMM it should be large (10-100x); for a streaming elementwise kernel it should be ~1.

7. Performance implications#

Optimization                        Typical gain
Tiling into shared memory           5-50x for reuse-heavy ops
Eliminating bank conflicts          up to 32x on the affected accesses
Avoiding register spills            2-5x if spilling was severe
Using async copy / TMA              10-30% (better overlap)
L2 persistence hints                5-20% for small hot tensors
Vectorized loads (float4/half8)     1.5-2x on bandwidth-bound kernels

That last one is easy and often overlooked: loading 16 bytes per instruction instead of 4 reduces instruction count and improves memory pipeline efficiency.

8. Production implications#

You mostly consume kernels rather than write them, but:

• Reading a memory-workload profile tells you whether a kernel is well-optimized. A kernel achieving 90% of HBM bandwidth is done; one at 30% has a problem.
• Choosing tensor layouts (Section IV.07) determines whether a kernel can coalesce.
• Understanding shared memory limits explains why certain configurations (very long sequences, unusual head dims) fall back to slow kernels — the fast one’s tiles don’t fit.
• Model architecture affects this: a head_dim of 256 uses twice the shared memory per tile as 128, potentially halving occupancy or forcing a different kernel.

9. Common mistakes#

Ignoring bank conflicts. A 32x slowdown hiding in plain sight.

Register spilling. Check lmem in ptxas output.

Not using vectorized loads. Free 1.5-2x on bandwidth-bound kernels.

Assuming L2 will save you. 50 MB doesn’t hold a 16 GB model.

Forgetting __syncthreads() after filling shared memory. Race condition; intermittent wrong answers.

Using shared memory when registers would do. Registers are 25x faster.

10. Hands-on exercise#

A. Measure the hierarchy. Implement the pointer-chase latency probe. Plot ns/access vs working-set size from 1 KB to 1 GB. Identify each level. Compare to the table in section 1.

B. Bank conflicts. Write a shared-memory transpose with [32][32] and [32][33]. Measure both and check the bank-conflict counter with ncu. Report the ratio.

C. Tiling. Write a naive matmul kernel and a shared-memory-tiled one. Measure both. Compute the arithmetic intensity of each and verify the improvement matches the roofline prediction.

D. Vectorized loads. Write a bandwidth-bound kernel with float loads and with float4 loads. Measure achieved bandwidth for each.

E. Register spilling. Deliberately write a kernel that spills (large local arrays). Confirm with ptxas output and measure the cost.

11. Interview questions#

1. Describe the GPU memory hierarchy with sizes and bandwidths.
2. What is a shared memory bank conflict and how do you fix it?
3. What is register spilling and how do you detect it?
4. What is TMA and why does it matter for modern kernels?
5. How does FlashAttention use the memory hierarchy?
6. Why does head_dim affect which attention kernel you get?
7. How would you tell from a profile whether a kernel is well-optimized for memory?

12. Further reading#

• [REFERENCE] CUDA C++ Best Practices Guide, “Memory Optimizations”
• [REFERENCE] Nsight Compute Memory Workload Analysis documentation
• [ESTABLISHED] CUTLASS documentation on shared memory layouts and swizzling
• Next: 04 — Bandwidth and the roofline on GPU