gpu memory hierarchy
**GPU Memory Hierarchy** is the **multi-level memory system in modern GPUs — from registers through shared memory/L1 cache, L2 cache, and HBM/GDDR main memory — that trades off capacity for bandwidth and latency** at each level, where understanding and exploiting this hierarchy is essential for achieving peak performance because GPU workloads are almost always memory-bandwidth-bound.
**NVIDIA A100 Memory Hierarchy**
| Level | Capacity | Bandwidth | Latency | Scope |
|-------|---------|-----------|---------|-------|
| Registers | 256 KB/SM (65536 × 32-bit) | ~20 TB/s (per SM) | 0 cycles | Per-thread |
| Shared Memory / L1 | 164 KB/SM (configurable) | ~19 TB/s (per SM) | ~20-30 cycles | Per-block (shared), per-SM (L1) |
| L2 Cache | 40 MB (total) | ~5 TB/s | ~200 cycles | Global (all SMs) |
| HBM2e (Main Memory) | 80 GB | 2 TB/s | ~400-600 cycles | Global |
**Register File**
- Fastest memory on GPU — zero latency operand access.
- 256 KB per SM × 108 SMs = ~27 MB total register file on A100.
- Register pressure: More registers per thread → fewer active warps → lower occupancy.
- **Register spilling**: When kernel uses too many registers → compiler spills to local memory (slow!).
**Shared Memory / L1 Cache**
- **Shared Memory**: Explicitly managed by programmer — `__shared__` in CUDA.
- **L1 Cache**: Hardware-managed cache for global memory accesses.
- A100: Combined 192 KB per SM, configurable split (e.g., 164 KB shared + 28 KB L1).
- Shared memory: ~19 TB/s bandwidth (32 banks, 4 bytes each, per cycle) — 30x faster than HBM.
**L2 Cache**
- Shared across all SMs. A100: 40 MB. H100: 50 MB.
- Caches global memory accesses — reduces HBM traffic.
- **L2 Cache Residency Control**: CUDA allows pinning data in L2 for persistent access.
- Important for: Reused data that doesn't fit in L1 but is accessed by many blocks.
**HBM (High Bandwidth Memory)**
- Main GPU memory. A100: 80 GB HBM2e at 2 TB/s. H100: 80 GB HBM3 at 3.35 TB/s.
- HBM uses 3D stacking of DRAM dies on silicon interposer adjacent to GPU die.
- Despite "high bandwidth" name: HBM bandwidth is still the bottleneck for most GPU kernels.
**Memory Access Optimization**
| Technique | How | Benefit |
|-----------|-----|--------|
| Coalesced access | Adjacent threads access adjacent addresses | Full memory transaction utilization |
| Shared memory tiling | Load tile into shared memory, compute from there | Replace many global reads with one |
| Register reuse | Keep values in registers across loop iterations | Avoid memory access entirely |
| L2 persistence | Pin working set in L2 | Avoid HBM accesses for reused data |
| Prefetching | `__ldg()` or async copy | Hide memory latency |
**Arithmetic Intensity**
- $\text{Arithmetic Intensity} = \frac{\text{FLOPs}}{\text{Bytes transferred}}$
- If AI < machine's ops:byte ratio → memory-bound → optimize memory access.
- A100: 312 TFLOPS FP16 / 2 TB/s = 156 ops/byte → most kernels are memory-bound.
The GPU memory hierarchy is **the single most important architectural concept for GPU performance optimization** — nearly every GPU kernel is limited by memory bandwidth rather than compute, making the ability to effectively use registers, shared memory, and cache the differentiating skill between mediocre and expert GPU programming.