Warp divergence penalties: Exercises — Hardware deep dive

This page is a self-test ladder for Warp divergence penalties (6.2.11). Work top to bottom: each level lifts the demand — from spotting divergence, to computing its cost, to analysing nested/loop structure, to designing fixes, to mastering the trade-offs.

Cover each solution and try first. Reveal only after committing to an answer.

Before we start, two symbols we will reuse everywhere, defined from zero:

The reason we add and not average: in SIMT there is one instruction decoder feeding 32 lanes. It can point at only one path's instructions at a time. Threads on the other path are still "there" — they are just masked: their outputs are thrown away, but the cycles still tick. So a path that even one thread takes costs its full length. Keep that in mind — it is the trap behind half the exercises.

The picture below is the mental model for every exercise on this page. Study it before you start: the top bar shows the ideal world where two paths (A = 12 cycles, B = 20 cycles) overlap in parallel, so the warp finishes at $max (12, 20) = 20$ . The bottom bar shows the diverged world where the same two paths run back-to-back, so the warp finishes at $12 + 20 = 32$ . The amber gap on the right is the pure divergence waste — exactly the quantity you compute as $T_{div} - T_{ideal}$ in the problems below.

Level 1 — Recognition

Recall

Solution L1.1 No divergence. Every thread index 0–31 satisfies < 32, so all 32 threads evaluate the condition to true and take path A together. The else path B is taken by nobody, so it is never issued. Since all threads share one path, there is nothing to serialize. Divergence needs disagreement inside one warp.

Worked example

L1.2 — Spot the divergent line Which line below causes intra-warp divergence? (warp = threads 0–31)

int idx = threadIdx.x;          // (a)
float x = data[idx];            // (b)
if (idx % 2 == 0) y = x*x;      // (c)
data[idx] = y;                  // (d)

Recall

Solution L1.2 Line (c). The predicate idx % 2 == 0 is true for even lanes (0,2,…,30) and false for odd lanes (1,3,…,31). Within the same warp the 32 threads split — 16 true, 16 false — so the hardware must issue the body once with the even-mask and once (empty here, the implicit else) with the odd-mask. Lines (a), (b), (d) have no branch: every thread does the same op on its own address, which is uniform control flow (see coalescing, a separate concern from divergence).

Level 2 — Application

Recall

Solution L2.1 Both paths are taken by at least one thread, so both count. $T_{ideal} = max (15, 25) = 25 cycles$ $T_{div} = 15 + 25 = 40 cycles$ $Efficiency = \frac{25}{40} = 0.625 = 62.5%$ The thread count (8 vs 24) never entered the arithmetic — only whether a path is taken and how long it is.

Worked example

L2.2 — Four-way switch One warp runs a 4-way if/else if chain. Path costs and how many threads take each:

path	cycles	threads
0	20	8
1	30	10
2	25	8
3	5	6

Find $T_{div}$ , $T_{ideal}$ , efficiency.

Recall

Solution L2.2 All four paths have at least one taker, so all four serialize: $T_{div} = 20 + 30 + 25 + 5 = 80 cycles$ $T_{ideal} = max (20, 30, 25, 5) = 30 cycles$ $Efficiency = \frac{30}{80} = 0.375 = 37.5%$ This matches Example 2 in the parent note — data-dependent 4-way branching is brutal.

Recall

Solution L2.3 Two paths are taken (lane 0 → path1; lanes 1–31 → path2), so both count fully: $T_{div} = 100 + 50 = 150 cycles, T_{ideal} = max (100, 50) = 100$ $Efficiency = \frac{100}{150} = 0.6 \overline{6} = 66.7%$ The "weighted" guess $50 + 100/32 \approx 53.1$ cycles is wrong because the 100-cycle path is not shared across 32 threads — it costs its full length even for a single lane.

Level 3 — Analysis

Worked example

L3.1 — Nested branches, worst case

if (A) { if (B) p1; else p2; }
else   { if (C) p3; else p4; }

Path costs: p1=10, p2=14, p3=8, p4=12. In one warp, the predicates split so that all four leaves get at least one thread. Give $T_{div}$ and efficiency. Then give the best-case $T_{div}$ if the warp is arranged so only p1 and p3 are ever reached.

Recall

Solution L3.1 Worst case — 4 distinct paths reached: $T_{div} = 10 + 14 + 8 + 12 = 44 cycles, T_{ideal} = max (10, 14, 8, 12) = 14 cycles$ $Efficiency = \frac{14}{44} \approx 0.318 = 31.8%$ Best arranged case — only p1 and p3 taken (e.g. B true whenever A, C true whenever not-A): $T_{div} = 10 + 8 = 18 cycles, T_{ideal} = max (10, 8) = 10 cycles$ $Efficiency = \frac{10}{18} \approx 0.556 = 55.6%$ Nesting depth did not set the cost — the number of distinct leaves actually reached did. See the tree figure.

Recall

Solution L3.2 The loop back-edge is a branch: on each iteration, lanes whose i has reached their counts[idx] drop out (masked), but the warp must keep running iterations until the max-count lane (20) finishes. So the warp executes $max = 20$ iterations worth of issue slots for all 32 lanes, of which only $\sum = 320$ are useful. $Utilisation = \frac{320}{32 \times 20} = \frac{320}{640} = 0.5 = 50%$ Half of every iteration slot after the short lanes exit is wasted on masked lanes.

Recall

