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.mdCh. 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:
- Why does
if (tid % 2)cost 2× butif (tid < 32)cost nothing (per warp)? - 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:
| Instr | Semantics |
|---|---|
SET rd, imm | rd = imm in every active lane |
TID rd | rd = lane_id (this is threadIdx.x % 32) |
ADD/SUB/MUL rd, ra, rb | lane-wise arithmetic on active lanes |
MOD rd, ra, imm | rd = ra % imm |
LT rd, ra, rb | rd = 1 if ra < rb else 0 |
BRZ ra, target | branch if ra == 0 — the divergence point |
JMP target | unconditional |
RECONV | reconvergence marker (pops the mask stack) |
HALT | warp 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:
- Push
(current_mask, reconvergence_pc)onto the mask stack. - Run the fall-through path with the not-taken lanes' mask.
- At
RECONV, switch to the taken path with the complementary mask. - At
RECONVagain, 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:
| Kernel | Branch pattern | Predicted cost | Why |
|---|---|---|---|
uniform | no branch | 1.0× | baseline |
warp_uniform | all lanes agree on the branch | ~1.0× | uniform branch — no split |
even_odd | tid % 2 | ~2× on the divergent region | two serialized half-warps |
max_divergent | 32-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_oddshows 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_divergentutilization 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 realswitchlowering on GPUs.
4. Extension exercises (do them — this is the lab)
- 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.)
VOTE_ANY rd, ra: warp-cooperative primitive —rd = 1in all active lanes if any active lane hasra != 0. This is__any_sync(). Then explain in a comment why warp votes need no memory traffic on real hardware.- Predication: add
SELP rd, ra, rb, rp(lane-wise select) and rewriteeven_oddwithout branches. Compare issue counts — you've just reproduced the compiler's if-conversion decision (WARMUP Ch. 4). - Latency hiding mini-model: simulate 4 resident warps where
LDstalls 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
- Forgetting to count issues for instructions with few active lanes — the whole point is that issue cost is mask-independent.
- Reconverging at the branch target instead of the post-dominator — your
even_oddwill produce wrong values in half the lanes. - Letting inactive lanes commit writes — masks must gate writes, not reads.
- Treating
JMPinside a divergent region as a reconvergence — it isn't; onlyRECONVpops.
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.