8. Schedulers, stalls, and execution pipes

Each SM has warp schedulers. On a cycle, a scheduler chooses an eligible warp and issues its next instruction. A warp is 32 threads executing the same instruction together in lockstep (SIMT); "issuing an instruction" for a warp means all 32 lanes step forward together. A warp may be active but ineligible because an operand, pipeline, barrier, or memory dependency is not ready — it is "not eligible" precisely when the next instruction it needs to run still has an unmet dependency: an operand not yet computed, a memory value not yet returned, a barrier not yet reached.

Figure 8.1 — The same surface symptoms can describe either a causal defect or an efficient producer/consumer pipeline.

Start with scheduler health

Only after establishing poor issue efficiency should stall reasons dominate the investigation.

Common stall reasons

Names vary slightly by architecture and tool release; use the installed Profiling Guide for exact definitions.

Turning a stall diagnosis into a fix

The table above names a direction — "coalescing," "reduce operations," "predication" — but a direction is not yet a technique. This section fills in the concrete mechanism behind each stall category's productive experiments, so a diagnosis has somewhere specific to go.

Long Scoreboard and MIO Throttle (memory-bound stalls). If a warp is repeatedly stalling because its next instruction needs a value that has not come back from memory yet, the instruction stream is asking for data too late relative to when it is used. Software pipelining, also called double buffering, restructures the loop so the load for tile $i+1$ is issued before the compute on tile $i$ finishes, using a second buffer (a second shared-memory tile or register set) so the two tiles do not collide. The warp then has independent work — compute on the already-resident tile — to issue while the new load is in flight, instead of stalling on every load it makes. Vectorized loads attack the same stall from the instruction side: replacing four scalar float loads with one float4 (or the int4 equivalent) load moves the same bytes in a quarter of the instructions, which reduces address-generation and issue-slot pressure on the memory pipe and, when the access is already contiguous per thread, tends to raise sector efficiency as a side effect (see ch. 9's coalescing discussion for the sector mechanics). On Ampere and newer, cp.async gives a warp a way to issue an asynchronous global-to-shared copy that never round-trips through the register file, freeing registers and letting the copy proceed while the warp does other work; on Hopper and newer, the Tensor Memory Accelerator (TMA) extends the same idea to bulk, whole-tile copies issued by a single thread on behalf of the CTA. These two mechanisms are the hardware underneath the "producer warps issue copies, consumer warps issue MMA" pattern described from the resource-budget side in ch. 10's warp-specialization section; this section is the same pattern described from the instruction-sequencing side — overlap is what turns a per-tile load stall into hidden latency.

Short Scoreboard and Barrier (on-chip dependency stalls). When the dependency is between lanes of the same warp — a reduction, a small exchange, a broadcast — routing the value through shared memory and a __syncthreads() barrier is often more machinery than the problem needs. Warp-level shuffle intrinsics (__shfl_sync and its variants) let the 32 lanes of one warp exchange register values directly, with no shared-memory write, no barrier, and no bank-conflict exposure, because the exchange happens inside the warp's own execution unit rather than through the shared memory system. This only applies within a single warp — cross-warp exchange still needs shared memory (and, if it must be ordered, a barrier), so shuffles are a targeted fix for warp-local dependency stalls, not a general replacement for shared-memory communication.

Math Pipe Throttle. Once a kernel's instruction stream is already minimal for the algorithm, "reduce operations" stops being available as a lever, and only two real levers remain. One is raising arithmetic intensity — doing more MACs per byte already resident, through tiling and reuse — which is the same roofline argument ch. 9 makes from the bandwidth side; the two are connected because a kernel that is compute-throttled on a low-precision pipe while its data is barely reused is often better addressed by fixing reuse than by chasing the compute pipe directly. The other is moving the work to a cheaper or lower-precision pipe: routing matrix-shaped math through a Tensor Core TF32/FP16/BF16/FP8 path instead of scalar FP32 CUDA-core instructions, where the workload's correctness budget tolerates the reduced precision. Neither lever is "delete instructions" — that option is rarely on the table once the algorithm itself is fixed.

Branch Resolving / Divergence. Divergence stalls come from lanes in a warp disagreeing about which path to take, forcing the hardware to execute both paths serially with lanes masked off. Two techniques address this at different levels. Grouping or sorting data so that the lanes assigned to one warp are likely to agree on the branch outcome — for example, sorting work items by the predicate before launch, or assigning a warp's lanes contiguous rather than interleaved elements — removes the disagreement at the data level, before the kernel ever runs. Predication removes it at the instruction level: rewriting a short conditional as arithmetic (min, max, a select, or a ternary) so every lane computes both outcomes and selects, rather than branching. Predication trades wasted arithmetic for guaranteed uniform control flow, so it pays off exactly when the conditional body is cheap enough that computing both sides costs less than the divergence did.

Sleeping (persistent and warp-specialized kernels). Sleeping stalls are not automatically a defect — ch. 10's warp-specialized GEMM is a case where they are the expected behavior of producer warps waiting for the next hand-off. The mechanism that produces this pattern deliberately is the persistent-kernel pattern: instead of launching a grid sized to cover the problem in one shot, the grid is sized to exactly the number of CTAs the device can keep resident, and each CTA loops internally over the problem's tiles — pulling work from a queue or a fixed schedule — until the problem is exhausted. Because the CTAs never exit and relaunch, a warp within one can go idle (sleeping) between hand-offs without paying a relaunch cost, which is what makes low occupancy and high sleeping time compatible with a near-saturated target pipe rather than being evidence of starvation.

Read stalls as a causal chain

A stall table only becomes a diagnosis once it is chained to occupancy and to the pipe the kernel is supposed to be feeding. Two kernels can report nearly the same low occupancy and mean opposite things:

Same headline number, opposite root cause: one kernel is starved because eligible warps are scarce and the ones present are waiting on memory; the other is idle by design because deep pipelining needs fewer, heavier CTAs. The same “low occupancy” headline supports opposite conclusions once scheduler and pipeline evidence is included.

Pipeline misalignment

SM utilization is an umbrella; inspect its breakdown. A kernel can issue many integer address instructions while barely using FP/Tensor units, or saturate Tensor units while overall SM percentage appears moderate.

Inspect executed instruction counts and source-correlated hot instructions before rewriting code around a single utilization percentage.