/gpu-memory-model

name: gpu-memory-model
description: GPU memory model skill for SIMT execution and memory hierarchy. Use when analyzing warp divergence, memory coalescing, shared memory bank conflicts, cache behavior, atomics, or occupancy tradeoffs. Activates on queries about SIMT, warp coalescing, bank conflicts, wavefront, GPU occupancy, or memory-bound kernels.

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