Warp & SIMT

Definition

A warp is a group of 32 threads (lanes) that the hardware schedules as the atomic unit of execution. SIMT (Single Instruction, Multiple Threads) is the model: all 32 lanes share one program counter and execute the same instruction in lockstep, each on its own registers and data. You program scalar per-thread code; the hardware bundles 32 threads into a warp behind the scenes. Nearly every GPU performance phenomenon — coalescing, bank conflicts, divergence — follows from this 32-wide bundling.

How it works

The transistor payoff: because 32 lanes share one fetch/decode and one PC, control logic is amortized 32× — that’s how you afford thousands of ALUs without paying for thousands of instruction streams.

SIMT vs SIMD: SIMD (AVX-512) exposes the vector width in your code; SIMT lets you write scalar per-thread code and lets lanes (in principle) follow different control paths — at a cost.

Warp divergence. When lanes in a warp take different branches, the warp cannot issue two instructions at once, so the hardware serializes the paths, masking off inactive lanes on each:
$t_{\text{divergent}}\approx t_A+t_B\text{vs.}{t}_{\text{uniform}}\approx \max (t_A,t_B)$
Worst case (32 different paths) loses up to 32× on the divergent region. Divergence is per-warp, not per-block: a branch on blockIdx.x, or any condition uniform across each 32-lane group, costs nothing. Fix data-dependent divergence by partitioning data so a warp handles a uniform chunk.

Predication. For short branches the compiler skips branching entirely: it executes both sides on all lanes and uses a per-lane predicate bit to decide whether each instruction commits. Cheaper than a real branch for small bodies; for large bodies, real divergent control flow wins (predication would run both full bodies everywhere).

Volta+ Independent Thread Scheduling: each thread now has its own PC/stack, so divergent lanes can interleave and make independent progress (fixed some starvation deadlocks) — but divergence is still not free (reduced lane utilization). Never assume implicit intra-warp sync; use __syncwarp() and *_sync intrinsics.

Why it matters

GEMM and the transformer’s dense projections have uniform control flow — every lane does the same FMA on different data — so SIMT is a near-perfect fit and warp divergence is zero. Divergence appears in the memory-bound glue (masking, variable-length sequences, sampling), where it silently caps throughput.

Taught in

• 01-gpu-architecture-and-simt — §5 the SIMT execution model, divergence, predication.

See also

• streaming-multiprocessor — where warps are scheduled
• cuda-execution-model — how threads map into warps
• occupancy — resident warps hide latency
• memory-coalescing — a warp’s 32 addresses → transactions
• gpu-systems-for-llms