Memory Hierarchy — Registers, Caches, HBM, GDDR, and Why Bandwidth Wins

Topics We'll Cover

Twelve slides on every level of the NVIDIA GPU memory hierarchy — how big each tier is, how fast, what it costs to miss it, and why the last level (HBM, GDDR, LPDDR) is the one that decides single-stream LLM tok/s.

• The Hierarchy in One Picture
• Registers — The Real Bandwidth Story
• Shared Memory + L1 — Configurable Per Kernel
• L2 Cache — The GPU-Wide Backstop
• TMA — Tensor Memory Accelerator (Hopper+)
• GDDR — The Consumer Memory
• HBM — The Datacenter Memory
• LPDDR5x + Unified Memory (Grace Systems)
• Bandwidth Ladder Visualised
• What Bandwidth Buys You — The Decode-Bound Insight
• Interactive: Decode-Speed Estimator

The Hierarchy in One Picture

Five tiers, five orders of magnitude. The trade is the usual one: as you climb down the pyramid, capacity grows by 10× or more, but latency grows by 10× and bandwidth drops by 4–10×. Everything kernel authors do — tiling, double-buffering, async copies, tensor cores — is in service of keeping the working set as high in the pyramid as possible.

The five tiers, summarised

Registers — The Real Bandwidth Story

Each SM has a 256 KB register file, split across four schedulers (partitions) of 64 KB each. That is — per chip — the largest, fastest, and most-bandwidth memory you have. An H100 with 132 SMs carries roughly 33 MB of register file on chip, two-thirds the size of its 50 MB L2.

Register pressure and occupancy

The register file is partitioned at kernel-launch time:

# Active threads per SM, given regs/thread
threads_per_SM = floor( 65536 regs / regs_per_thread )

# 65536 = 256 KB / 4 bytes per register, per partition × 4 partitions
# A warp = 32 threads. Active warps = threads_per_SM / 32.

# Examples
regs_per_thread = 32  → 2048 threads → 64 warps  # max occupancy on most arches
regs_per_thread = 64  → 1024 threads → 32 warps  # half-occupancy
regs_per_thread = 128 → 512  threads → 16 warps  # quarter
regs_per_thread = 255 → 256  threads → 8  warps  # tensor-core kernel

High-occupancy vs tensor-core kernels

• High occupancy — Kernels with ≤ 32 regs/thread keep many warps in flight, which lets the SM hide latency by switching warps when one stalls. Good for memory-bound work, simple math, embeddings.
• Tensor-core kernels — Modern matmul kernels (CUTLASS, cuBLASLt, FlashAttention) routinely use 128–255 regs/thread. The accumulators for a 64×64 or 128×128 tile sit entirely in registers. Occupancy is intentionally low because each warp is doing more work per cycle and the math itself does the latency hiding.
• Spilling — Exceed the register file and the compiler "spills" to local memory (which is just a slice of L1/L2/HBM marked as thread-private). Spills are catastrophic on hot paths — they convert register access into HBM access. nvcc -Xptxas -v tells you how many regs/thread and how many spill bytes.

Reading the ptxas verbose output

ptxas info    : 132 bytes gmem
ptxas info    : Compiling entry function '_Z6my_gemmPfS_S_' for 'sm_90'
ptxas info    : Function properties for _Z6my_gemmPfS_S_
                0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info    : Used 168 registers, 8704 bytes smem, 408 bytes cmem[0]
                                ^^^^                ^^^^
                          regs/thread        shared memory/block

168 regs/thread × 32 threads/warp = 5376 regs/warp. With four schedulers' worth of 16384 regs each (65536 total per SM), this kernel can host 65536 / 5376 ≈ 12 warps per SM — well below the 64-warp ceiling. That is fine for a tensor-core kernel where each warp is doing a 64×64×16 wgmma.async; bad if you intended a memory-bound kernel.

Shared Memory + L1 — Configurable Per Kernel

