Warps, SIMT and divergence

32 threads share one instruction stream; a branch they disagree on runs both ways, with lanes masked off.

Loading the animation…

Concept

A GPU runs threads in groups of 32 called warps. A warp issues one instruction at a time for all of its threads, which the CUDA Programming Guide calls SIMT, single instruction, multiple threads: "full efficiency is realized when all 32 threads of a warp agree on their execution path."

When they do not agree, at an if whose condition differs between lanes, the warp diverges: "the warp executes each branch path taken, disabling threads that are not on that path." In the animation, each row is one issued instruction. With i < 8, lanes 0–7 run the if-path while lanes 8–31 sit masked off (hatched); then lanes 8–31 run the else-path while 0–7 wait; then all 32 carry on together. Both paths cost the whole warp time.

Try the conditions:

i < k splits the warp into two runs of lanes. Make $k$ a multiple of 32 and the whole warp agrees: no divergence, and the untaken path is never issued.
i % k == 0 spreads the active lanes out. With $k = 2$ the even and odd lanes take turns: half the warp idles at every branch instruction.
x[i] < k branches on data, here 32 pseudo-random values from 0 to 99 drawn by the model. Data-dependent branches are the usual source of divergence, and how much it costs depends on the data.
uniform branches on blockIdx.x, the same for every thread of a block: the warp never diverges, whatever the paths' lengths.

Divergence happens only within a warp: "different warps execute independently." So a branch that splits threads at warp boundaries, multiples of 32, costs nothing extra.

Maths

If the warp issues $S$ instructions and $a_s$ lanes are active for instruction $s$ , the fraction of the warp's issue slots doing useful work is

\eta = \frac{\sum_s a_s}{32\,S}.

For a branch taken by $m$ of 32 lanes, with $L_A$ instructions on the if-path, $L_B$ on the else-path and $L_0$ outside it, a divergent warp issues $S = L_0 + L_A + L_B$ instructions for

\eta = \frac{32 L_0 + m L_A + (32 - m) L_B}{32 (L_0 + L_A + L_B)} ,

while a warp that does not diverge issues only one path, and $\eta = 1$ . The default in the animation, $m = 8$ , $L_A = 3$ , $L_B = 2$ and $L_0 = 4$ (the index, the branch and two instructions after it), gives $\eta = (128 + 24 + 48)/288 \approx 69\%$ .

Code

The model's trace of one warp, cut from src/lib/gpu/model.ts: a path no lane takes is not issued.

if (taken)
  for (let j = 0; j < lenA; j++) steps.push({ phase: "A", j, mask: taken });
if (notTaken)
  for (let j = 0; j < lenB; j++) steps.push({ phase: "B", j, mask: notTaken });
for (let j = 0; j < lenC; j++) steps.push({ phase: "C", j, mask: FULL_MASK });

Avoiding divergence

Branch on warp-aligned boundaries. if (threadIdx.x / 32 < k) keeps whole warps together.
Sort or bin the work so threads in a warp take the same path (for example, group the elements that need the slow path).
Keep divergent sections short. The compiler may turn a short branch into predicated instructions, which still issue for the whole warp but avoid the branch itself (see the predication slide below).