The if. The moment a warp reaches
if (x[i] < 0) on data where some lanes say yes and some say no,
the fiction of lockstep meets its awkward question: one program counter, two opinions. The
hardware's answer is blunt and beautiful — run both paths, one after the other, switching
off the lanes that shouldn't be there — and it is enforced by two small mechanisms: the
active mask and the reconvergence stack. Understanding them is
understanding why some GPU code runs at full speed and some at a thirty-second of it, on the same
hardware, for the same instruction count.
Alongside the warp's shared program counter lives a 32-bit active mask — one bit
per lane. Bit
Now the branch. Suppose lanes 0–15 evaluate the condition true and lanes 16–31 false. The warp cannot go both ways at once, so it goes both ways in turn:
0x0000FFFF (bits 0–15 set): lanes 0–15 execute the if-body while lanes 16–31 ride
along disabled.0xFFFF0000: lanes 16–31
execute the else-body while lanes 0–15 are switched off.0xFFFFFFFF and the warp marches on, whole again.
Count the cost. If each body is
Here is the whole life cycle on an 8-lane mini-warp (everything scales by 4 for the real thing). The data gives lanes 1, 3, 5 and 7 negative values, so the warp splits down the middle:
Note what the diagram is not showing: parallelism between the two paths. There is none. The then-pass and the else-pass are serial — the same 8-lane engine runs twice, each time half-empty. That serialisation, not the branch instruction itself, is the cost of divergence.
How does the hardware know where to put the warp back together? At a divergent branch it pushes an entry onto a small per-warp reconvergence stack: the address of the join point plus the masks still owed. The join point chosen is the branch's immediate post-dominator — the first instruction that every path out of the branch must eventually reach. The warp then runs one side; when it reaches the join address it pops the stack, swaps in the other side's mask, and jumps back to run that; when the second side arrives at the join, the pop restores the full mask and execution continues as one.
A stack, not a register — because divergence nests. An if inside an
if splits an already-split warp: 16 active lanes become 8 and 8, each pass now
running at a quarter occupancy, with two entries stacked. Every level of disagreement halves the
active fraction and doubles the pass count, and the floor is grim: worst case, the lanes disagree
about everything and the warp decays into 32 serial passes of one lane each —
switch on the lane id. It runs; it just runs like a very expensive
single-threaded machine, thirty-two times over.
Loops diverge too, and more sneakily. A data-dependent trip count —
while (v[i] > eps) — keeps the warp iterating until the last lane
is done. A lane that finishes after 3 trips waits, masked out, while a stubborn neighbour grinds
through 40; the warp's loop cost is the maximum trip count, not the average. One outlier
lane holds thirty-one hostages.
The mechanism fits in a page of TypeScript: an interpreter with an active mask, executing both sides of every divergent branch under complementary masks, recursing for nesting (the recursion depth is the reconvergence stack). The kernel below branches on sign and then, inside the negative side, branches again on magnitude — watch the mask thin out and the passes pile up:
Same kernel, same lane count — the uniform warp does the whole job in half the issues at 100% efficiency, because it never has to pay for a path its lanes didn't take. Try making all eight values negative: the warp diverges only at the inner branch, and the numbers land in between.
Since the Volta generation, each thread carries its own program counter and call stack, and the hardware may interleave the two sides of a divergent branch rather than strictly finishing one first. This exists for correctness, not speed: with per-thread PCs, a lane spinning on a lock held by a lane in the other path can make progress instead of deadlocking — a progress guarantee that pre-Volta warps could not offer. Do not mistake it for free divergence. Execution is still SIMT: in any given cycle the warp issues one instruction under one mask, and lanes on the other path still sit out. The bookkeeping got smarter; the physics — serial passes, thinned masks, lane-slots burned — is exactly the arithmetic you just learned.
Divergence is a data-layout problem wearing a control-flow costume, and the standard cures all amount to making warps agree:
if (i < N) is uniform within almost every warp — only the
final partial warp diverges.y = c ? a : b compiles to a predicated select — every lane does a little wasted
work, but there is one pass, no split, no stack. Two serialised passes only beat that when the
bodies are big.
Because then you would have built a CPU — thirty-two of them. Fetch, decode, branch handling and
scheduling per lane is exactly the overhead
SIMT exists to amortise: one front end serving 32 ALUs is the deal that buys the GPU its lane
count. The mask-and-stack trick is the minimal patch that lets one front end pretend each
lane has free will, paying only when lanes actually disagree. It is also a genuinely old idea —
the ILLIAC IV in the 1970s let each of its 64 processing elements switch itself off with a
mode bit so they could sit out an instruction, and vector machines carried mask registers
for the same reason. The GPU's contribution was automating the bookkeeping: the compiler and the
reconvergence stack manage the masks so the programmer can just write if.
"My kernel has a big if — half the threads take each side — so I lose half my
performance." Not necessarily! The mask machinery only engages when lanes of the same
warp disagree. If warp 3's thirty-two threads all take the if-side and warp 7's all take the
else-side, nothing is serialised: each warp runs exactly one path under a full
mask, and the scheduler interleaves them like any other pair of warps. Divergence between warps is
free. That is why the avoidance patterns above are all really alignment patterns: you
rarely eliminate the branch — you arrange the data so the disagreement falls between
warps instead of inside them. A condition like threadIdx.x < 32 splits
cleanly on a warp boundary and costs nothing; the same test on threadIdx.x < 16
cuts a warp in half and costs a pass.