Divergence and the Reconvergence Stack

A warp has one program counter and thirty-two lanes. When those lanes evaluate a condition and disagree, the program counter cannot go to two places, so the hardware serializes: it runs one side of the branch with the disagreeing lanes disabled, then the other side, then rejoins. The cost is not the arithmetic of the extra instructions. It is that a warp executing with a small active mask occupies a full issue slot while delivering a fraction of the lane work, and the fraction can be as low as one lane in thirty-two.

The mechanism that decides who is disabled and when the paths rejoin is a per-warp stack of masks. Understanding it explains which code shapes diverge expensively, which diverge harmlessly, and why the compiler’s opinion about reconvergence sometimes differs from the source’s structure.

Serialization is a mask change, not a second execution

Consider the smallest case: if (a[i] > 0) x = f(x); else x = g(x); with the condition true for half the lanes. The warp executes the compare, and the result is a thirty-two-bit value with half its bits set: a ballot. The branch instruction consumes that mask. Because some lanes want the taken path and some do not, the warp pushes an entry onto its reconvergence stack recording the mask of lanes that have been put aside and the address at which they should resume, then jumps into the taken side with only the true lanes active.

While executing f, the warp’s active mask is that half of the lanes. Every instruction issued costs one issue slot and delivers work for sixteen lanes. When the taken side finishes, the hardware pops the stack entry, restores the full mask, and jumps to the else side with the complementary lanes active: the sixteen that took the false path. Then it pops again and resumes after the branch with all lanes active.

The warp pays four phases: the compare, executed at full mask; the true path; the false path; and the join, again at full mask. If the two sides are each a single instruction, the warp spends two slots on the two paths where a non-divergent warp would spend one. If the sides are twenty instructions each, it spends forty instead of twenty. The ratio is the number of distinct paths, not the size of the divergence. That is why divergence is cheap when the paths are short and expensive when they are long. The cost is the serialized path length, not the number of lanes that disagree.

/* Divergence cost is path length times distinct paths, not lane count. */
__global__ void weighted(float *out, const float *in, int n)
{
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i >= n) return;                  /* one path, mostly uniform */
    float v = in[i];
    if (v > 0.0f)                        /* divergent here */
        out[i] = v * 2.0f;               /* path A: 1 instruction */
    else
        out[i] = -v;                     /* path B: 1 instruction */
}

The guard clause at the top deserves separate comment because it is the one divergence every kernel has. The final warp of a grid usually has inactive lanes, so the bounds check produces a mask that is not all-ones and the return path is executed with the survivors. The cost is one extra instruction for one extra phase, which is why the guard is cheap and should not be removed in the name of avoiding divergence.

What the stack actually holds

Each entry in a warp’s reconvergence stack records three things: the address the warp is executing, the mask of lanes active there, and the reconvergence address at which the lanes that were put aside expect to resume. The third field is the one the compiler fills in, and its value decides where the paths meet again.

For an if/else, the compiler knows the join point (the instruction after the else block) and emits it as the reconvergence address in the push. Both sides are entered with that address on the stack, so whichever side runs, the pop returns to the join. This is why structured conditionals reconverge exactly where the source says they do.

For a loop with a divergent trip count, the structure is different. The stack entry for the loop records the mask of lanes still iterating. Lanes that exit early are removed from the active mask on each iteration test, and the loop continues until the mask is empty, at which point the entry is popped. The consequence is that a loop runs for as many iterations as its slowest lane requires, with the other lanes masked off for the tail. A loop where lane 0 iterates once and lane 31 iterates a thousand times costs a thousand iterations of the loop body, not the average of a thousand and one.

The mask stack is shallow. Only a few divergence points can be live in a warp at once, and the depth is a hardware constant that no programming model exposes. The compiler lowers control flow so that the live count stays within it, which is one reason machine code’s control-flow graph can look quite different from the source’s nesting: what the hardware counts is how many distinct reconvergence points are outstanding at an instruction, not how many braces the source has.

Where reconvergence points are placed

The compiler chooses push and pop sites, and its choice is visible in the cost. Two heuristics dominate.

The first is immediate post-dominator reconvergence: join at the earliest point that all paths from the branch must reach. For a diamond if/else, that is the instruction after the else block. For an if with no else, it is the instruction after the if body, with the false lanes masked during the body. This is optimal when the paths are balanced and the join is genuinely common.

The second is to hoist the join later. If one side of the branch is much longer than the other, or if the short side’s work could be profitably merged with work after the join, the compiler may keep the lanes split across a longer region so that instructions common to both sides are not issued twice. The trade is between issuing shared instructions once with a partial mask and issuing them twice with a fuller mask; neither dominates universally, which is why the choice is heuristic.

A third effect is worth naming because it confuses intuition: a branch whose condition is uniform across the warp still costs the branch instruction but no serialization. When the ballot is all-ones or all-zeros, the hardware takes the branch like a scalar machine and no stack entry is pushed. Kernels whose branches test block-level or warp-level quantities (threadIdx.x < k with k a multiple of thirty-two, or a flag in shared memory that the whole block reads) therefore diverge in the source and not in the machine. This is the basis of the standard trick of arranging work so that conditions align with warp boundaries, discussed in predication and branch selection.

Why divergence is not the same as wasted work

The serialization model says the warp issues each side of a divergent branch in turn. It does not say the inactive lanes are doing useful work, but it also does not say the time is wasted in proportion to the mask. The distinction matters for two reasons.

First, issue slots are the scarce resource, and a partially active warp consumes a slot that a fully active warp could have used. On a machine with several warp schedulers per SM, a scheduler whose only ready warps are half-active issues at half the lane throughput it could otherwise reach. That is the sense in which divergence is a throughput loss: it converts lane capacity into instruction slots. Occupancy can hide it, because another warp’s full-mask instruction can occupy the datapath in the same cycle. The scheduler picks one warp per cycle, so a partially masked warp does not block others; it merely occupies the slot it takes.

Second, memory behaviour is mostly unchanged. The lanes that are masked off do not issue loads, so they generate no addresses and no transactions. A divergent branch that halves the active mask halves the memory traffic as well as the arithmetic, which means the bandwidth cost is proportional to useful work rather than to warp width. The exception is instruction fetch: when a warp diverges, both paths execute, so both paths’ instructions are fetched and issued even though each is executed by only part of the warp. Instruction cache pressure therefore tracks the total code size of the paths a warp takes, not the number of lanes that take them.

The modern complication: lanes that do not stay together

On architectures with independent thread scheduling, each lane carries its own program counter, and the hardware is not obliged to keep the lanes of a warp in lockstep. Divergence still costs, but the cost is no longer a strict two-phase serialization; the scheduler may interleave lanes from different paths in the same warp, and the compiler inserts explicit reconvergence points where the source requires lanes to be together. A warp-level shuffle after a divergent region is the clearest case: it is meaningless if the participating lanes are at different instructions.

The practical consequences are subtle. Code that relies on implicit reconvergence after a divergent region needs an explicit synchronization primitive on those architectures. And the performance model changes: divergence can be partly hidden by interleaving, so the naive path-length multiplication overestimates the cost. Both effects are covered in per-lane program counters and independent scheduling.

The divergence that remains unavoidable is the one where lanes need genuinely different data at the same instruction. That is not control divergence at all but a consequence of the memory system, and the fix belongs to the memory hierarchy rather than the control-flow model.

What the stack mechanism leaves is a design rule with a precise justification: divergence costs the sum of the distinct path lengths, so the cheapest control flow is the one with the fewest distinct paths per warp, and the next cheapest is the one where those paths are short. Both properties are properties of how work is assigned to lanes, not of how the branch is written.