Pre-Volta, shared memory and L1 cache were separate SRAMs in the SM. Volta unified them into a single block of fast on-chip storage that the kernel can configure at launch time: how much for shared, how much for L1.

Sizes per generation

So a Hopper SM with the largest split is L1 = 28 KB, Shmem = 228 KB; tile-heavy kernels run that way. A graphics-style kernel that benefits from caching reads will pick a bigger L1 instead.

Bank conflicts — the one footgun

Shared memory is split into 32 banks of 4 bytes. A warp issues 32 lanes per cycle; if all 32 lanes hit different banks, the access completes in one cycle. If two lanes hit the same bank at different addresses, the access serialises — a 2-way conflict costs 2×, a 32-way conflict costs 32×. Same address on the same bank is fine: it is broadcast.

Common uses

• Tile-based matmul — load a 64×64 or 128×128 sub-tile of A and B once, reuse it across many warp-level MMAs.
• Reductions and prefix sums — warp shuffles handle 32 lanes; shmem stitches across warps inside a block.
• Stencils, transposes, FFT butterflies — anything with non-coalesced access patterns gets staged in shmem.
• Distributed Shared Memory (Hopper+) — threads in one block can read shmem of another block in the same cluster, over the SM-to-SM network. Up to 8 blocks per cluster on H100.

L2 Cache — The GPU-Wide Backstop

L2 sits between the SMs and the memory controllers. Every HBM/GDDR transaction traverses L2; everything an SM evicts from its L1/shmem ends up there. Hits cost ~200 cycles instead of ~500–700 for an HBM round trip, and the bandwidth is roughly an order of magnitude higher.

L2 size by family

Why it matters — per workload

L2 residency control (Ampere+)

Ampere introduced explicit hints for L2: a kernel can mark a buffer as cudaAccessPropertyPersisting so the cache prefers to keep it resident across kernel launches. Useful for KV caches and constant tables that get touched by every kernel in a pipeline.

cudaStreamAttrValue attr = {};
attr.accessPolicyWindow.base_ptr  = kv_cache_ptr;
attr.accessPolicyWindow.num_bytes = 32 * 1024 * 1024;  // 32 MB window
attr.accessPolicyWindow.hitRatio  = 1.0f;
attr.accessPolicyWindow.hitProp   = cudaAccessPropertyPersisting;
attr.accessPolicyWindow.missProp  = cudaAccessPropertyStreaming;
cudaStreamSetAttribute(stream, cudaStreamAttributeAccessPolicyWindow, &attr);

The driver sets aside a fraction of L2 (configurable via cudaDeviceSetLimit(cudaLimitPersistingL2CacheSize, ...)) for the persisting set; the rest serves normal streaming traffic. Get this wrong and you starve the streaming path; get it right on a workload that actually has hot reuse and you can save ~10–20% bandwidth.

TMA — Tensor Memory Accelerator (Hopper+)

The single biggest memory-system change since unified shmem. Pre-Hopper, every thread that wanted to bring data from HBM into shared memory issued its own load instruction. A 128×128 BF16 tile is 32 KB, 16,384 elements — a warp of 32 threads needs 512 instructions just to issue the loads, plus address arithmetic, plus the bookkeeping to coalesce.

What TMA actually does

• The CPU (or kernel) builds a tile descriptor: base pointer, multi-dim shape, strides, element type, swizzle mode.
• One thread issues a single cp.async.bulk.tensor instruction with the descriptor and a destination shared-memory address.
• A dedicated TMA engine on the SM walks the tile in HBM/L2, reorders into the swizzle layout shmem expects, and writes the whole tile.
• The SM is freed to keep computing on the previous tile. A barrier (mbarrier) signals when the new tile lands.

Why it is a big deal for LLMs

