Predication and Branch Selection

Every conditional in a kernel is a choice between two implementations. The compiler can emit a branch, which costs a serialized execution of whichever paths the warp actually reaches, or it can emit predicated instructions, which execute unconditionally but let individual lanes discard the results. The two are not interchangeable, and the compiler’s decision between them is based on a cost model that has nothing to do with how a scalar processor would evaluate the same code.

The reason the choice exists at all is that a GPU already has a per-lane validity mechanism: the active mask. Predication reuses it. An instruction predicated on a per-lane condition issues for the whole warp, and the lanes whose condition is false write nothing. No stack entry is pushed, no program counter changes, and the warp stays converged. What is given up is the ability to skip work.

Predicated execution and the mask it reuses

The instruction set provides predicate registers, written by comparisons and consumed by other instructions. A compare instruction like setp.lt.f32 p, a, b produces a predicate that is true on some lanes and false on others; a following instruction may carry a predicate guard so that only lanes where the predicate holds commit a result. The masked-off lanes are not branch-skipped. They occupy the same instruction issue and the same datapath for the same cycle, and simply do not write back.

The same hardware is used for the branch mechanism described in divergence and the reconvergence stack, which is why the two approaches compose. Predication is the degenerate case where the mask narrows and widens without the program counter ever leaving the straight-line path.

Two details distinguish a predicated instruction from a branch. First, a false predicate suppresses the write but not the issue: the datapath is busy for both outcomes. Second, a false predicate suppresses memory operations entirely. A predicated load whose predicate is false does not compute an address, does not access memory, and cannot fault. That property is what makes guard clauses and bounds-checked loads safe to write in straight-line form: the out-of-range lanes issue the instruction and are masked off before the address is formed.

A select instruction is the second half of the idiom. Rather than predicating two separate assignments, the compiler computes both candidate values into registers and uses a select to choose per lane. On architectures that provide both min/max and select, a clamp, an absolute value, or a conditional sign flip becomes two or three unconditional instructions with no branch anywhere.

The cost comparison that drives if-conversion

If-conversion is the transformation of a conditional region into straight-line predicated code. The compiler applies it when the straight-line cost is lower than the expected branching cost, and both sides are estimable.

Let the taken path be L_t instructions and the not-taken path L_f. A branch that the warp splits incurs roughly L_t + L_f issue slots for the two phases, plus the branch and reconvergence overhead, and the mask is partial during each phase. The predicated version incurs L_t + L_f issue slots as well if both paths are materialized, or fewer if the compiler can merge common subexpressions and use selects. When the paths are equal in length and the branch is genuinely divergent, the two are close, and predication avoids the branch overhead and keeps the warp at full mask width for surrounding code.

The comparison separates when the branch is not divergent. If the condition is uniform across the warp, a branch costs only the taken path, L_t or L_f, whichever is taken, because the hardware detects the all-ones or all-zeros ballot and jumps without pushing a stack entry. Predication costs both paths always. So for uniform conditions, branching is strictly better and the gap grows with path length. This is the reason that conditions testing warp-level or block-level quantities should stay as branches, and the reason compilers apply if-conversion mainly to small, data-dependent regions.

/* Two forms of the same clamp. The branchless form issues unconditionally
 * for every lane; the branched form skips the clamp when no lane needs it. */
__global__ void clamp_branchless(float *out, const float *in, float lo, float hi)
{
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    float v = in[i];
    out[i] = fminf(fmaxf(v, lo), hi);       /* no divergence, always issued */
}

__global__ void clamp_branched(float *out, const float *in, float lo, float hi)
{
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    float v = in[i];
    if (v < lo || v > hi)                    /* two tests, often uniform */
        v = fminf(fmaxf(v, lo), hi);
    out[i] = v;
}

The second kernel pays the two comparisons and, on the lanes that need it, the clamp; the first pays the clamp everywhere. Which wins depends on the fraction of lanes in range and on how the data is distributed: a question of data layout, not of instruction count.

The most common real threshold is a small constant. Compilers typically if-convert regions of a handful of instructions and branch around larger ones, because beyond that size the probability that all lanes take the same path is low enough that the serialized cost approaches the sum of both paths, and the predicated version at least keeps the pipeline without branch penalties. The exact limit is a compiler heuristic and varies by version and optimization level; it is not part of any programming model contract.

Short-circuit evaluation and why it reintroduces branches

The C and C++ languages require short-circuit evaluation of && and ||: the right operand is evaluated only if the left does not determine the result. On a warp, that requirement cannot be honoured by predication alone, because whether the right operand is evaluated differs per lane. The compiler must emit branches to satisfy the language, and the region between the two operands becomes divergent.

The effect is that a compound condition of the form if (p && q) costs two branch points rather than one, and each one is a potential divergence site. Rewriting as if (p & q) on integral predicates evaluates both sides unconditionally and produces a single mask, which is often cheaper when q is a cheap register comparison but wrong when q has side effects or is expensive. The same applies to ternaries: a nested ternary chain is usually converted into selects, while a ternary with a function call on one side is not.

Early return is the other common shape. A guard clause if (i >= n) return; produces a divergent region of exactly one instruction, which is the cheapest possible divergence, and it prevents all downstream instructions from issuing for inactive lanes. The alternative, predication of the entire body, would issue the body for lanes that have no work, so the guard is the right form even though it introduces a branch.

Loop-carried conditions and trip-count divergence

Loops differ from conditionals because the test runs every iteration. A loop whose trip count is uniform across the warp never diverges at the test, so the branch is effectively scalar and the cost is the trip count. A loop whose trip count varies per lane diverges at the exit test on every iteration for the lanes that have finished, and the warp runs the full body until the longest lane’s count is exhausted.

The usual fix is not predication but redistribution: instead of giving each lane a variable number of iterations, give each lane the same number and mask the excess work off at the end. This is the standard pattern for kernel tails: a grid-stride or bounds-checked loop where the extra lanes execute a predicated no-op rather than exiting early. It trades wasted lane slots for a converged loop.

Where the trip count cannot be equalized, unrolling the interior loop reduces the number of divergent exit tests proportionally to the unroll factor, at the cost of register pressure. That trade belongs with the occupancy arithmetic in register allocation and the occupancy ceiling: the unrolled body fits fewer warps, and the divergence saved may be worth less than the latency hiding lost.

What selection costs in the memory system

The two forms also differ in their access patterns. A predicated load with a partially true predicate issues the load for the active lanes only, so the transaction is formed from fewer addresses; a branch that skips the load entirely issues nothing. Conversely, a branch that sends lanes to different addresses on the two paths can produce two separate sets of transactions where a predicated version would have produced one.

The pattern to watch for is a conditional access to a location that only some lanes need. Predicating it costs a transaction carrying the addresses of the true lanes; branching around it costs nothing on the false side but reorganizes the warp’s address stream when the sides differ. Because the memory system charges per transaction rather than per lane, the branching version is usually better when the false-side fraction is large, and the predicated version is better when both sides touch nearby addresses.

None of this changes the accounting that governs how much the choice matters: what a warp issues with a partial mask is a lost fraction of a lane slot, described in active-mask accounting. Predication reduces the number of divergent phases but never the unconditional work, so it helps most exactly where divergence is shallow and the paths are cheap, which is the case it was designed for.