3.1.2–3.1.3SIMT Deadlock, Stackless SIMT & Warp Scheduling

3.1.2 SIMT Deadlock and Stackless SIMT Architectures

The SIMT stack is elegant — and it has a failure mode serious enough that NVIDIA redesigned divergence handling for Volta. The stack’s rule is “finish the TOS path before anything else.” Combine that rule with a spin lock and you get a program that is correct on any MIMD machine but hangs on a stack-based GPU. ElTantawy and Aamodt named it SIMT deadlockSIMT deadlockStack-induced deadlock: a lock-winner waits at the reconvergence point while spinning losers hold the TOS, so the lock is never released.Glossary →:

A: *mutex = 0;
B: while (!atomicCAS(mutex, 0, 1));   // spin until we take the lock
C: // critical section
   atomicExch(mutex, 0);              // release

atomicCAS reads mutex, compares with 0, and — atomically, serialized across the warp’s threads — swaps in 1 if it matched, returning the old value. So exactly one thread sees 0 (wins); the rest see 1 (spin). Watch the stack paint itself into a corner:

The winner waits at the reconvergence point, below the TOS; the spinners own the TOS and can never succeed until the winner releases — which it can’t, because it isn’t in the TOS active mask. Circular dependence, purely an artifact of the divergence mechanism.

Convergence barriers: divergence handling without a stack

NVIDIA’s fix (disclosed in a patent application and consistent with Volta’s Independent Thread SchedulingIndependent Thread SchedulingNVIDIA's (Volta) name for stackless divergence management that lets diverged paths interleave and locks make progress.Glossary →; equivalent in spirit to academic Multi-Path IPDOM) is to replace the stack with per-warp convergence barriersconvergence barrierVolta-style stackless divergence handling: barrier participation mask + per-thread state; ADD arms, WAIT blocks, scheduler interleaves warp splits freely.Glossary →. The per-warp state (Fig 3.6) is a small set of registers:

• Barrier Participation Mask — which threads of the warp take part in a given convergence barrier (several may exist at once, one per nesting level);
• Barrier State — which participants have already arrived;
• Thread State (per thread) — ready / blocked at a barrier / yielded (the escape hatch that lets other threads make progress in would-be-deadlock situations);
• Thread rPC (per inactive thread) — its next instruction;
• Thread Active — one bit per thread.

Two instructions drive it (Figs 3.7/3.8): ADD arms a barrier — every thread active at that point sets its bit in the named barrier’s participation mask; WAIT stops an arriving warp splitwarp splitSubset of a warp's threads sharing a PC after divergence; the schedulable unit in stackless/multi-path designs.Glossary → at the barrier — its threads turn blocked and join the barrier state. When all participants have executed the WAIT, the scheduler flips them ready together: reconvergence. (A YIELD instruction and indirect-branch support round out the mechanism.)

The crucial difference from the stack: between ADD and WAIT, the scheduler may freely switch among warp splits. Divergence still splits the warp into subsets with a common PC — but no entry “owns” the warp anymore. Compare the book’s two timing figures on the same 8-thread warp:

if (threadIdx.x < 4) { A; B; }
else                 { X; Y; }
Z;   // Volta: preceded by __syncwarp()

Interleaving doesn’t raise SIMD efficiency — each issue still uses part of the datapath — but it lets one split’s memory stalls be covered by the other split’s work, and, decisively, it makes the spin lock finish:

3.1.3 Warp Scheduling

A core hosts many warps — in what order should they issue? (For now, assume each warp issues one instruction and can’t issue again until it completes; §3.2 relaxes this.)

With an “ideal” fixed-latency memory, you’d support just enough warps to hide that latency and schedule them round-robin (RR): every warp gets equal time; if warps × issue-time exceeds the memory latency, the execution units never idle. But warps aren’t free: each thread needs its own registers (that’s what makes zero-cost warp switching possible), so more warps = more register file = fewer cores per chip. The right number is a throughput-per-area balance, not “as many as possible.”

Real memory latency depends on locality — and locality is exactly where scheduling policy bites:

The deeper question — how scheduling interacts with the memory system — returns in force once Chapter 4 builds one, and in §5.1’s research survey.