gpu memory hierarchy
**GPU Memory Hierarchy** is the **multi-level storage system in GPU architectures that provides different capacity-bandwidth-latency trade-offs**, from per-thread registers (fastest, smallest) through shared memory and caches to global device memory (slowest, largest) — and understanding this hierarchy is the single most important factor in GPU kernel optimization.
GPU performance is overwhelmingly determined by memory access patterns. A kernel that reads from registers runs at ~100 TB/s effective bandwidth; the same kernel reading from global memory achieves ~1-3 TB/s. The 100x difference between these levels makes memory hierarchy optimization the dominant concern in GPU programming.
**Memory Levels (NVIDIA Architecture)**:
| Memory | Scope | Size | Latency | Bandwidth |
|--------|-------|------|---------|----------|
| **Registers** | Per-thread | ~256 x 32-bit per thread | 0 cycles | ~100+ TB/s |
| **Shared memory** | Per-SM (block) | 48-228 KB | ~20-30 cycles | ~20-100 TB/s |
| **L1 cache** | Per-SM | Unified with shared mem | ~30 cycles | ~20 TB/s |
| **L2 cache** | Chip-wide | 6-96 MB | ~200 cycles | ~6-12 TB/s |
| **Global (HBM)** | Device | 16-80 GB | ~400-600 cycles | 1-3.3 TB/s |
| **Constant memory** | Device, cached | 64 KB + cache | ~5 cycles (hit) | Broadcast to warp |
| **Texture memory** | Device, cached | Through L1/L2 | ~400 cycles (miss) | Spatial locality optimized |
**Register Optimization**: Registers are the fastest storage but are finite per SM (~65K 32-bit registers per SM on modern GPUs). If a kernel uses too many registers, occupancy drops (fewer concurrent warps per SM). **Register spilling** to local memory (which resides in slow global memory, cached through L1) can cause 10-50x slowdown for spilled accesses. Compiler flags (`-maxrregcount`) and algorithmic refactoring (reducing live variables) manage register pressure.
**Shared Memory**: Programmer-managed scratchpad memory shared across threads in a block. Critical for: **data reuse** (load from global memory once, access from shared memory many times — matrix tiling achieves near-peak throughput this way), **inter-thread communication** (threads in the same block exchange data via shared memory + `__syncthreads()`), and **reduction/scan** (tree-based parallel reductions). **Bank conflicts**: shared memory is organized into 32 banks; if multiple threads in a warp access different addresses in the same bank, accesses serialize. Padding shared memory arrays avoids conflicts.
**Memory Coalescing**: Global memory is accessed in transactions (32/64/128 bytes). When threads in a warp access consecutive addresses (stride-1 pattern), the hardware coalesces these into minimal transactions — achieving peak bandwidth. Scattered or strided access patterns cause multiple transactions per warp, wasting bandwidth by up to 32x. **Array-of-Structures to Structure-of-Arrays (AoS→SoA)** transformation is the most common optimization to achieve coalesced access.
**L2 Cache Management**: Modern GPUs (Ampere+) support **L2 cache residency control** (`cudaAccessPolicyWindow`) to pin frequently accessed data in L2, and **L2 persistence** to keep streaming data from evicting resident data. This is critical for workloads with mixed access patterns (frequent small reads + streaming large buffers).
**The GPU memory hierarchy is the defining constraint of GPU programming — every kernel optimization reduces to moving data closer to the compute units and accessing it in patterns that match the hardware, making memory hierarchy mastery the essential skill for achieving peak GPU performance.**