• Instruction-issue savings — the warp scheduler is freed from address arithmetic and tile loads, leaving more issue slots for tensor-core MMAs.
• True async double-buffer — while tile n is computing, tile n+1 is loading. FlashAttention-3 is built on this.
• Cluster-wide sharing — one TMA can write a tile that all blocks in a thread-block cluster consume via Distributed Shared Memory. One HBM read, eight blocks of compute.
• Multi-cast — a single TMA can broadcast the same tile to multiple SMs in the cluster — particularly useful for the B-matrix in batched matmul.

GDDR — The Consumer Memory

GDDR is the descendant of DDR for graphics: faster signalling per pin, slightly looser latency, much wider point-to-point buses. It lives on the PCB next to the GPU die, soldered on a 256- or 384-bit bus.

Generation by generation

Signalling, briefly

• NRZ / DDR — 1 bit per UI; the obvious mapping. Transition density limits the per-pin rate.
• PAM4 — 4 voltage levels, 2 bits per UI. Doubles capacity at the same baud but eye-height is smaller; needs sophisticated equalisation, careful PCB design, and better thermals. GDDR6X uses it.
• PAM3 — 3 levels, ~1.58 bits per UI. A compromise: less efficient than PAM4 but with much larger noise margins. GDDR7's choice was driven by the consumer-card SI/PI constraints — PAM4 is hard at 28+ Gbps without an interposer.

Bandwidth maths

Aggregate bandwidth = bus_width × per-pin_rate / 8. The 4090's GDDR6X runs at 21 Gbps on a 384-bit bus → 384 × 21e9 / 8 = ~1008 GB/s. The 5090's GDDR7 runs at 28 Gbps on a 512-bit bus → 512 × 28e9 / 8 = ~1792 GB/s. Two levers, neither of them changes lightly: bus width is set by the die's memory-PHY count, per-pin rate is set by the signalling tech.

What ECC costs you on consumer GDDR

Datacenter cards (A100, H100, L40S, RTX 6000 Ada) have ECC; the consumer 4090/5090 do not. On HBM, ECC is "side-band" — extra dies in the stack store the ECC bits without consuming user bandwidth. On GDDR, ECC is "in-line" — ECC consumes ~6.25% of the user bandwidth and capacity, which is why Quadro / RTX 6000 GDDR cards advertise slightly less peak BW than their consumer twins running at the same per-pin rate.

HBM — The Datacenter Memory

HBM (High-Bandwidth Memory) is 3D-stacked DRAM — 8 to 12 dies stacked vertically, connected through the stack with through-silicon vias (TSVs), and bonded to a base logic die. The whole stack sits on a silicon interposer right next to the GPU die, usually under the same heatspreader.

The trade compared to GDDR

• Bus width — HBM is 1024+ bits per stack, 5–8 stacks per GPU → 5,000–8,000-bit aggregate bus. GDDR is 384–512 bits.
• Frequency — HBM runs much slower per pin (a few Gbps) — the wide bus does the work. Less SI/PI pain at the price of expensive packaging.
• Power — HBM moves ~3–5× the bandwidth at ~half the energy/bit of GDDR. Critical at datacenter scale.
• Cost — an HBM stack costs more than the entire GDDR PCB on a consumer card. The interposer alone is expensive silicon.
• Capacity — goes up by adding dies per stack; H200 (141 GB) = 6 stacks of 8-high HBM3e 24Gb dies, B200 (192 GB) = 8 stacks of 8-high HBM3e 24Gb dies.

HBM in the wild

LPDDR5x + Unified Memory (Grace Systems)

The third memory tier in the modern lineup is the one most people miss. NVIDIA's Grace CPU pairs with Hopper or Blackwell GPUs over NVLink-C2C, a 900 GB/s coherent chip-to-chip link, and lets the GPU treat the CPU's LPDDR5x as EGM — Extended GPU Memory.

What this gives you

