GPU Memory Hierarchy
A GPU has four places a byte can live, and they differ by a thousandfold in bandwidth: registers, shared memory on the SM, a chip-wide L2, and HBM off-chip. Almost every kernel optimization is a decision about which level a value is read from and how many times. Knowing the sizes and bandwidths for an H100 cold is what lets you say why a kernel is slow before you profile it.
TL;DR: On an H100: 65,536 registers per SM at register speed, up to 228 KB of shared memory per SM at roughly 30 TB/s aggregate, a 50 MB L2 shared by all SMs at several TB/s, and 80 GB of HBM3 at 3.35 TB/s. A byte that stays in registers or shared memory is nearly free; a byte that goes to HBM costs about 500 cycles and a slice of a budget the whole chip shares. Fast kernels are the ones that read each HBM byte once and do everything else higher up.
The four levels, with numbers
| Level | Where | Size (H100 SXM) | Bandwidth | Latency | Who shares it |
|---|---|---|---|---|---|
| Registers | inside each SM | 256 KB per SM (65,536 × 32-bit) | register speed | ~1 cycle | one thread |
| Shared memory / L1 | inside each SM | up to 228 KB per SM, configurable split with L1 | on the order of 30 TB/s across the chip | ~20 to 30 cycles | one thread block |
| L2 cache | on the die, between SMs and HBM | 50 MB, in two partitions | several TB/s | ~200 cycles | every SM |
| HBM3 | stacked DRAM next to the die | 80 GB | 3.35 TB/s | ~500 cycles or more | the whole GPU, and the host over PCIe |
Two things about the table matter more than the exact numbers. The bandwidth gap between shared memory and HBM is roughly an order of magnitude, and the capacity gap is roughly four orders. So the game is always the same: bring a tile of data up from HBM once, reuse it from shared memory or registers as many times as the algorithm allows, and write the result back once. Tiling, fusion and FlashAttention are all this one move applied to different operations.
What each level is for
Registers hold a thread's working values and are allocated per thread at compile time. They are the fastest storage on the chip and also the scarcest per thread: a kernel that wants 128 registers per thread can only keep 512 threads resident per SM, a quarter of the maximum, and latency hiding suffers (see GPU Execution Model). When the compiler runs out it "spills" to local memory, which lives in HBM behind the L1, and a spill in an inner loop can double a kernel's time.
Shared memory is the programmer-managed scratchpad that a thread block owns for its lifetime. It is the level that makes cooperation possible: threads load a tile from HBM into it, synchronize with __syncthreads(), then read each other's values. It is physically the same SRAM as the L1 cache, and the split is configurable per kernel (on Hopper up to 228 KB of the 256 KB can be shared memory). It is organized in 32 banks, and two threads hitting the same bank in one instruction serialize; padding a tile by one column is the standard fix (Shared Memory and Bank Conflicts).
The L2 is a hardware-managed cache in front of HBM, shared by every SM. It is large enough to hold a whole KV block or a slice of weights, and Hopper lets a kernel pin lines with residency hints. It matters most in two cases: when several blocks read the same data (an L2 hit costs less than half the latency and none of the HBM bandwidth), and when a kernel's working set almost fits (a 40 MB set on a 50 MB L2 is fast; a 60 MB set thrashes). It is split into two partitions with a crossbar between them, so "50 MB" is not quite one uniform pool.
HBM is where models live. Eighty gigabytes on an H100, 141 on an H200, 180 on a B200, 192 on an MI300X. Its bandwidth is the slope of the roofline and the single number that decides decode throughput: a 70B model in bf16 is 141 GB, and streaming that once through 3.35 TB/s takes 42 ms, so single-stream decode cannot exceed about 24 tokens per second on one card whatever the software does. Capacity decides what fits; bandwidth decides how fast it runs.
Two transactions that explain most slow kernels
A warp's 32 threads issue a load together, and the memory system serves it in 128-byte segments (32-byte sectors within). If the 32 addresses are contiguous 4-byte words, one segment satisfies the whole warp: one transaction, 100% useful bytes. If they are strided by 128 bytes, thirty-two segments are fetched to deliver 128 useful bytes, and the effective bandwidth drops 32-fold. This is memory coalescing, and it is the first thing to check when a kernel is far below the HBM roof.
The second is the round trip. An unfused chain of elementwise ops (multiply, add bias, GELU) writes its intermediate tensor to HBM after each step and reads it back for the next. Three ops means three full passes over the data, all at intensity under 1 FLOP per byte. Keeping the intermediate in registers turns three passes into one, which is why fusion is the standard fix for anything bandwidth-bound and why compilers like Inductor spend most of their effort on it.
Hopper additions worth naming
Hopper added two mechanisms that move data between levels without going through registers. cp.async lets a thread issue a copy from global memory straight into shared memory and keep computing while it lands. The Tensor Memory Accelerator (TMA) goes further: a single thread describes a multi-dimensional tile, and dedicated hardware streams it into shared memory with the address arithmetic done for it. FlashAttention-3's speedup on H100 comes largely from using TMA and warp specialization so that some warps only move data while others only do matrix math. Distributed shared memory lets blocks in a cluster read each other's shared memory directly, a small crack in the "blocks cannot communicate" rule.
What interviewers are listening for
The question is usually "where does the time go in this kernel?", and the strong answer walks the hierarchy: how many bytes come from HBM, whether the loads coalesce, whether the working set fits in shared memory or L2, and whether anything spills. Candidates who only talk about FLOPs have not run a profiler. The number to have ready is the H100 quartet: 228 KB shared per SM, 50 MB L2, 80 GB HBM at 3.35 TB/s. The follow-up in reserve is "what changes on the H200?", and the answer is that the compute die and the on-chip levels are identical, only HBM grew (141 GB, 4.8 TB/s), so memory-bound kernels get 1.4x faster and compute-bound ones do not move.
Key takeaways
- Four levels: registers, shared memory (up to 228 KB per SM), L2 (50 MB), HBM (80 GB at 3.35 TB/s on H100). Bandwidth falls roughly tenfold per level; capacity grows roughly a thousandfold.
- Read each HBM byte once, reuse it from shared memory or registers. Tiling, fusion and FlashAttention are that rule applied three ways.
- Coalescing decides how much of HBM bandwidth a warp's load actually gets; a 128-byte stride wastes 97% of it.
- HBM capacity decides what fits; HBM bandwidth decides decode speed. The H200 changes only those two numbers.
- Hopper's cp.async and TMA move tiles into shared memory without touching registers; FA3 and modern GEMMs depend on them.