Warp Scheduling & SIMT — How the SM Actually Issues Instructions

Topics We'll Cover

The micro-architecture of warp issue: how the SM picks one warp out of dozens each cycle, why the picks decide your kernel's runtime, and the mental models that turn occupancy and divergence from mysteries into arithmetic.

• The Warp — 32 Threads, One Program Counter (Mostly)
• SM Partitions and Schedulers
• The Issue Rule
• Instruction Latencies (Approximate)
• Occupancy & ILP — Two Ways to Hide Latency
• Divergence & Re-convergence
• Predication vs Branching
• Independent Thread Scheduling (Volta+)
• Occupancy Maths — Worked Out
• Common Stall Patterns and Cures
• Interactive: Issue-Rate Predictor

The Warp — 32 Threads, One Program Counter (Mostly)

A warp is the unit of scheduling on an NVIDIA SM: 32 threads grouped together, historically sharing a single program counter and executing the same instruction in lockstep. This is the “Single-Instruction, Multiple-Thread” (SIMT) model — the SM fetches one instruction and broadcasts it to all 32 lanes, each holding a different thread's register state.

Volta (2017) changed the deal: it gave each thread its own PC and call stack so that re-convergence after divergence works automatically and starvation-free. But warps still issue at warp granularity for performance; divergent warps are slower, not broken.

SIMT vs SIMD — the precise difference

• SIMD (CPU AVX-512, ARM SVE): one instruction operates on a vector register; lanes have no per-lane PC, no per-lane control flow.
• SIMT: each thread keeps its own register state, and Volta+ each its own PC. The compiler emits scalar-looking code; the hardware bundles 32 threads at issue time. Predication and reconvergence handle divergence transparently.
• The practical consequence: you write a per-thread C++ kernel, but you must think in warps. A 32-thread warp that touches 32 different cache lines costs 32 transactions; a coalesced one costs 1.

SM Partitions and Schedulers

From Volta onwards an SM is split into 4 partitions (also called “sub-cores” or “processing blocks”). Each partition is essentially its own little SIMT machine, with one warp scheduler that handles a private pool of warps and feeds private functional units.

One Ampere SM partition, drawn out

Issue rate, drawn at SM and GPU scope

• Per partition: 1 instruction / cycle (some Ada/Ampere consumer SKUs effectively 2 FP32 ops/cycle thanks to dual-issue FP32, but only one scheduler slot).
• Per SM: 4 instructions / cycle.
• Per H100: 4 IPC × 132 SMs = 528 issued instructions per cycle. At 1.83 GHz that is ~0.97 trillion warp-instructions/s, of which any one HMMA represents many tensor-core math ops.
• The SM is the smallest unit of scheduling work. CUDA blocks are placed onto whole SMs; warps within a block are spread roughly evenly across the four partitions when the block is launched.

The Issue Rule

Each cycle every scheduler looks at its pool of resident warps (up to 16 per partition, so up to 64 per SM at full Volta+ occupancy) and picks one that is ready. “Ready” is a conjunction of conditions; every condition is a hardware scoreboard or pipeline status bit.

If more than one warp is ready, the scheduler picks one. The exact policy is undocumented and varies by architecture; it has been described as a loose round-robin biased by age and by recently-issued warps. The key invariant is “greedy then fair”: a warp that just got an instruction issued is not preferred next cycle.

# Captured on a memory-bound GEMM kernel, A100, fp16
ncu --set full --section WarpStateStats ./gemm

Section: Warp State Statistics
 
Stall — Long Scoreboard          27.4 cycles/instr   # LDG miss waiting on L2/HBM
Stall — Wait                     14.2 cycles/instr   # dep on a fixed-latency op (FFMA)
Stall — Short Scoreboard         11.8 cycles/instr   # LDS waiting on shared mem
Stall — Math Pipe Throttle        4.1 cycles/instr   # FFMA pipe saturated
Stall — Drain                     3.6 cycles/instr   # EXIT or BRA flushing
Stall — MIO Throttle              2.9 cycles/instr   # shared-mem bank conflicts
Stall — IMC Miss                  1.3 cycles/instr   # constant-cache miss
Stall — Branch Resolving          0.7 cycles/instr   # indirect / divergent branch
Stall — Tex Throttle              0.2 cycles/instr   # tex pipe full
Stall — Membar                    0.0 cycles/instr   # fence in flight
 
# Selected (eligible) warp count: 1.6 / cycle  — only 1.6 warps were ready,
# so “Issued” instructions / cycle were limited even though 4 schedulers exist.

