mohitmishra786/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]);

Mitigations:

• Predication: compute both, select with ?: (trade compute for uniformity)
• Separate kernels for different code paths
• Branch only on block-level data (uniform within warp)
• Loop over bins instead of if (data[i] < threshold) per thread with scattered outcomes

Divergence cost ≈ sum of paths taken (not max).

4. Global memory coalescing

NVIDIA coalescing rule (simplified): threads in a warp accessing consecutive 4-byte words → single 128-byte transaction.

// Coalesced: consecutive threads → consecutive addresses
int idx = blockIdx.x * blockDim.x + threadIdx.x;
float val = data[idx];

// Uncoalesced: stride access
float val = data[threadIdx.x * stride];  // stride > 1

// Partially coalesced: misaligned start
float val = data[base + threadIdx.x * 3];

AoS vs SoA impact:

// AoS — poor coalescing when reading one field
struct Particle { float x, y, z; };
float x = particles[i].x;  // threads read with stride 3

// SoA — coalesced
float x = pos_x[i];

5. Shared memory bank conflicts

Shared memory is divided into 32 banks (4-byte words). Simultaneous accesses to different addresses in the same bank serialize.

__shared__ float tile[32][32];

// Bank conflict: all threads access tile[threadIdx.x][0]
// 32 threads, 32 banks, but column 0 → same bank per row offset
float val = tile[threadIdx.x][0];

// Fix: pad columns to break bank alignment
__shared__ float tile[32][33];  // +1 padding

Detection: Nsight Compute l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_ld.sum or NCU shared load conflict metrics.

6. L1/L2 cache behavior

| Level | Scope | Notes | |-------|-------|-------| | L1 | Per-SM | Often unified with shared mem; configurable split | | L2 | Device-wide | Cache lines typically 128 bytes | | Texture/L1 readonly | Per-SM | Cached read-only path for uniform access |

Cache-friendly patterns:

• Spatial locality: consecutive threads access consecutive memory
• Temporal locality: reuse data in shared memory before re-fetching global
• Avoid random scatter: atomic updates and pointer chasing defeat caches

// Cache-friendly tile load
for (int t = 0; t < num_tiles; t++) {
    __shared__ float smem[TILE][TILE];
    smem[ty][tx] = global[row * N + t * TILE + tx];
    __syncthreads();
    // compute from smem — L1/L2 only hit on first load per tile
}

7. Atomics and memory ordering

GPU atomics (atomicAdd, atomicCAS, atomicExch) provide sequential consistency among threads targeting the same address, but high contention serializes execution.

// Bad: all threads atomic to one counter
atomicAdd(&global_sum, local_val);

// Better: per-block reduction, one atomic per block
__shared__ float block_sum;
// ... warp reduce to block_sum ...
if (threadIdx.x == 0)
    atomicAdd(&global_sum, block_sum);

HIP/CUDA memory fences:

__threadfence_block();  // visible to threads in same block
__threadfence();        // visible to all threads on device
__threadfence_system(); // visible to host (expensive)

8. Occupancy vs latency hiding

Occupancy tradeoff
├── High occupancy → more warps to hide memory latency
│   └── Costs: fewer registers/SM, less shared mem per block
└── Low occupancy + high ILP → enough independent instructions per warp
    └── Works for compute-bound kernels with deep pipelines

Decision tree:

Memory-bound kernel?
├── Yes → maximize active warps (occupancy), coalesce, tile with shared mem
└── No (compute-bound) → may lower occupancy if registers enable more ILP

# Measure achieved occupancy
ncu --metrics sm__warps_active.avg.pct_of_peak_sustained_active ./app

Rule of thumb: memory-bound kernels need occupancy ≥ 50%; compute-bound may run well at 25% with sufficient instruction-level parallelism.

9. Quick reference table

| Symptom | Likely cause | Check | |---------|--------------|-------| | Low DRAM throughput | Uncoalesced access | NCU memory workload analysis | | Shared load stalls | Bank conflicts | Pad arrays; change access pattern | | High stall: barrier | Missing __syncthreads or divergence at barrier | synccheck | | Atomic bottleneck | Too many contended atomics | Hierarchical reduction | | Low occupancy | Register/shared mem pressure | launch__occupancy_limit_* metrics |

Common Problems

| Symptom | Cause | Fix | |---------|-------|-----| | 10x slower after SoA→AoS change | Strided coalescing broken | Keep hot fields in SoA layout | | Tiled matmul slower than naive | Bank conflicts in shared tile | Pad shared array columns | | Identical code, different perf NVIDIA vs AMD | Warp 32 vs wavefront 64 | Retune block size and reductions | | Atomics correct but slow | Global contention | Block-level reduce first | | High L2 hit but still slow | L2 bandwidth saturated | Reduce total bytes moved |

Related Skills

• skills/gpu/cuda — practical kernel patterns using this model
• skills/gpu/cuda-profiling — metrics to validate coalescing and occupancy
• skills/gpu/hip-rocm — AMD wavefront-specific tuning
• skills/low-level-programming/cpu-cache-opt — CPU cache concepts (analogous)
• skills/low-level-programming/simd-intrinsics — vector width on CPU side
• skills/profilers/hardware-counters — general cache miss measurement