Warp divergence penalties: Foundations — Hardware deep dive

This page assumes you have seen none of the words in the parent note. We build every one — thread, warp, SIMT, mask, cycle, branch — from a picture, in an order where each idea only uses ideas already built. When you finish, go read the parent note and every symbol will already mean something.

1. A thread — one worker following one recipe

Why does the topic need this? Because a GPU's whole trick is running thousands of these workers at once. Before we can talk about many, we must be crystal clear on one: a thread is exactly one worker, one recipe, one step at a time.

Each worker in the figure has the same recipe card but a different idx — a personal ID number. That ID is the only thing that makes worker #3 behave differently from worker #17. Hold onto idx; it is the source of almost all divergence.

Recall What single value makes one thread behave differently from its neighbour running the same code?

Its thread index idx ::: the per-thread ID number, e.g. threadIdx.x.

2. `idx` — the personal ID number

Picture 32 workers standing in a row, wearing numbered jerseys 0 through 31. When the recipe says if (idx < 16), workers 0–15 answer "yes, that's me" and workers 16–31 answer "no". The recipe is identical; only the jersey number differs. That single fact will later cause the whole warp to split.

3. The warp — 32 workers marching in lockstep

Here is the picture that makes everything else make sense:

Look at the single conductor's baton (the instruction pointer, in burnt orange) pointing at one step. All 32 workers must be on that step. There is not one baton per worker — there is one baton for the whole team of 32. This is the entire reason divergence hurts: a team sharing one baton cannot have two workers on two different steps simultaneously.

(AMD builds the same idea with 64 threads and calls it a wavefront. Same picture, wider team.)

4. SIMT — "one instruction, told to many workers"

Why this exact tool and not plain "run each thread independently"? Because independent threads would each need their own baton, decoder, and clock — expensive. SIMT answers the question "how do I command 32 workers with the cost of commanding one?" Answer: give them one shared command and let each apply it to different data.

The catch, which the whole topic is about: one shared command means every worker must be able to obey the same step. If step 5 is "turn left" but worker #17's recipe says "turn right," the shared baton is stuck — and Section 6 shows what the hardware does about it.

Recall SIMT lets 32 threads share what, saving hardware cost?

One instruction decoder / one instruction pointer (baton) ::: broadcast to all 32, each applying it to its own data[idx].

Related deeper notes: 6.2.8-SIMT-execution-model and 6.2.9-Warp-scheduling.

5. A branch — the fork in the recipe

Picture the 32 marching workers reaching a fork. A sign asks a question — "is your idx less than 16?" Each worker answers with its own jersey number, so they don't all point the same way.

If every worker answers the same (all "yes" or all "no"), the team turns together down one path. No problem — the shared baton still works, everyone on the same step. This is a uniform branch.
If workers answer differently, the team is split. This is divergence, and the shared baton cannot point two ways at once.

That impossibility is the entire crisis. We need one more tool to see how the hardware escapes it.

6. The mask — a row of on/off switches

This is the hardware's clever escape. It cannot point the baton two ways, so instead it does this:

Point the baton at path X. Turn the mask on only for the threads that chose X; turn everyone else off. Run path X. The off threads idle — present, powered, but their results are thrown away.
Point the baton at path Y. Flip the mask: on for the Y-threads, off for the X-threads. Run path Y.
Both paths done → turn all switches back on → reconverge and march together again.

Read the two rows of switches in the figure. Notice path X and path Y run one after the other, never together. That is the serialization: the two paths that could have been simultaneous are now stacked in time. The greyed-out (masked-off) workers in each row are the wasted capacity — the whole cost of divergence lives in those grey squares.

Recall When threads diverge, does the hardware run both paths at once?

No — it runs them one after another, using the mask to switch which threads are active on each pass ::: the masked-off threads idle and burn a clock tick.

7. A cycle — the tick of the hardware clock

Why measure in cycles rather than seconds? Because seconds depend on the specific chip's speed, but the ratio of wasted-to-useful cycles is the same on every chip. When the parent note writes $T_{A} = 10$ cycles and $T_{B} = 8$ cycles, it is counting these ticks. Two paths that could overlap now cost $10 + 8 = 18$ ticks instead of $max (10, 8) = 10$ .

You now have every symbol the parent note uses. idx, warp, SIMT, branch, mask, cycle, sum-vs-max — each is a picture. The parent's formulas are just arithmetic on these pictures.

8. How the foundations feed the topic

Read top to bottom: nothing on this map uses a box above it that wasn't already built. The two roots (a single thread, and its idx) grow into the warp, which needs SIMT to be worth building, which makes branches dangerous, which forces masks, whose cost we count in cycles — and that cost is the topic.

9. Where to go next

The lockstep rule in depth: 6.2.8-SIMT-execution-model
How warps get picked to run each cycle: 6.2.9-Warp-scheduling
The other way memory access order hurts you: 6.2.10-Memory-coalescing
How CPUs face a related fork problem differently: 8.4.2-Branch-prediction
Designing algorithms so warps don't diverge: 9.3.5-Parallel-algorithmsdesign
Keeping enough warps alive to hide the penalty: 6.2.12-Occupancy-optimization
Practical fixes on real GPUs: 10.1.3-CUDA-optimization-patterns

Equipment checklist

Cover the right side; can you answer before revealing?

A thread is ::: one stream of instructions — a single worker following one recipe, one step at a time.
idx is ::: the unique whole-number ID (0,1,2,…) that makes otherwise-identical threads touch different data.
A warp is ::: a fixed team of 32 threads forced to execute the same instruction at the same instant, sharing one instruction pointer.
The "one baton for the whole team" picture represents ::: the single shared instruction pointer — the whole warp is always on the same recipe step.
SIMT stands for and means ::: Single Instruction, Multiple Thread — one broadcast command applied by many threads to their own data.
A branch is ::: a fork in the recipe (if/else, while) where threads may head different ways.
Warp divergence is ::: threads in one warp answering a branch differently, so the shared baton cannot point two ways.
An execution mask is ::: a row of 32 on/off switches deciding which threads act on the current instruction; off threads idle.
Reconvergence is ::: turning every mask switch back on after all divergent paths finish, so the warp marches together again.
A cycle is ::: one clock tick, the unit we count divergence cost in.
Diverged cost versus ideal cost is ::: sum of all path lengths ( $\sum T_{i}$ ) versus the maximum single path ( $max T_{i}$ ).