Instruction Latencies (Approximate)

Latency in this context means “how many cycles after issue can a dependent instruction issue?”. Throughput is “how many of the same instruction can issue per cycle?”. The SM hides latency by issuing other warps' instructions in the gap.

The latency-hiding rule of thumb

To hide an op of latency L cycles when its throughput is 1/cycle, you need L independent in-flight instances. With one warp giving one instance per cycle (back-to-back), that is L resident warps with at least one ready instruction each, or fewer warps with ILP ≥ L.

• FFMA, L=4: 4 ready warps per partition trivially saturate the FP32 pipe.
• LDS, L=30: ~30 in-flight loads per partition needed. At 1 warp per LDS that is 30 warps — nearly the full 16-warp pool plus some ILP.
• LDG to HBM, L=600: hopeless to hide with warp count alone. Solve it with prefetch (cp.async), wider tiles, and software pipelining of compute over loads.

Occupancy & ILP — Two Ways to Hide Latency

The scheduler keeps the SM busy by always finding something ready to issue. There are two complementary ways to give it work:

Both tactics are valid and modern CUTLASS / cuBLASLt kernels lean heavily on ILP rather than chasing 100% occupancy. Volta-and-later kernels often deliver near-peak FLOPS at 25–50% occupancy because they pack four independent FFMAs and two prefetches into each warp's loop body.

Two implementations of the same SAXPY-like kernel

__global__ void saxpy_high_occ(float a, const float* x, float* y, int n)
{
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) y[i] = a * x[i] + y[i];
}
// Launch: 256 threads/block, 32 regs/thread → 8 blocks/SM = 64 warps/SM = 100% occ.
// One LDG, one FFMA, one STG per thread. Per-warp ILP = 1.

__global__ void saxpy_high_ilp(float a, const float4* x4, float4* y4, int n4)
{
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i >= n4) return;
    // 128-bit load: one LDG.E.128 brings 16 bytes/thread → 512 B/warp aligned
    float4 xv = x4[i];
    float4 yv = y4[i];
    yv.x = a * xv.x + yv.x;
    yv.y = a * xv.y + yv.y;
    yv.z = a * xv.z + yv.z;
    yv.w = a * xv.w + yv.w;   // 4 independent FFMAs per warp per loop iter
    y4[i] = yv;
}
// Launch: 128 threads/block, 64 regs/thread → 4 blocks/SM = 16 warps/SM = 25% occ.
// Per-warp ILP = 4 FFMAs; one wide LDG hides — achieves 95+% of memory peak.

Divergence & Re-convergence

When threads in a warp take different paths through an if/else, the warp diverges. The hardware tracks per-lane activity with an active mask; the warp issues the if-branch with one mask, then the else-branch with the inverse, both serialised. After the branches re-meet at the immediate post-dominator, they reconverge and the active mask returns to all-ones.

Predication vs Branching

The compiler has two ways to implement an if in a warp:

• Predication. Emit both sides of the branch unconditionally, but tag each instruction with a predicate register. Lanes whose predicate is false simply discard their result — the issue slot is still spent.
• Branching. Emit a real BRA. If the lanes agree (uniform branch), only one side runs. If they disagree, the warp diverges and both sides run sequentially with masks — same cost as predication, plus branch overhead.

The ptxas heuristic: short branches (1–3 instructions per side) get predicated; long branches get a real branch because predication wastes too many issue slots and can hurt instruction-cache density. You can force one or the other with #pragma unroll, by using ?: for short branches, and by structuring loops so the compiler sees them as uniform.

Worked SASS — ptxas chose predication for a tiny if/else

__device__ unsigned choose(unsigned a, unsigned b)
{
    return (a > b) ? a + b : a + 0x10u;
}

// R0 = a, R1 = b
ISETP.GT.AND.U32 P0, PT, R0, R1, PT      // P0 = (a > b)
@P0  IADD3 R2, R0, R1, RZ                // if-true:  R2 = a + b
@!P0 IADD3 R2, R0, 0x10, RZ              // if-false: R2 = a + 0x10
// no BRA emitted — both IADD3s issue, predicates filter their commit

Both IADD3 instructions issue in every lane. Lanes where P0 is false on the first instruction discard the result; the same lanes where P0 is true on the second instruction discard that result. The cost is two issue slots but no divergence overhead. For three-instruction bodies it is exactly break-even with a real branch when the condition is fully divergent.

Independent Thread Scheduling (Volta+)

