Lab 01 — SIMT Simulator: Warps, Masks, and Divergence

Phase: 01 — GPU Architecture | Difficulty: ⭐⭐⭐☆☆ | Time: 3–4 hours Language: Python (stdlib only) | Hardware: none

Concept primer: ../WARMUP.md Ch. 3–4 (SIMT, active mask, divergence), ../HITCHHIKERS-GUIDE.md §2.

Run

python solution.py        # runs the simulator + divergence benchmark + asserts

0. The mission

Build (and then extend) a tiny SIMT machine: one program counter per warp, 32 lanes, an active mask — and use it to measure what divergence costs instead of taking the textbook's word for it.

The two questions you must be able to answer at the end:

  1. Why does if (tid % 2) cost 2× but if (tid < 32) cost nothing (per warp)?
  2. Why is the worst case exactly 32× — and what real code patterns approach it?

1. The ISA

The simulator executes a list of instructions over 32 lanes. Each lane has its own register file; the warp has one PC and a mask stack:

InstrSemantics
SET rd, immrd = imm in every active lane
TID rdrd = lane_id (this is threadIdx.x % 32)
ADD/SUB/MUL rd, ra, rblane-wise arithmetic on active lanes
MOD rd, ra, immrd = ra % imm
LT rd, ra, rbrd = 1 if ra < rb else 0
BRZ ra, targetbranch if ra == 0the divergence point
JMP targetunconditional
RECONVreconvergence marker (pops the mask stack)
HALTwarp done

This mirrors real hardware closely enough to be honest: PTX has @p bra (predicated branch) and the hardware tracks reconvergence via a stack of (mask, pc) entries (classic pre-Volta SSY/SYNC discipline — see WARMUP Ch. 3 for the Volta+ nuance).

2. How divergence is executed

When BRZ finds that some active lanes branch and some don't:

  1. Push (current_mask, reconvergence_pc) onto the mask stack.
  2. Run the fall-through path with the not-taken lanes' mask.
  3. At RECONV, switch to the taken path with the complementary mask.
  4. At RECONV again, pop the stack — full mask restored, paths rejoined.

Every instruction executed counts once per issue, regardless of how many lanes were active — that's the physical cost model: an instruction with 1 live lane costs the same issue slot as one with 32.

3. What the benchmark shows

solution.py runs the same workload four ways and reports issue counts:

KernelBranch patternPredicted costWhy
uniformno branch1.0×baseline
warp_uniformall lanes agree on the branch~1.0×uniform branch — no split
even_oddtid % 2~2× on the divergent regiontwo serialized half-warps
max_divergent32-way switch on tid~32×each per-lane path issues with 1 live lane

Expected output (yours will match exactly — it's deterministic):

kernel            issues   util%  slowdown
uniform                6   100.0     1.00x
warp_uniform           7   100.0     1.17x
even_odd              10    70.0     1.67x
max_divergent        226    58.4    37.67x

util% = lane-cycles doing committed work ÷ (issues × 32) — the simulator's version of Nsight's "warp execution efficiency" metric you'll meet in Phase 02.

Two readings worth pausing on:

  • even_odd shows 1.67×, not 2.0× — because only the divergent region doubles; the shared prologue (TID/MOD) issues once. Real kernels behave the same way: divergence cost is proportional to how much code sits inside the split. Isolate the divergent region and you'll measure the clean 2×.
  • max_divergent utilization is 58%, not 3% — the 32-way dispatch chain (compare/branch per case) runs fully populated; only each one-lane body wastes 31 lanes. The slowdown is still catastrophic (37×). This mirrors real switch lowering on GPUs.

4. Extension exercises (do them — this is the lab)

  1. Reconvergence stack depth: nest two divergent branches and assert the stack reaches depth 2 and unwinds correctly. (The skeleton handles one level; nested divergence is where naive implementations break.)
  2. VOTE_ANY rd, ra: warp-cooperative primitive — rd = 1 in all active lanes if any active lane has ra != 0. This is __any_sync(). Then explain in a comment why warp votes need no memory traffic on real hardware.
  3. Predication: add SELP rd, ra, rb, rp (lane-wise select) and rewrite even_odd without branches. Compare issue counts — you've just reproduced the compiler's if-conversion decision (WARMUP Ch. 4).
  4. Latency hiding mini-model: simulate 4 resident warps where LD stalls a warp for 40 cycles; a scheduler issues from any non-stalled warp each cycle. Measure utilization with 1, 2, 4 warps and observe the occupancy curve flatten — Chapter 5 made quantitative.

5. Common pitfalls

  1. Forgetting to count issues for instructions with few active lanes — the whole point is that issue cost is mask-independent.
  2. Reconverging at the branch target instead of the post-dominator — your even_odd will produce wrong values in half the lanes.
  3. Letting inactive lanes commit writes — masks must gate writes, not reads.
  4. Treating JMP inside a divergent region as a reconvergence — it isn't; only RECONV pops.

6. What this lab proves about you

You can explain SIMT execution mechanically — masks, stacks, issue costs — which is the difference between "GPUs don't like branches" hand-waving and the engineering-leader version: "this kernel's branch is warp-uniform, it's free; that one is per-lane data-dependent and our profile shows 41% warp efficiency, here's the restructuring." Phase 02's profiler sessions will show you these exact numbers on real silicon.