Solution L3.3 $Eff_{X} = \frac{m a x ( 40 , 40 , 40 )}{120} = \frac{40}{120} = 33.3%$ $Eff_{Y} = \frac{m a x ( 10 , 10 , 10 , 10 , 10 )}{120} = \frac{10}{120} = 8.3%$ Warp Y is worse. Same total serialized time, but Y's ideal (best single path) is far shorter, so more of Y's time is pure divergence waste. Lesson: more, shorter, equally-populated paths are the most divergence-inefficient shape — the ideal shrinks while the sum stays put.

Level 4 — Synthesis

Recall

Solution L4.1 idx % 2 alternates every lane, guaranteeing a split. Replace with a warp-granular predicate:

int warp_id = idx / 32;
if (warp_id % 2 == 0) heavy(); else light();

Now all 32 lanes of a warp share the same warp_id, so a warp is entirely heavy or entirely light — one path each. For an all-heavy warp: only path taken is 60 cycles, $T_{div} = 60 = T_{ideal}$ : $Efficiency = \frac{60}{60} = 100%$ The total work is unchanged, but the divergence penalty is eliminated. This is the alignment strategy from the parent note, and it feeds directly into occupancy planning.

Recall

Solution L4.2 First, where does the rule come from? Two ways to run a two-path conditional:

Branch (serialize): pay $T_{A} + T_{B}$ — each path issued once, back-to-back.
Predicate (no branch): every lane runs both paths and masks the unwanted result, so it also pays $T_{A} + T_{B}$ of work — but it deletes the branch instruction, the mask-stack push/pop, and the reconvergence, which cost roughly one "extra max-path" worth of overhead. So predication's effective charge is about $2 max (T_{A}, T_{B})$ : you do the bigger path's work twice (once real, once masked) with no branch machinery. Set predication's charge below branching's charge: $branch charge T_{A} + T_{B} > predication charge 2 max (T_{A}, T_{B}) \Rightarrow predicate is cheaper.$ The inequality is exactly the crossover: it holds only when the two paths are short and balanced (so their sum exceeds twice the larger), and fails when one path dwarfs the other (then $2 max$ blows past the sum and the real branch wins).

Now apply it with $T_{A} = 4$ , $T_{B} = 3$ : $T_{A} + T_{B} = 7, 2 max (T_{A}, T_{B}) = 2 \times 4 = 8.$ Since $7 < 8$ , the serialized-branch sum is already less than the predication charge — this branch is so short and balanced that a real branch is the raw-cost equal-or-better choice, yet in practice compilers still predicate here because it removes the divergence stall and reconvergence entirely at essentially the same cycle count. For very lopsided long branches ( $T_{A} + T_{B} > 2 max$ ) you'd instead prefer the real branch so the big path isn't paid by lanes that don't need it. Related idea on CPUs: branch prediction, which GPUs lack.

Recall

Solution L4.3 Sorted, each warp is uniform → single path → $T_{div} = T_{ideal}$ , efficiency $= 100%$ for every warp. The four warp-flavours cost 20, 30, 25, 5 respectively; the average per-warp time is $\frac{20 + 30 + 25 + 5}{4} = \frac{80}{4} = 20 cycles/warp .$ Compare: the mixed warp spent 80 cycles to do the same four kinds of work. Sorting bought a $4 \times$ speed-up here purely by killing divergence — the classic parallel-algorithm preprocessing trick, also a staple of CUDA optimization patterns.

Level 5 — Mastery

Recall

Solution L5.1 Price each warp with the two-number method. A warp that splits into two taken paths costs the sum; a uniform warp costs its single path.

Warp 0 (uniform, one 30-cycle path): $T_{div} = 30$ .
Warps 1, 2, 3 (each splits into 30 and 18, both taken): $T_{div} = 30 + 18 = 48$ each.

Sum the issued cycles across all four warps: $Total issued = 30 + 3 \times 48 = 30 + 144 = 174 cycles.$ The all-uniform ideal charges each warp only its longest single path once ( $max (30, 18) = 30$ ), across four warps: $T_{ideal}^{block} = 4 \times 30 = 120 cycles.$ $Block efficiency = \frac{120}{174} \approx 0.690 = 69.0%.$ One clean warp cannot rescue three dirty ones — divergence is per-warp and the penalties simply add across the block.

Recall

Solution L5.2 Option A — every warp splits, both paths taken: per-warp $= 40 + 10 = 50$ ; four warps $= 4 \times 50 = 200$ cycles. Option B — 2 warps uniform-heavy ( $40$ each) + 2 warps uniform-light ( $10$ each): $2 \times 40 + 2 \times 10 = 80 + 20 = 100 cycles.$ $Speed-up = \frac{200}{100} = 2.0 \times .$ Sorting halved the issued cycles by making every warp internally uniform. This is the synthesis of everything above: recognise divergence (L1), price it (L2), analyse the split (L3), pick a fix (L4), and prove the win at scale (L5).

Recall

Solution L5.3 Divergent region for both: $25 + 15 = 40$ cycles.

Warp P: $40$ cycles total.
Warp Q: after the paths reconverge, all 32 lanes are active again and share one instruction stream, so the 10-cycle tail is issued once: $40 + 10 = 50$ cycles. The tail is uniform control flow — no split, no serialization — so it is added once, not summed per path. Reconvergence is exactly the point where the divergence stack pops and the warp goes back to full-width SIMT.

Recall

Self-test: the one formula to leave with How do you price any divergent region? ::: $T_{div} = \sum (lengths of paths with \geq 1 taker)$ ; $T_{ideal} = max (those lengths)$ ; efficiency $= T_{ideal} / T_{div}$ . Populations don't affect timing; uniform code before/after runs once.