The CUDA Handbook ▸ Chapter 8. Streaming Multiprocessors

8.4 Conditional Code

The hardware implements “condition code” or CC registers that contain the usual 4-bit state vector (sign, carry, zero, overflow) used for integer comparison. These CC registers can be set using comparison instructions such as ISETP, and they can direct the flow of execution via predication or divergence. Predication allows (or suppresses) the execution of instructions on a per-thread basis within a warp, while divergence is the conditional execution of longer instruction sequences. Because the processors within an SM execute instructions in SIMD fashion at warp granularity (32 threads at a time, 64 on AMD’s wave64 hardware), divergence can result in fewer instructions executed, provided all threads within a warp take the same code path.

To ground the discussion, Listing 8-1 gives the SASS that Hopper (SM 9.0) generates for a small kernel that counts the steps of the Collatz iteration for each element of an array, written with the grid-stride loop this book’s kernels use throughout. The inner per-element loop runs a data-dependent number of times, which makes it a natural source of both predication and divergence:

__global__ void
collatz( unsigned *out, const unsigned *in, size_t N )
{
    for ( size_t i = blockIdx.x*blockDim.x + threadIdx.x;
                 i < N;
                 i += blockDim.x*gridDim.x ) {
        unsigned n = in[i], steps = 0;
        while ( n > 1 ) {
            n = (n & 1) ? (3*n + 1) : (n >> 1);
            ++steps;
        }
        out[i] = steps;
    }
}

/*0160*/       BSSY  B0, 0x280 ;/*0170*/       IMAD.MOV.U32 R5, RZ, RZ, RZ ;/*0180*/       IMAD.X R7, RZ, RZ, R7, P0 ;/*0190*/       ISETP.GE.U32.AND P0, PT, R4, UR8, PT ;/*01a0*/       ISETP.GE.U32.AND.EX P0, PT, R7, UR9, PT, P0 ;/*01b0*/       ISETP.GE.U32.AND P1, PT, R2, 0x2, PT ;/*01c0*/  @!P1 BRA   0x270 ;/*01d0*/       IMAD.MOV.U32 R5, RZ, RZ, RZ ;/*01e0*/       IMAD.MOV.U32 R3, RZ, RZ, 0x1 ;/*01f0*/       LOP3.LUT R0, R2, 0x1, RZ, 0xc0, !PT ;/*0200*/       VIADD R5, R5, 0x1 ;/*0210*/       ISETP.NE.U32.AND P1, PT, R0, 0x1, PT ;/*0220*/       IMAD  R0, R2, 0x3, R3 ;/*0230*/   @P1 SHF.R.U32.HI R0, RZ, 0x1, R2 ;/*0240*/       ISETP.GT.U32.AND P1, PT, R0, 0x1, PT ;/*0250*/       IMAD.MOV.U32 R2, RZ, RZ, R0 ;/*0260*/   @P1 BRA   0x1f0 ;/*0270*/       BSYNC B0 ;

8.4.1 Predication

Due to the additional overhead of managing divergence and convergence, the compiler uses predication for short instruction sequences. The effect of most instructions can be predicated on a condition; if the condition is not TRUE, the instruction is suppressed. This suppression occurs early enough that predicated execution of instructions such as load/store and TEX inhibits the memory traffic that the instruction would otherwise generate. Note that predication has no effect on the eligibility of memory traffic for global load/store coalescing: the addresses specified to all load/store instructions in a warp must reference consecutive memory locations, even if they are predicated.

Predication is used when the number of instructions that vary depending on a condition is small; the compiler uses heuristics that favor predication up to about 7 instructions. Besides avoiding the overhead of the divergence-management machinery described below, predication also gives the compiler more optimization opportunities (such as instruction scheduling) when emitting microcode.

The ternary operator in C (? :) is considered a compiler hint to favor predication.

The Collatz loop body of Listing 8-1 is a compact example. The ternary (n & 1) ? (3*n + 1) : (n >> 1) becomes a single predicated instruction rather than a branch: the compiler computes the odd-case value 3*n + 1 unconditionally, then lets a predicated shift overwrite it for the even lanes.

/*0210*/       ISETP.NE.U32.AND P1, PT, R0, 0x1, PT ;
/*0220*/       IMAD  R0, R2, 0x3, R3 ;
/*0230*/   @P1 SHF.R.U32.HI R0, RZ, 0x1, R2 ;

R0 has already been set to n & 1; the ISETP sets P1 true for the even lanes, IMAD forms 3*n + 1, and @P1 SHF – a right shift by one, i.e. n >> 1 – executes only on the lanes whose predicate is set, overwriting the odd-case value there. No branch is taken, and both cases cost the warp the same handful of instructions.

