gpu memory hierarchy optimization

**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**.