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):

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.