8.4.2 Divergence and Convergence

Predication works well for small fragments of conditional code, especially if statements with no corresponding else. For larger amounts of conditional code, predication becomes inefficient because every instruction is executed, regardless of whether it will affect the computation. When the larger number of instructions causes the costs of predication to exceed the benefits, the compiler will use conditional branches. When the flow of execution within a warp takes different paths depending on a condition, the code is called divergent.

NVIDIA has never fully documented how the hardware supports divergent code paths, and has changed the implementation across generations. The hardware tracks which threads in a warp are active with a bit mask and suppresses execution for the inactive threads, much as predication does. Through the Pascal generation, divergence was managed with a branch synchronization stack: before a divergent branch, an SSY instruction pushed the active mask and a reconvergence address onto the stack, and a .S suffix on a later instruction popped it, steering any threads that had not taken the branch down the other path. Lindholm et al. describe the model:

If threads of a warp diverge via a data-dependent conditional branch, the warp serially executes each branch path taken, disabling threads that are not on that path, and when all paths complete, the threads reconverge to the original execution path. The SM uses a branch synchronization stack to manage independent threads that diverge and converge. Branch divergence only occurs within a warp; different warps execute independently regardless of whether they are executing common or disjoint code paths¹⁶.

Volta reworked this scheme. Independent thread scheduling (Section 7.5) gave each thread its own program counter, so reconvergence is no longer a single stack of masks but is coordinated through explicit convergence barriers. A BSSY instruction establishes a barrier and names the point at which the divergent threads should reconverge; a matching BSYNC waits on the barrier, releasing the warp once the participating threads have all arrived. As with the older stack, none of this appears in PTX – it is visible only in the SASS emitted by cuobjdump – and the compiler omits it entirely where it can prove a branch cannot diverge.

The inner loop in Listing 8-1 shows the modern form. A BSSY (offset 0x0160) establishes barrier B0, naming 0x0280 – the instruction just past the loop – as the reconvergence point. Threads whose n is already 1 fail the entry test and branch straight to the exit (@!P1 BRA 0x270); the remaining threads iterate, each leaving the loop through the predicated back-edge (@P1 BRA 0x1f0) as its own sequence reaches 1. The BSYNC at 0x0270 holds the reconverged threads until the stragglers arrive, so the warp proceeds in lockstep from there.

Not every loop earns a barrier. The enclosing grid-stride loop compiles to a plain conditional back-edge (@!P0 BRA, outside the fragment shown) with no BSSY/BSYNC around it: because every lane runs the same, count-controlled number of iterations, the compiler can prove the loop does not diverge and omits the convergence machinery entirely. The barriers appear only where the hardware actually needs them – here, around the data-dependent inner loop.

The important thing to realize about branching in CUDA is that in all cases, it is most efficient for all threads within a warp to follow the same execution path.

8.4.3 Special Cases: Min, Max and Absolute Value

Some conditional operations are so common that they are supported natively by the hardware. Minimum and maximum operations are supported for both integer and floating point operands and are translated to a single instruction. Additionally, floating point instructions include modifiers that can negate or take the absolute value of a source operand.

The compiler does a good job of detecting when min/max operations are being expressed, but if you want to take no chances, call the min()/max() intrinsics for integers, or fmin()/fmax() for floating-point values.

Clamping to the unit interval is handled the same way. __saturatef(x) returns 0 when x is negative, 1 when x exceeds 1, and x itself in between (and 0 for a NaN) – a single instruction, exposed directly from the floating-point saturation (.sat) modifier that many instructions can apply to their output. Saturation to [0.0, 1.0] is pervasive in graphics, where it keeps colors in range.

Deep learning has since motivated a wave of additional conditional primitives, since neural networks lean heavily on clamped, low-precision arithmetic. Native min/max were extended to the 16-bit floating-point types – __hmax/__hmin and their NaN-propagating __hmax_nan/__hmin_nan, the packed __hmax2/__hmin2, and __nv_bfloat16 equivalents in cuda_fp16.h and cuda_bf16.h – and execute natively on Ampere and later. The integer instruction set gained a family of fused conditional operations, each collapsing a common idiom into a single instruction: three-way __vimax3_s32/__vimin3_s32, add-then-compare __viaddmax_s32 (max(a + b, c)), __vibmax_s32/__vibmin_s32 (which return the larger or smaller operand and set a predicate identifying which one won), and _relu-suffixed forms that clamp the result at zero – __vimax_s32_relu computes max(max(a, b), 0), the rectified-linear activation itself. Signed, unsigned, and packed 16-bit (_s16x2) overloads round out the set; the SIMD-intrinsics header has the full roster.