Branch Divergence and Predication Performance in Warps

Hello! Last lesson separated where a value lives from how threads safely observe it. This time, the question is which threads are doing useful work when an instruction executes. You have already seen that an SM schedules warps and that resident warps help hide latency. But a warp’s 32 threads do not necessarily follow the same path through a kernel.

In this lesson, you will predict when a conditional makes a warp execute multiple paths, when the compiler can instead use predicated instructions, and why neither “a branch exists” nor “a divergence counter is low” is enough to predict performance.

Divergence is a property of a warp, not a branch in the source

A warp executes a common instruction for its participating threads. If every thread in a warp makes the same decision at an if, that warp follows one path. A neighboring warp is free to choose the other path: different warps taking different branches is not, by itself, branch divergence.

If threads within one warp choose different paths, the warp must execute the instructions needed by both groups, with threads not participating in the current path disabled. Watch this short explanation for the basic execution picture.

The programming guide supplies two qualifications that matter when applying that picture to modern hardware.

The supplied image makes the two-path case concrete:

For the condition shown, threadIdx.x < 4, only a warp containing both indices below 4 and indices at least 4 has to cover both paths. In an ordinary block’s first warp, four threads take A, B and the remaining 28 take X, Y. During A and B, the latter group does no work on those instructions; the roles reverse for X and Y.

On GPUs with Independent Thread Scheduling, threads can remain at different program points longer than the simple “split, execute, immediately rejoin” cartoon suggests. That flexibility does not erase the throughput cost of issuing instructions for distinct paths. It also does not justify assuming that threads have reconverged at a particular source line. As the previous lesson noted, code that relies on warp-wide participation must use an appropriate explicit synchronization and participation mask rather than an old lockstep assumption.

Predict the instruction cost before predicting elapsed time

A useful first approximation is to count warp-issued path instructions, not branches in source code or the number of threads choosing each side. Suppose path A requires $C_A$ issued instructions and path B requires $C_B$, ignoring their common condition and continuation. A uniform warp needs approximately $C_A$ or $C_B$. A warp with threads on both paths needs approximately

This is a model of path work, not a formula for kernel runtime. Instruction mix, pipeline throughput, dependencies, cache behavior, and overlap with other warps can all change elapsed time.

Consider a block with two warps and two ways to assign the same number of threads to two tasks. Suppose task A takes 40 arithmetic instructions and task B takes 60:

Both assignments perform the same per-thread arithmetic. The mixed assignment issues instructions for both paths in each warp, often with only half its threads useful on a path. If arithmetic issue is the dominant limit, grouping can have a substantial advantage. Do not infer a guaranteed twofold kernel speedup: other costs may dominate or overlap with that extra issue work.

The split need not be 16–16 to be expensive. A 31–1 split still requires both substantial paths if that one thread must execute its path. Conversely, a short, infrequent branch may add too little work to matter. For example, a bounds check affecting only the final partial warp of a large grid is unlikely to justify reorganizing the entire computation around that one warp.

Memory paths demand another layer of reasoning. Threads disabled on a branch path do not request that path’s loads, but the active threads may access addresses with poor locality or make little use of fetched data. A change that groups threads by branch outcome may improve instruction efficiency while worsening memory access patterns, or vice versa. Carry forward last lesson’s separation between instructions issued, addresses requested, and where those requests are served.

Branch predication trades control flow for masked instruction issue

A source-level if does not guarantee a hardware branch. For a short conditional, the compiler may schedule instructions guarded by per-thread predicates. An instruction with a false predicate has no result write; the CUDA best-practices guide also specifies that it does not evaluate addresses or read operands. Thus a predicated load is not a load by every thread merely because the instruction is issued.

Predication avoids a control-flow split for that conditional, but it does not make the other path’s instructions free. If a short if/else becomes predicated instructions for both sides, the warp may still issue instructions belonging to both sides. The benefit is avoiding branch-control overhead and making a short sequence easier to schedule; the cost is issuing conditional instructions even when their predicates are false for some threads.

This yields an important contrast:

• Mixed decisions, short paths: predication may be cheaper than managing a divergent branch, although both sides can consume issue capacity.
• Uniform decision, substantial unused path: a real branch can skip the unused path for the whole warp. Predicating a long path instead could waste issue capacity.

The compiler chooses whether and how to predicate based on the generated instruction sequence, not merely the apparent length of the C++ block. It may also transform the surrounding code. Treat “the compiler will predicate this” as a hypothesis until you inspect the generated code in the later PTX and SASS module.

Be careful with hand-written “branchless” rewrites. A conditional selection whose two values are computed eagerly can perform both calculations, including loads you meant to avoid. That differs from genuinely predicated instructions with false predicates, for which addresses are not evaluated. A branchless rewrite must preserve correctness as well as improve the relevant hardware cost.

Place decisions at the right granularity

When the algorithm permits it, arranging work so each warp largely agrees on a decision avoids repeatedly paying for two paths. These conditions illustrate the distinction:

// Warps within the block can make different decisions,
// but each complete warp agrees internally.
if (threadIdx.x / warpSize == 0) {
    do_A();
} else {
    do_B();
}

// In a complete warp, neighboring threads choose different paths.
if (threadIdx.x & 1) {
    do_A();
} else {
    do_B();
}

In the first example, the block’s first warp takes A and later warps take B; no complete warp splits on that condition. In the second, each complete warp has threads on both paths. This is about how work is assigned, not a recommendation to replace an algorithm’s actual decision with a thread-ID test.

Reassignment has a price. Grouping like work might require sorting or compaction, an extra kernel launch, temporary storage, or less favorable memory access. The grouping pays off only when its saved path work exceeds those costs. If decisions depend on data, evaluate how often a warp is actually mixed rather than assuming that irregular data always produces maximum divergence.

Loops follow the same principle. If threads in a warp run a common loop body for different iteration counts, the warp can keep issuing body instructions while some threads have already finished. A long tail handled by one remaining thread can therefore have poor useful-lane utilization. A rough model for a simple shared loop body is the largest relevant iteration count in the warp, with fewer participating threads toward the end—not the sum of all 32 iteration counts. Actual compiled control flow and scheduling can complicate that model.

Finally, keep branching distinct from participation. An instruction may have few participating threads because of a divergent path, an early exit, a partial final warp, or a false predicate. A low count of divergent branch events does not prove that most issued instructions have 32 useful threads; predicated code can have low divergence yet substantial masked work. Equally, seeing divergence does not prove it is the kernel’s limiting cost. For a profiling hypothesis, ask which path instructions are issued, how many threads use them, and what other resource limits the kernel.

Takeaways

To predict divergence costs, identify decisions made within each warp. If a warp needs both branch paths, estimate the instruction work of both paths; if whole warps take different paths, each warp needs only its chosen path. Then qualify that estimate using path length, memory behavior, and overlap with other warps.

Predication can remove a control-flow split for a short conditional, but predicated instructions still consume scheduling opportunities; false-predicate threads do not perform their guarded operation or evaluate its address. Neither source syntax nor a single divergence metric tells the whole story.

Next, you will examine scoreboards, dependencies, and pipeline availability—other reasons an SM may have resident warps but still be unable to issue the instruction you expect.