gpu-memory-model

GPU Memory Model

Purpose

Explain the GPU execution and memory model for agents optimizing kernels: SIMT execution, warp (32) vs wavefront (64) divergence costs, global memory coalescing rules, shared memory bank conflicts, L1/L2 cache behavior, atomic memory ordering, and the occupancy-vs-latency-hiding tradeoff.

When to Use

• Diagnosing why a kernel is memory-bound despite high theoretical bandwidth
• Understanding warp divergence from branching
• Fixing shared memory bank conflicts in tiled algorithms
• Choosing block size for occupancy vs register pressure
• Porting kernels between NVIDIA (warp 32) and AMD (wavefront 64)
• Reasoning about atomic contention in parallel reductions

Workflow

1. SIMT execution model

GPU hardware
├── Device
│   └── SM / CU (Streaming Multiprocessor / Compute Unit)
│       ├── Warp schedulers (NVIDIA) or Wavefront schedulers (AMD)
│       │   └── Warp/Wavefront (32 or 64 threads in lockstep)
│       ├── Register file (partitioned per thread)
│       ├── Shared memory / LDS (per SM)
│       └── L1 cache (often shared with shared memory)
└── L2 cache (device-wide) → DRAM/HBM

SIMT (Single Instruction, Multiple Threads): one instruction stream drives a warp/wavefront; each thread has its own registers and thread ID but executes the same instruction in lockstep.

2. Warp vs wavefront

Vendor	Unit size	Name
NVIDIA	32 threads	Warp
AMD	64 threads	Wavefront

Implications:

• Reduction trees: NVIDIA halves at 16→8→4→2→1; AMD at 32→16→8→4→2→1
• Block sizes: prefer multiples of 32 (NVIDIA) or 64 (AMD)
• Occupancy counters report active warps/wavefronts per SM

3. Warp divergence cost model

When threads in a warp take different branches, the hardware serializes paths:

// Divergent: half warp does A, half does B → 2x instruction issue
if (threadIdx.x % 2 == 0) {
    result = expensive_a(data[idx]);
} else {
    result = expensive_b(data[idx]);
}

// Non-divergent: all threads same path
result = expensive_a(data[idx]);