• Huge effective capacity — a GH200 has 96 GB HBM3 + 480 GB LPDDR5x = 576 GB the GPU can address. A GB200 superchip (1 Grace + 2 B200) has 384 GB HBM3e + 480 GB LPDDR5x = 864 GB.
• Fast load times — load a 405B model in a few seconds because LPDDR is right next to the system, not over PCIe.
• Coherent paging — pages move on demand. Hot pages migrate to HBM under driver control; cold pages stay in LPDDR. Modest BW (480 GB/s LPDDR5x) but huge capacity.
• No PCIe bottleneck — C2C is 7× PCIe 5.0 x16. The CPU/GPU look like one address space.

Where it appears

Bandwidth Ladder Visualised

Memory bandwidth from low to high, across every memory technology covered. The shape of this chart is the shape of single-stream tok/s on a fixed model size.

Same shape, different axis

If you replace the x-axis with "single-stream 70B FP8 decode tok/s" you get almost the same shape, divided by 70 (since FP8 is 1 byte/param and 70B is 70 GB). That is the entire content of the next slide.

What Bandwidth Buys You — The Decode-Bound Insight

Autoregressive decoding generates one token at a time. To produce that token, the model reads every weight, exactly once, multiplies by the previous activation, accumulates, and writes the next token. The arithmetic intensity of decode is ~1 (one multiply-add per byte read). That is firmly memory-bound on every modern GPU.

tok/s ≈ min( bandwidth / weight_bytes,  compute_limit )

# For a decode-bound model (compute_limit very high), the right term drops:
tok/s ≈ bandwidth / weight_bytes

# weight_bytes = num_params × bytes_per_param
#   FP16/BF16 → 2.0 bytes/param
#   FP8       → 1.0 bytes/param
#   INT4 / q4 → 0.5 bytes/param
#   MX-FP4    → 0.5 bytes/param + tiny scale overhead

Worked examples — 70B model

Where compute starts to matter

• Prefill — Reading a prompt of 4 K tokens is fundamentally a batched matmul (4 K rows × weight matrix). Arithmetic intensity climbs into the hundreds; tensor cores run flat out; bandwidth stops being the bottleneck.
• Big-batch decode — If you batch 32 concurrent decode streams together, every weight read is shared across 32 streams. Effective intensity goes from 1 to 32, and you re-enter compute-bound territory.
• MoE — Only the routed experts get read per token. A 405B-param model with 8 experts of 50B activated per token reads "only" 50B of weights — not 405B. Sparsity changes the bandwidth requirement, not the compute one.

The KV cache, the other big bandwidth consumer

So far we have only counted weight reads. There is a second per-token cost: reading the KV cache — the past keys and values for every token in the context window. Per layer, per head, this is 2 × context_len × head_dim × bytes_per_elem. For a 70B model with 64 layers, 8 KV heads, 128 head_dim, BF16, 4 K context: 2 × 4096 × 1024 × 2 × 64 = ~1 GB per stream. At 32 K context that becomes 8 GB.

• So the actual decode bandwidth requirement is weight_bytes + kv_bytes_per_token × context_len. At long context, KV reads can equal or exceed weight reads.
• This is why FP8 KV cache is now standard in vLLM and TRT-LLM — it halves the second term and is essentially free in quality if done with per-token scales.
• It is also why GQA (grouped-query attention) and MLA (multi-head latent attention) matter so much for serving — they cut the KV cache by 4–16× without changing weight count.

Interactive: Decode-Speed Estimator

Pick a model size, a precision, and a GPU. The estimator computes the weight bytes, checks they fit in VRAM, and applies tok/s ≈ bandwidth / weight_bytes.

Prefill is FLOPS-bound and behaves very differently — a 4 K-token prompt on a 70B FP8 model exercises the tensor cores, not the memory bus, and a B200 will out-prefill an H200 by roughly the FLOPS ratio (~2×) regardless of the bandwidth gap. Big-batch serving is similar: at batch 32+, weight reads are amortised across streams and the workload re-enters compute-bound territory. The estimator below is for the single-stream decode case where bandwidth dominates.