GPU memory model

Contract

Field	Bound contract
Trigger	Kernel behavior needs explaining or predicting through the SIMT execution model and the memory hierarchy: divergence costs, coalescing rules, bank conflicts, cache levels, atomic contention, or occupancy versus latency hiding.
Authority	Read-only. The skill emits explanations and code patterns; it writes no files. Nothing to roll back. No remote mutation.
Side effect	Explanations, access-pattern rewrites, and metric names to confirm each hypothesis.
Done	The observed behavior is derived from the model, and a profiler metric that can confirm or refute it is named.

Inputs

1. Kernel or pattern (required): the code whose behavior needs explaining.
2. Vendor and architecture family (required): NVIDIA (warp) or AMD (wavefront), CDNA or RDNA for AMD.
3. Symptoms (optional): profiler readings that motivated the question.

Procedure

1. Ground the execution model. One instruction stream drives a warp or wavefront; each thread keeps its own registers and ID. The SM or CU holds the schedulers, the register file, shared memory or LDS, and usually the L1; L2 is device-wide. Done when: the question is located at one level of that hierarchy.
2. Size the scheduling unit per vendor:

Vendor	Unit	Notes
NVIDIA	32 threads (warp)	All current parts
AMD CDNA (MI series)	64 threads (wavefront)	Fixed width
AMD RDNA (consumer)	32 or 64 threads	Compiler picks per kernel

Consequences: reduction trees halve 16, 8, 4, 2, 1 on NVIDIA and 32, 16, 8, 4, 2, 1 on 64-lane AMD; block sizes should be multiples of the unit; occupancy counters report active units per SM. Done when: reductions and block sizes match the recorded unit. 3. Price divergence. When lanes take different branches, the hardware serializes the paths, so cost follows the sum of the taken paths, not the maximum:

// Divergent: the two halves issue one after the other
if (threadIdx.x % 2 == 0) { result = expensive_a(d[i]); }
else                      { result = expensive_b(d[i]); }

Mitigations, in order of preference: branch on data uniform within a warp; predicate both sides and select with ?: when the sides are cheap; split into separate kernels when the paths are heavy; regroup work so outcomes land in shared bins instead of per-thread branches. Done when: every hot divergent branch has a mitigation or a measured justification. 4. Check coalescing. Consecutive threads reading consecutive 4-byte words collapse into one wide transaction (a 128-byte sector run on NVIDIA); stride or misalignment multiplies transactions:

int idx = blockIdx.x * blockDim.x + threadIdx.x;
float good = data[idx];                 // coalesced
float bad  = data[threadIdx.x * stride]; // strided

Array-of-structs reads of one field stride across the struct size; struct-of-arrays restores contiguity:

struct Particle { float x, y, z; };
float aos_bad = particles[i].x;  // stride 3
float soa_good = pos_x[i];       // stride 1

Done when: every hot global access is contiguous across the warp, or the stride is deliberate. 5. Check shared memory bank conflicts. Shared memory and AMD LDS are organized in 32 banks of 4 bytes. Simultaneous accesses to different words in one bank serialize; a broadcast of one address to all lanes does not:

__shared__ float tile[32][32];
float conflict = tile[threadIdx.x][0]; // column 0 maps many lanes to one bank

__shared__ float padded[32][33];       // +1 column breaks the bank alignment
float ok = padded[threadIdx.x][0];

Confirm with Nsight Compute shared-load conflict metrics such as l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_ld.sum. Done when: hot shared reads are conflict-free or the conflict count is measured and accepted. 6. Read cache behavior by level: L1 is per SM and often shares capacity with shared memory; L2 is device-wide with 128-byte lines; the read-only path serves uniform lookups through the texture pipeline. Patterns that help: consecutive threads on consecutive addresses; reuse tiles in shared memory before re-fetching; avoid pointer chasing and heavy scatter, which defeat both caches. Done when: the data reuse plan survives one pass through shared memory. 7. Keep atomics off the critical path. atomicAdd, atomicCAS, and atomicExch serialize updates to the same address across the device; ordering beyond that serialization comes from fences and scope:

__threadfence_block();   // visible within the block
__threadfence();         // visible across the device
__threadfence_system();  // visible to the host; expensive

Reduce hierarchically: warp reduction, block reduction into shared memory, then one atomic per block:

if (threadIdx.x == 0)
    atomicAdd(&global_sum, block_sum);

Done when: contended atomics are reduced to one per block or the contention is measured and justified. 8. Balance occupancy against latency hiding. High occupancy gives the scheduler more warps to hide memory latency and costs register and shared memory headroom per block. Low occupancy with high instruction-level parallelism can still saturate a compute-bound kernel. Decide from measurement, not a fixed occupancy target: run the kernel, then read sm__warps_active.avg.pct_of_peak_sustained_active plus the stall reasons in Nsight Compute, and treat low achieved occupancy as a problem only when stalls show memory latency not being hidden. Done when: the occupancy choice is tied to measured stalls. 9. Derive the answer and its confirming metric:

Symptom	Likely cause	Confirm with
Low DRAM throughput	Uncoalesced access	Memory workload analysis
Shared load stalls	Bank conflicts	Shared conflict counters
Barrier stalls	Missing __syncthreads or divergence at the barrier	synccheck
Atomic bottleneck	Global contention	Warp stall sampling
Low occupancy	Register or shared memory pressure	launch__occupancy_limit_*

Done when: the explanation and its metric are both stated.

Failure and recovery

Failure class	Behavior
Symptom fits two causes	Name both, rank by the metric evidence, and name the measurement that separates them. Do not pick silently.
Vendor behavior unknown	State the model for the recorded vendor and unit width; do not transfer an NVIDIA rule to AMD without noting the unit difference.
Model and measurement disagree	Trust the measurement; revise the explanation and say which assumption failed.

Output

An explanation tied to the model: the execution or memory rule, the rewritten pattern where one applies, and the profiler metric that confirms it. Every claim names its vendor and unit width.