GPU Memory Hierarchy Optimization is the practice of strategically utilizing the multi-level memory system of modern GPUs — from fast but small shared memory and L1 cache (20 TB/s, 128 KB per SM) to large but slower global memory (1-3 TB/s, 40-80 GB) — to maximize data reuse, minimize memory latency, and achieve peak computational throughput by keeping data as close to the compute units as possible.

Memory Hierarchy Levels:

• Registers: fastest storage (per-thread private registers, ~20 TB/s effective bandwidth); each SM on NVIDIA Ampere/Hopper has 65,536 32-bit registers shared across all active threads; register spilling to local memory (cached in L1) occurs when kernel uses >255 registers per thread, causing 10-100× slowdown
• Shared Memory/L1 Cache: 128-192 KB per SM configurable between shared memory (programmer-managed) and L1 cache (hardware-managed); shared memory provides 20 TB/s bandwidth with ~20 cycle latency — 10-20× faster than global memory for data shared across thread block
• L2 Cache: 40-50 MB unified cache (A100) or 50 MB (H100) shared across all SMs; 4-6 TB/s bandwidth; automatically caches global memory accesses; residency hints (cudaAccessPolicyWindow) allow programmer control over L2 caching for streaming vs reused data
• Global Memory (HBM): 40-80 GB capacity with 1.5-3 TB/s bandwidth (A100: 1.9 TB/s, H100: 3.35 TB/s); 200-400 cycle latency; all data must initially reside here; optimizing global memory access patterns is the primary performance bottleneck for memory-bound kernels

Shared Memory Programming Patterns:

• Tiling/Blocking: decompose computation into tiles that fit in shared memory; load tile from global memory cooperatively, compute on tile data (reused many times), write results back; matrix multiplication achieves 10-20× speedup by reusing each matrix element across multiple dot products
• Cooperative Loading: threads in a block collaboratively load data into shared memory using coalesced access patterns; each thread loads one or more elements; __syncthreads() barrier ensures all data is loaded before computation begins
• Reduction Trees: parallel reduction (sum, max, min) uses shared memory to accumulate partial results; each iteration halves active threads and combines pairs; log₂(N) iterations reduce N elements with O(N) work instead of O(N²) atomic operations to global memory
• Halo Regions: stencil computations load neighboring elements (halo) into shared memory along with the tile; enables each thread to access neighbors without additional global memory reads; 3D stencils with radius R require loading (TILE_SIZE + 2R)³ elements for TILE_SIZE³ output

Memory Access Optimization:

• Coalescing: threads in a warp accessing consecutive memory addresses (stride-1 pattern) are coalesced into a single 128-byte transaction; non-coalesced access (stride > 1, random access) generates 32 separate transactions — 32× bandwidth waste; structure-of-arrays (SoA) layout enables coalescing vs array-of-structures (AoS)
• Bank Conflict Avoidance: shared memory is divided into 32 banks (4-byte width); simultaneous access to the same bank by multiple threads serializes the access; padding arrays by 1 element (e.g., [TILE_SIZE][TILE_SIZE+1]) shifts columns to different banks, eliminating conflicts in transpose operations
• Alignment: global memory transactions are 32, 64, or 128 bytes; misaligned access (address not multiple of transaction size) requires multiple transactions; cudaMalloc guarantees 256-byte alignment; manual allocation should align to at least 128 bytes
• Streaming vs Caching: streaming data (accessed once) should bypass L1/L2 to avoid cache pollution; use __ldg() intrinsic or const __restrict__ pointers to hint read-only caching; cudaAccessPolicyWindow API explicitly controls L2 residency for persistent data

Performance Metrics:

• Memory Bandwidth Utilization: achieved_bandwidth / peak_bandwidth; well-optimized kernels reach 70-90% of peak HBM bandwidth; below 50% indicates access pattern issues (non-coalesced, bank conflicts, insufficient parallelism)
• Cache Hit Rates: L1 hit rate >80% and L2 hit rate >60% indicate good data locality; low hit rates suggest working set exceeds cache capacity or poor temporal locality
• Occupancy Impact: higher occupancy (more active warps per SM) hides memory latency through warp scheduling; memory-bound kernels benefit from high occupancy (>50%) to overlap memory access with computation from other warps

GPU memory hierarchy optimization is the most critical factor determining real-world GPU performance — the 100-1000× speed difference between memory levels means that algorithmic changes to improve data locality often provide larger speedups than low-level instruction tuning, making memory access pattern design the primary focus of high-performance GPU programming.