3.1–3.1.1One-Loop Approximation & SIMT Execution Masking

This chapter and the next split GPU architecture down its natural seam: compute cores here, the memory system in Chapter 4. GPUs run tens of thousands of threads concurrently — per-thread on-chip storage is tiny, yet caches still work (adjacent pixels, adjacent data), and everything about the core is shaped by sustaining huge off-chip bandwidth behind programmability.

The chapter’s roadmap is a single picture. A SIMT core is a SIMT front-end feeding a SIMD datapath, organized as three coupled scheduling loops — instruction fetch, instruction issue, and register access. The book builds understanding as three increasingly accurate “approximation loops,” and this site follows suit: this page is the one-loop view; §3.2–3.3 add the other two.

3.1 One-Loop Approximation

Pretend the core has a single scheduler (roughly the mental model you’d form from the CUDA manual alone). Each cycle it selects one warpwarpNVIDIA's group of 32 scalar threads executed in lockstep on SIMD hardware; the book's generic unit of SIMT execution.Glossary →; the warp’s single program counter finds the next instruction; the instruction is decoded; source operands are fetched from the register file — and, in parallel, the SIMT execution mask for the instruction is determined (how, is §3.1.1’s story). With mask and operands ready, execution proceeds SIMD-style: each thread runs on the function-unit lane it owns, if its mask bit is set.

Function units are heterogeneous, as in CPUs: NVIDIA cores contain special function units (SFUs — transcendentals), load/store units, floating-point units, integer units, and (from Volta) Tensor CoresTensor CoreVolta-generation unit specialized for machine-learning matrix operations.Glossary →. Nominally every unit has one lane per warp thread; several real GPUs instead run a warp across a narrower, faster-clocked unit over several cycles — better performance per area, at an energy cost.

3.1.1 SIMT Execution Masking

The SIMT model’s promise: each thread behaves functionally as if it executed independently (performance is another matter). Pure predicationpredicationExecuting an instruction conditionally under a predicate register (@%p1 in PTX, @P0 in SASS) instead of branching.Glossary → could deliver that promise — but real GPUs combine predication with a stack of predicate masks: the SIMT stackSIMT stackPer-warp stack of (reconvergence PC, next PC, active mask) entries that serializes divergent paths and reconverges them; TOS decides what issues next.Glossary →. The stack elegantly handles the two hard cases:

1. Nested control flow — a branch inside a branch, where one path’s mask must be carved out of another’s;
2. Skipped computation — when no thread in the warp needs a path, the warp skips it entirely (a big win in complex control flow, and awkward to get from plain predication).

(The book describes a slightly simplified, hardware-managed stack from the academic literature; real GPUs manage theirs partly with special instructions — you saw AMD’s software version with exec-mask arithmetic in §2.2.2.)

The worked example

Two branches nested in a do-while loop, four threads per warp. CUDA first, then the PTX it compiles to (labels A–G mark the basic blocks):

do {
  t1 = tid*N;          // A
  t2 = t1 + i;
  t3 = data1[t2];
  t4 = 0;
  if( t3 != t4 ) {
      t5 = data2[t2];  // B
      if( t5 != t4 ) {
          x += 1;      // C
      } else {
          y += 2;      // D
      }
  } else {
      z += 3;          // F
  }
  i++;                 // G
} while( i < N );

A:  mul.lo.u32  t1, tid, N;
    add.u32     t2, t1, i;
    ld.global.u32  t3, [t2];
    mov.u32     t4, 0;
    setp.eq.u32 p1, t3, t4;
@p1 bra F;
B:  ld.global.u32  t5, [t2];
    setp.eq.u32 p2, t5, t4;
@p2 bra D;
C:  add.u32     x, x, 1;
    bra E;
D:  add.u32     y, y, 2;
E:  bra G;
F:  add.u32     z, z, 3;
G:  add.u32     i, i, 1;
    setp.le.u32 p3, i, N;
@p3 bra A;

How can a SIMD datapath — one instruction per cycle for the whole warp — let these four threads take three different routes? By serializing the divergent paths and reconverging them. Step through the entire dance: the control-flow graph, the stack, and the lanes move together.

Three things to internalize from the animation:

• Divergence is a stack edit. At a divergent branch the current entry’s Next PC is rewritten to the branch’s reconvergence pointreconvergence point / immediate post-dominatorEarliest program point where diverged threads are compile-time guaranteed to rejoin; where stack entries pop.Glossary →, and one entry is pushed per outgoing path. The active maskactive maskBit per thread of a warp saying who executes the current instruction (e.g. 1110 = threads 1–3 on, thread 4 masked).Glossary → in each entry says exactly who participates.
• Reconvergence is a pop. A path executes until its PC reaches its entry’s RPC — then the entry pops and the next path runs. The reconvergence point used is the branch’s immediate post-dominator: the earliest spot where diverged threads are compile-time guaranteed to be able to run in lockstep again. (Runtime mechanisms can sometimes do even better — §3.4.2.)
• Divergence costs throughput. While block C ran, three of four lanes were hollow — 25% SIMD efficiencySIMD efficiencyFraction of SIMD lanes doing useful work; divergence's main casualty.Glossary →. Branch divergencebranch divergenceThreads of one warp taking different targets at a data-dependent branch, forcing serialized execution of each path.Glossary → never breaks correctness; it quietly eats performance, which is why §3.4’s research menagerie exists.

One last subtlety the book flags: push order matters for area. Pushing the more-populated path first means the path executed next always has at most half the threads — bounding worst-case stack depth to log₂(warp size) entries instead of warp-size entries (AMD documents exactly this). Fig 3.4’s parts (c) and (d) deliberately show both orders.