gpu memory hierarchy optimization

**GPU Memory Hierarchy Optimization** is **the systematic tuning of data placement and access patterns across GPU's multi-level memory system to maximize bandwidth utilization and minimize latency** — where understanding the hierarchy from registers (20,000 GB/s effective bandwidth) through shared memory (19 TB/s on H100), L1/L2 caches (10-15 TB/s), to global HBM memory (1.5-3 TB/s) enables 5-20× performance improvements through techniques like shared memory tiling that reduces global memory accesses by 80-95%, register blocking that keeps frequently accessed data in fastest storage, and memory coalescing that achieves 80-100% of theoretical bandwidth, making memory hierarchy optimization the most impactful optimization for memory-bound kernels that dominate GPU workloads where 60-80% of kernels are memory-limited rather than compute-limited. **Memory Hierarchy Levels:** - **Registers**: fastest storage; 32-bit registers; 65,536 registers per SM on A100; 20,000+ GB/s effective bandwidth; private to each thread; limited quantity (255 registers per thread max); excessive usage reduces occupancy - **Shared Memory**: on-chip SRAM; 164KB per SM on A100, 228KB on H100; 19 TB/s bandwidth on H100; shared across thread block; explicit programmer control; 32 banks for parallel access; 100× faster than global memory - **L1 Cache**: 128KB per SM on A100; combined with shared memory; automatic caching; benefits from spatial and temporal locality; cache line size 128 bytes; write-through to L2 - **L2 Cache**: 40MB on A100, 50MB on H100; shared across all SMs; 10-15 TB/s bandwidth; benefits from reuse across thread blocks; victim cache for L1; configurable persistence for critical data - **Global Memory**: 40-80GB HBM2/HBM3; 1.5-3 TB/s bandwidth; highest capacity but slowest; 400-800 cycle latency; requires coalescing for efficiency; all threads can access **Shared Memory Optimization:** - **Tiling Strategy**: divide data into tiles that fit in shared memory; load tile cooperatively; reuse across threads; reduces global memory accesses by 80-95%; matrix multiplication: 5-20× speedup with tiling - **Bank Conflicts**: 32 banks on modern GPUs; simultaneous access to same bank serializes; stride by 33 elements to avoid conflicts; padding arrays prevents conflicts; 2-10× slowdown from conflicts - **Cooperative Loading**: all threads in block load data collaboratively; maximizes memory bandwidth; coalesced global loads; synchronize with __syncthreads() after loading - **Double Buffering**: overlap computation with next tile load; use two shared memory buffers; hide memory latency; 20-40% performance improvement; requires careful synchronization - **Capacity Planning**: 48KB per block typical; balance between occupancy and tile size; larger tiles reduce global accesses but limit occupancy; profile to find optimal size **Register Optimization:** - **Register Pressure**: monitor with nvcc --ptxas-options=-v; shows registers per thread; high usage limits occupancy; target 32-64 registers per thread for good occupancy - **Register Spilling**: when exceeding register limit, spills to local memory (slow); 10-100× slowdown for spilled accesses; reduce by simplifying code, using fewer variables - **Loop Unrolling**: #pragma unroll increases register usage but improves ILP; unroll factor 2-4 typical; balance between ILP and occupancy; measure impact with profiler - **Constant Memory**: use __constant__ for read-only data; 64KB per kernel; cached; broadcast to all threads; 2-5× faster than global memory for uniform access - **Texture Memory**: use for spatial locality; 2D/3D access patterns; cached; interpolation hardware; 2-10× speedup for irregular access patterns **Cache Optimization:** - **L1 Cache Hints**: use __ldg() for read-only data; forces L1 caching; improves temporal locality; 20-50% speedup for reused data - **L2 Persistence**: cudaStreamSetAttribute() sets L2 persistence; keeps critical data in L2; benefits data reused across kernels; 30-60% speedup for multi-kernel workloads - **Cache Line Utilization**: 128-byte cache lines; access consecutive data to utilize full line; 4-8× improvement vs scattered access; structure data for sequential access - **Streaming Access**: use streaming loads for data accessed once; bypasses L1 cache; prevents cache pollution; improves performance for other data **Memory Access Patterns:** - **Coalescing**: threads in warp access consecutive addresses; 128-byte aligned; achieves 100% bandwidth; stride-1 access optimal; stride-2 achieves 50%; stride-32 achieves 3% - **Structure of Arrays (SoA)**: prefer SoA over AoS; enables coalesced access; 5-10× memory bandwidth improvement; example: x[N], y[N], z[N] instead of point[N].x, point[N].y, point[N].z - **Alignment**: align data to 128 bytes; cudaMalloc provides automatic alignment; manual alignment with __align__(128); misalignment causes 2-10× slowdown - **Padding**: add padding to avoid bank conflicts and improve coalescing; 1-2 elements padding typical; 10-30% performance improvement **Bandwidth Optimization:** - **Measure Bandwidth**: use Nsight Compute; reports achieved bandwidth vs peak; target 80-100% for memory-bound kernels; identifies bottlenecks - **Vectorized Loads**: use float4, int4 for 128-bit loads; 2-4× fewer transactions; improves bandwidth utilization; requires aligned data - **Asynchronous Copy**: async memory copy (compute capability 8.0+); overlaps with compute; 20-50% speedup; uses copy engines separate from compute - **Prefetching**: load next iteration's data while computing current; hides latency; software pipelining; 15-30% improvement **Latency Hiding:** - **High Occupancy**: more active warps hide memory latency; target 50-100% occupancy; balance register and shared memory usage; 256 threads per block typical - **Instruction-Level Parallelism**: independent operations hide latency; reorder instructions; multiple accumulators; 20-40% improvement - **Warp Scheduling**: GPU schedules ready warps while others wait for memory; sufficient warps (8-16 per SM) ensure full utilization - **Memory-Compute Overlap**: structure kernels to overlap memory access with computation; double buffering; asynchronous operations **Unified Memory:** - **Automatic Migration**: CUDA Unified Memory migrates pages between CPU and GPU; convenient but slower than explicit management; 2-5× overhead vs explicit - **Prefetching**: cudaMemPrefetchAsync() prefetches to GPU; reduces page faults; 50-80% of explicit performance; good for prototyping - **Access Counters**: track which processor accesses data; optimizes placement; reduces migration overhead; improves performance by 30-60% - **When to Use**: rapid prototyping, irregular access patterns, CPU-GPU collaboration; production code prefers explicit management for performance **Memory Bandwidth Bottlenecks:** - **Identification**: Nsight Compute shows memory throughput; <50% of peak indicates memory bound; optimize memory access patterns first - **Arithmetic Intensity**: FLOPs per byte; low intensity (<10) is memory bound; high intensity (>50) is compute bound; tiling increases intensity - **Roofline Model**: plots performance vs arithmetic intensity; shows whether memory or compute limited; guides optimization strategy - **Bandwidth Saturation**: achieved bandwidth / peak bandwidth; target 80-100%; below 50% indicates access pattern problems **Advanced Techniques:** - **Shared Memory Atomics**: faster than global atomics; 10-100× speedup; use for reductions within block; warp-level primitives even faster - **Warp Shuffle**: exchange data between threads in warp; no shared memory needed; 2-5× faster than shared memory; __shfl_sync(), __shfl_down_sync() - **Cooperative Groups**: flexible synchronization; grid-wide sync; warp-level operations; more expressive than __syncthreads() - **Multi-Level Tiling**: tile at multiple levels (L2, shared memory, registers); maximizes reuse at each level; 10-30× speedup for complex algorithms **Profiling and Tuning:** - **Nsight Compute Metrics**: Memory Throughput, L1/L2 Hit Rate, Global Load/Store Efficiency, Shared Memory Bank Conflicts; guide optimization - **Memory Replay**: indicates uncoalesced access; high replay (>1.5) means poor coalescing; restructure data layout - **Occupancy vs Performance**: higher occupancy doesn't always mean better performance; balance with resource usage; profile to find optimal - **Iterative Optimization**: optimize one aspect at a time; measure impact; memory coalescing first, then shared memory, then registers **Common Patterns:** - **Matrix Multiplication**: shared memory tiling; 80-95% of peak; 10-20 TFLOPS on A100; load tiles into shared memory, compute, repeat - **Reduction**: warp primitives + shared memory; 60-80% of peak bandwidth; 500-1000 GB/s; minimize global memory accesses - **Stencil**: shared memory halo; load neighbors into shared memory; 70-90% of peak; 1-2 TB/s; reduces redundant global loads - **Histogram**: shared memory atomics + global atomics; 40-60% of peak; 500-800 GB/s; balance between shared and global atomics **Best Practices:** - **Profile First**: identify bottleneck before optimizing; memory or compute bound; use Nsight Compute - **Coalesce Always**: ensure coalesced access; SoA layout; aligned data; 5-10× improvement - **Use Shared Memory**: for data reused across threads; 100× faster than global; tile algorithms - **Balance Resources**: registers, shared memory, occupancy; find optimal trade-off; profile-guided tuning - **Measure Impact**: verify each optimization improves performance; some optimizations hurt; iterate based on data GPU Memory Hierarchy Optimization is **the art of data orchestration across multiple storage levels** — by understanding the 1000× performance difference between registers and global memory and applying techniques like shared memory tiling, memory coalescing, and register blocking, developers achieve 5-20× performance improvements and 80-100% of theoretical bandwidth, making memory hierarchy optimization the most critical skill for GPU programming where the vast majority of kernels are memory-bound and proper data placement determines whether applications achieve 5% or 80% of peak performance.