Lockstep warp execution made nested critical sections impossible to write safely. A thread holding a lock could not release it while siblings span on a sibling lock; both lanes would be pinned to the same PC, both holding their respective locks forever. Several decades of multi-threaded textbook code simply did not port to pre-Volta GPUs.

Volta's promise: per-thread forward progress

• Each thread owns its own PC and call stack.
• The scheduler is allowed to interleave divergent halves of a warp at instruction granularity.
• The hardware guarantees that every thread eventually makes forward progress (no starvation), so spin-wait + release patterns terminate.
• This unlocks idiomatic concurrent code:

__device__ int mutex = 0;

__device__ void acquire() {
    while (atomicCAS(&mutex, 0, 1) != 0) { /* spin */ }
}
__device__ void release() { atomicExch(&mutex, 0); }

__global__ void k(float* shared_acc, float v) {
    acquire();
    *shared_acc += v;       // critical section
    release();
    // On Pascal: lanes that lost the CAS spin forever because the winning lane
    // is at the same PC and cannot make progress to the release.
    // On Volta+ ITS: the winner can advance to release, the loser then makes
    // progress on its own retry, and the warp converges naturally.
}

The cost: explicit warp synchronisation now required

Code that relied on lockstep, especially warp-level intrinsics, broke. Volta+ introduced the _sync family with a 32-bit mask argument: the mask names participating lanes and forces them to converge before the op.

// Pascal era — relied on implicit lockstep
float v = __shfl(x, lane);                       // deprecated, removed in CUDA 9+

// Volta+ era — explicit mask, explicit convergence
float v = __shfl_sync(0xffffffff, x, lane);     // all 32 lanes participate
unsigned mask = __activemask();                   // query who is here right now
int b      = __ballot_sync(mask, predicate);
int sum    = __reduce_add_sync(mask, x);          // Volta+ HW reduction

Occupancy Maths — Worked Out

Occupancy = (resident warps per SM) / (max warps per SM). On Volta and later the cap is 64 warps/SM (= 2048 threads). The actual resident count is the minimum of three independent limits.

Worked example A — register-bound

• Block: 256 threads, 64 regs/thread.
• Per-block regs: 256 × 64 = 16 384 regs = 16 K → 4 blocks/SM fit in 64 K registers.
• 4 blocks × 256 threads = 1024 threads = 32 warps/SM = 50% occupancy.
• The warp limit (64) and block limit (32) are not binding here; registers are.

Worked example B — warp-cap-bound

• Same block size, but reduce regs/thread to 32.
• Per-block: 256 × 32 = 8 192 regs → up to 8 blocks/SM by registers alone.
• But 8 × 256 = 2 048 threads = 64 warps — at the warp cap.
• So 8 blocks × 256 = 64 warps = 100% occupancy.

Worked example C — SMEM-bound

• Block: 128 threads, 48 regs/thread, 48 KB SMEM/block.
• Registers: 128 × 48 = 6 144 regs → 64K/6144 ≈ 10 blocks; not binding.
• Warps: 128/32 = 4 warps/block → warp cap allows 64/4 = 16 blocks; not binding.
• SMEM: 100 KB / 48 KB = 2 blocks/SM; SMEM is binding.
• 2 blocks × 4 warps = 8 warps/SM = 12.5% occupancy. Probably bandwidth-bound and will need either smaller tiles or async-copy double-buffering to hit peak.

int blocks = 0, threads_per_block = 256;
size_t shared_bytes = 48 * 1024;
cudaOccupancyMaxActiveBlocksPerMultiprocessor(
    &blocks,                       // out
    my_kernel,                     // kernel function pointer
    threads_per_block,             // per-block thread count
    shared_bytes);                 // dynamic SMEM

// Or pick a configuration that maximises occupancy:
int minGrid = 0, optBlock = 0;
cudaOccupancyMaxPotentialBlockSize(&minGrid, &optBlock,
    my_kernel, /*dyn smem fn*/0, /*max block*/1024);

// And to query the actual achieved occupancy at run time:
int active = 0;
cudaOccupancyMaxActiveBlocksPerMultiprocessorWithFlags(
    &active, my_kernel, threads_per_block, shared_bytes,
    cudaOccupancyDefault);

Common Stall Patterns and Cures

Nsight Compute names ten or so warp-stall reasons. Most kernels show one or two dominant ones; once you know which, the cure is usually mechanical.

Interactive: Issue-Rate Predictor

Pick a GPU class and your block size / register pressure / shared memory. The predictor computes the binding limit and estimates achievable IPC plus a likely dominant stall.