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

<svg viewBox="0 0 760 470" xmlns="http://www.w3.org/2000/svg" font-family="-apple-system,Segoe UI,Roboto,sans-serif">
  <rect x="0" y="0" width="760" height="470" fill="#0d1117"/>
  <text x="380" y="28" fill="#e6edf3" font-size="21" font-weight="700" text-anchor="middle">GPU Memory Hierarchy — CUDA Optimization</text>
  <text x="380" y="48" fill="#8b98a5" font-size="12" text-anchor="middle">registers → shared memory → L1/L2 → HBM — each level trades capacity for speed</text>

  <!-- Memory hierarchy pyramid -->
  <rect x="30" y="65" width="350" height="340" rx="6" fill="#080d14" stroke="#233043" stroke-width="1"/>
  <text x="205" y="84" fill="#e6edf3" font-size="10" font-weight="600" text-anchor="middle">Memory Hierarchy (H100)</text>

  <!-- Registers (fastest, smallest) -->
  <rect x="145" y="98" width="120" height="34" rx="4" fill="#14261f" stroke="#34d399" stroke-width="1.2"/>
  <text x="205" y="114" fill="#6ee7b7" font-size="9" font-weight="600" text-anchor="middle">Registers</text>
  <text x="205" y="127" fill="#6b7684" font-size="7" text-anchor="middle">255 per thread · 0 cycles</text>

  <!-- Shared memory -->
  <rect x="115" y="142" width="180" height="34" rx="4" fill="#0b1220" stroke="#60a5fa" stroke-width="1.2"/>
  <text x="205" y="158" fill="#93c5fd" font-size="9" font-weight="600" text-anchor="middle">Shared Memory (SMEM)</text>
  <text x="205" y="171" fill="#6b7684" font-size="7" text-anchor="middle">228 KB/SM · ~20 cycles · 19 TB/s</text>

  <!-- L1 cache -->
  <rect x="95" y="186" width="220" height="34" rx="4" fill="#0b1220" stroke="#a78bfa" stroke-width="1"/>
  <text x="205" y="202" fill="#c4b5fd" font-size="9" font-weight="600" text-anchor="middle">L1 Cache</text>
  <text x="205" y="215" fill="#6b7684" font-size="7" text-anchor="middle">256 KB/SM · ~30 cycles · shared with SMEM</text>

  <!-- L2 cache -->
  <rect x="75" y="230" width="260" height="34" rx="4" fill="#0b1220" stroke="#f59e0b" stroke-width="1"/>
  <text x="205" y="246" fill="#fbbf24" font-size="9" font-weight="600" text-anchor="middle">L2 Cache</text>
  <text x="205" y="259" fill="#6b7684" font-size="7" text-anchor="middle">50 MB · ~200 cycles · 12 TB/s</text>

  <!-- HBM (global memory) -->
  <rect x="55" y="274" width="300" height="40" rx="4" fill="#0b1220" stroke="#f87171" stroke-width="1.2"/>
  <text x="205" y="292" fill="#fca5a5" font-size="9" font-weight="600" text-anchor="middle">HBM3e (Global Memory)</text>
  <text x="205" y="306" fill="#6b7684" font-size="7" text-anchor="middle">80 GB · ~400 cycles · 3.35 TB/s</text>

  <!-- Host/PCIe -->
  <rect x="55" y="324" width="300" height="30" rx="4" fill="#1a0a0a" stroke="#475569" stroke-width="0.8"/>
  <text x="205" y="343" fill="#6b7684" font-size="8" text-anchor="middle">CPU Host (PCIe 5.0: 64 GB/s · NVLink: 900 GB/s)</text>

  <!-- Speed/capacity arrows -->
  <text x="50" y="190" fill="#34d399" font-size="7" text-anchor="middle">faster</text>
  <text x="50" y="205" fill="#34d399" font-size="9" text-anchor="middle">↑</text>
  <text x="50" y="265" fill="#f87171" font-size="9" text-anchor="middle">↓</text>
  <text x="50" y="280" fill="#f87171" font-size="7" text-anchor="middle">bigger</text>

  <text x="205" y="375" fill="#8b98a5" font-size="8" text-anchor="middle">key insight: keep data in SMEM/registers</text>
  <text x="205" y="390" fill="#8b98a5" font-size="8" text-anchor="middle">to avoid slow HBM round-trips</text>

  <!-- Right: Optimization techniques -->
  <rect x="400" y="65" width="330" height="340" rx="6" fill="#080d14" stroke="#233043" stroke-width="1"/>
  <text x="565" y="84" fill="#e6edf3" font-size="10" font-weight="600" text-anchor="middle">Optimization Strategies</text>

  <text x="420" y="110" fill="#34d399" font-size="9" font-weight="600">1. Maximize shared memory reuse</text>
  <text x="420" y="126" fill="#8b98a5" font-size="8">tile GEMM into SMEM blocks (128×128 tiles)</text>
  <text x="420" y="140" fill="#6b7684" font-size="7.5">FlashAttention: fuse Q×K×V in SRAM, never write attn matrix to HBM</text>

  <text x="420" y="164" fill="#60a5fa" font-size="9" font-weight="600">2. Coalesced global memory access</text>
  <text x="420" y="180" fill="#8b98a5" font-size="8">adjacent threads access adjacent addresses</text>
  <text x="420" y="194" fill="#6b7684" font-size="7.5">128-byte transaction: 32 threads × 4 bytes = 1 cache line</text>

  <text x="420" y="218" fill="#a78bfa" font-size="9" font-weight="600">3. Occupancy tuning</text>
  <text x="420" y="234" fill="#8b98a5" font-size="8">balance registers/SMEM per block vs active warps</text>
  <text x="420" y="248" fill="#6b7684" font-size="7.5">more warps = hide latency, fewer = more resources/thread</text>

  <text x="420" y="272" fill="#f59e0b" font-size="9" font-weight="600">4. Async memory copies (cp.async)</text>
  <text x="420" y="288" fill="#8b98a5" font-size="8">overlap HBM→SMEM with compute</text>
  <text x="420" y="302" fill="#6b7684" font-size="7.5">software pipelining: load next tile while computing current</text>

  <text x="420" y="326" fill="#f87171" font-size="9" font-weight="600">5. Avoid bank conflicts</text>
  <text x="420" y="342" fill="#8b98a5" font-size="8">SMEM has 32 banks — pad arrays to avoid collisions</text>
  <text x="420" y="356" fill="#6b7684" font-size="7.5">conflict = serialization = wasted cycles</text>

  <text x="420" y="380" fill="#38bdf8" font-size="9" font-weight="600">6. L2 persistence (H100+)</text>
  <text x="420" y="396" fill="#8b98a5" font-size="8">pin hot data in L2 across kernel launches</text>

  <!-- Footer -->
  <text x="380" y="432" fill="#8b98a5" font-size="8.5" text-anchor="middle">FlashAttention's entire insight: keep everything in SRAM (19 TB/s) instead of HBM (3.35 TB/s) → 5× faster attention</text>
  <text x="380" y="452" fill="#6b7684" font-size="11" text-anchor="middle">GPU programming is memory programming — the fastest kernel is the one that never touches global memory.</text>
</svg>

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.