Warp divergence and what it costs

This course's own first draft named day 22 like this:

"Warp divergence. Why if (threadIdx.x % 2) halves throughput and if (threadIdx.x < 16) does not."

research/CURRICULUM.v0.md, day 22, corrected before anything was written

Both of those predicates diverge. A warp is 32 threads issuing one instruction together, warps are cut from consecutive threadIdx.x values, and a condition that changes inside one of those runs of 32 splits the warp however it is spelled. threadIdx.x < 16 splits the warp holding thread 16, and in a block of 256 threads it also sends 16 threads down the first path instead of 128, so it is not a fix, and it is not even the same amount of work.

A branch costs nothing when all 32 lanes agree. This page measures what it costs when they do not, across two predicates that split the warp and two that do not, with the work held identical across all four.

What happens when 32 lanes disagree

The SIMT rule is one sentence long, and the guide states both halves of it:

"A warp executes one common instruction at a time, so full efficiency is realized when all 32 threads of a warp agree on their execution path. If threads of a warp diverge via a data-dependent conditional branch, the warp executes each branch path taken, disabling threads that are not on that path."

CUDA Programming Guide 3.2.2.1, "SIMT Execution Model", https://docs.nvidia.com/cuda/cuda-programming-guide/03-advanced/advanced-kernel-programming.html (checked 2026-08-30)

Read "each branch path taken" as arithmetic. A warp where every lane takes the if runs the if and stops. A warp where one lane takes the else runs the if, then runs the else, and the lanes that are not on the current path sit there switched off.

The warp pays the sum of the paths its lanes take, not the average and not the larger of the two. One dissenting lane costs the same as sixteen.

That is warp divergence, and the second half of the rule is what makes it fixable:

"Branch divergence occurs only within a warp; different warps execute independently regardless of whether they are executing common or disjoint code paths."

Same section, same page.

So the unit that can be wasted is the warp, and the warp scheduler is happy to send neighbouring warps down completely different code. Take a block of 256 threads, 8 warps, where half the threads run a long path A and half run an equally long path B. If the split falls inside each warp, all 8 warps run both paths: 16 path-runs.

If the split falls on warp boundaries, 4 warps run A and 4 run B: 8 path-runs. The same threads did the same arithmetic, and one arrangement issued twice the instructions.

After the branch the warp reconverges and goes back to one instruction for 32 lanes.

Since Volta that is less automatic than it sounds, which independent thread scheduling covers: "Starting with the Volta architecture, Independent Thread Scheduling allows a warp to remain diverged outside of the data-dependent conditional block. An explicit __syncwarp() can be used to guarantee that the warp has reconverged for subsequent instructions." (CUDA C++ Best Practices Guide 13.1, "Branching and Divergence", https://docs.nvidia.com/cuda/cuda-c-best-practices-guide/index.html , checked 2026-08-30.)

For a kernel that only does arithmetic this changes nothing. For the warp shuffles and votes in the next lesson it changes everything, which is why every one of them carries a _sync in its name.

The predicate is the thing, not the branch

The intuition people bring is that branches are expensive on a GPU, so the skill is removing them. That is the wrong target. A branch every thread in the warp agrees on is an ordinary branch.

It takes one instruction and one path, with no waste. Thirty-two warps can take thirty-two different paths and none of them pays for the others.

NVIDIA writes the fix as a formula rather than as advice:

"A trivial example is when the controlling condition depends only on (threadIdx / WSIZE) where WSIZE is the warp size. In this case, no warp diverges because the controlling condition is perfectly aligned with the warps."

Best Practices Guide 13.1, same page, checked 2026-08-30

So run your predicate through that test rather than counting ifs. The mapping it tests never moves: "The way a block is partitioned into warps is always the same; each warp contains threads of consecutive, increasing thread IDs with the first warp containing thread 0" (same section). In a two-dimensional block the thread ID is x + y * blockDim.x, from the guide's Thread Hierarchy section, so a 32 by 8 block puts one row in each warp and a condition on threadIdx.y is warp-uniform while the same condition on threadIdx.x is not.

In a one-dimensional block of 256 threads, with warps cut at 0 to 31, 32 to 63 and so on, which day 21 read back out of the hardware rather than assuming:

Predicate	Warps that split	Threads on path A
threadIdx.x % 2 == 0	8 of 8	128
(threadIdx.x % 32) < 16	8 of 8	128
threadIdx.x < 16	1 of 8	16
(threadIdx.x / 32) % 2 == 0	0 of 8	128
blockIdx.x % 2 == 0	0 of 8	0 or 256

Row 3 is the draft's fix, and the table says both things that are wrong with it at once. It does split a warp, so it diverges. And it puts 16 threads on path A where the row above puts 128, so if it had been timed against row 1 it would have looked faster for a reason that has nothing to do with divergence: it does less work.

Row 4 is the shape you want, and you have already written its cheap cousin. Most kernels in this course end with if (i < n), which splits at most one warp in the whole grid, the one that straddles the end of the data, and guards a single store. Divergence is a cost per warp per branch, so one warp in a grid of thousands is not a cost at all.

Measuring it without changing the work

Full program in code/day22-divergence/divergence.cu. It runs on four rules.

One kernel, one variable. The predicate arrives as a runtime argument, so every row of the first table executes the same instructions in the same order and only the identity of the threads on each path changes. A separate kernel per predicate would let the compiler specialise each one and the rows would stop being comparable.

Every row does the same work. Each of the four predicates sends exactly half the grid's threads down each path, and the two paths are the same length, so the aggregate arithmetic is identical and only its grouping moves. This is the discipline day 11 exists to teach, applied to instructions rather than bytes.

One predicate function, three consumers. The kernels, the CPU reference and the compile-time checks all call the same constexpr function, so the page, the check and the hardware cannot end up describing different programs.

Events, and a warm-up per kernel. Day 9 covers why a host clock around a launch measures the launch, and why the first launch of each kernel pays its own module load.

Here is the predicate:

__host__ __device__ constexpr bool takesPathA(int mode, unsigned int tid,
                                              unsigned int block,
                                              unsigned int phase) {
    switch (mode) {
        case kLaneParity:
            return ((tid + phase) & 1u) == 0u;
        case kLaneHalf:
            return ((tid + phase) % kWarpSize) < (kWarpSize / 2u);
        case kWarpParity:
            return ((tid / kWarpSize + phase) & 1u) == 0u;
        case kBlockParity:
            return ((block + phase) & 1u) == 0u;
        default:
            return true;
    }
}

mode is a kernel argument, so it holds the same value for every thread in the grid and the switch itself never diverges. phase is 0 for the whole of part 1.

Because that function is constexpr, the claim the page rests on is checked by the compiler rather than asserted in prose. splitWarpsPerBlock walks the 8 warps of a block and counts the ones whose lanes disagree:

static_assert(splitWarpsPerBlock(kLaneParity) == kWarpsPerBlock,
              "threadIdx.x & 1 must split every warp in the block");
static_assert(splitWarpsPerBlock(kLaneHalf) == kWarpsPerBlock,
              "a threshold inside the warp splits every warp too");
static_assert(splitWarpsPerBlock(kWarpParity) == 0,
              "threadIdx.x / 32 is constant across a warp");
static_assert(splitWarpsPerBlock(kBlockParity) == 0,
              "blockIdx.x is constant across a warp");
static_assert(splitWarpsPerBlock(kAlwaysA) == 0, "no branch, no split");

A second block of static_asserts in the file checks the other half of the fairness condition, that every mode sends half the block down path A. The file does not build if either is untrue.

The kernel is one branch around two equal chains of fused multiply-adds:

__global__ void mixOneBranch(const float* __restrict__ in,
                             float* __restrict__ out, size_t n, int mode) {
    const size_t i = blockIdx.x * static_cast<size_t>(blockDim.x) + threadIdx.x;
    float x = (i < n) ? in[i] : 0.0f;

    if (takesPathA(mode, threadIdx.x, blockIdx.x, 0u)) {
        for (int k = 0; k < kIters; ++k) {
            x = fmaf(x, kMulA, kAddA);
        }
    } else {
        for (int k = 0; k < kIters; ++k) {
            x = fmaf(x, kMulB, kAddB);
        }
    }

    if (i < n) {
        out[i] = x;
    }
}

Part 2 of the program is that kernel with the if moved inside the loop, so each side is one instruction, plus a third kernel with no branch at all as the baseline.

The Best Practices Guide leaves this case to the compiler: "The compiler replaces a branch instruction with predicated instructions only if the number of instructions controlled by the branch condition is less than or equal to a certain threshold." (Section 13.2, "Branch Predication", same page, checked 2026-08-30. The threshold is not published.)

Note. Nothing here can stop the compiler from removing a branch, and it is allowed to. If all four rows of the first table come back equal, that is what happened, and the fix is to read the SASS rather than to argue with the timer. The README gives the cuobjdump -sass line, and day 46 is where that skill is taught properly.

Results

Measured. Tesla T4, driver 595.84, CUDA 12.6 (V12.6.85), built with nvcc -O3 -arch=sm_75. Captured 2026-08-30 on the project's verification node; full transcript in the page's evidence file.

GPU: Tesla T4 (compute capability 7.5)
40 SMs, warp size 32, 4096 blocks x 256 threads = 1048576 threads
4096 multiply-adds per thread, 8388608 bytes moved per launch

Part 1: one branch, 4096 multiply-adds on each side
predicate                 warps split    time (ms)   vs uniform
------------------------  -----------   ----------   ----------
threadIdx.x & 1                     8        3.272         2.02
(threadIdx.x & 31) < 16             8        3.235         2.00
threadIdx.x / 32 & 1                0        1.616         1.00
blockIdx.x & 1                      0        1.617         1.00
warps split is out of the 8 warps in a block, computed from the predicate

Part 2: the same branch inside the loop, one multiply-add on each side
kernel                          time (ms) vs no branch
----------------------------   ---------- ------------
no branch                           1.615         1.00
per-iteration, lane parity          3.703         2.29
per-iteration, warp parity          2.457         1.52

Every row above moved the same 8388608 bytes and ran the same 4096
multiply-adds per thread. Only the branch changed.

The CUDA 13.0 re-run preserved Part 1 exactly at the level that matters: both split predicates took 2.00x the uniform rows. Part 2 moved materially. The no-branch baseline fell from 1.615 to 1.093 ms, while lane parity measured 2.27x and warp parity measured 2.15x.

In that run, evaluating a uniform branch in the loop accounted for nearly all of the slowdown and splitting the warp added only another 1.06x. The full CUDA 13 table is in the second evidence file.

The predicate decides everything, and the two that look different are the same. threadIdx.x & 1 splits all 8 warps in a block and costs 2.02x. (threadIdx.x & 31) < 16 also splits all 8 and costs 2.00x. They are the same operation: both cut every warp into two halves that must be executed one after the other.

This matters because an earlier draft of this very lesson contrasted those two as though the second were the fast case. It is not, and the measurement above is what settles it. A condition on threadIdx.x is warp-uniform only when it cannot change within a group of 32 consecutive lanes.

threadIdx.x / 32 & 1 and blockIdx.x & 1 split zero warps and cost 1.00x. Both branch, both run the same arithmetic, and neither costs anything, because every lane in a warp takes the same side.

Where the branch sits changes what the timer includes. In the CUDA 12.6 run, moving the same one-line branch inside the loop took the split row from 2.02x to 2.29x. CUDA 13.0 measured a similar 2.27x, but its uniform branch was already at 2.15x. A split branch cannot be priced against the no-branch baseline without also measuring the matching uniform branch.

Part 2 also prices something that is not divergence at all. The warp-uniform per-iteration row splits no warps and still cost 1.52x the no-branch baseline on CUDA 12.6 and 2.15x on CUDA 13.0, because the predicate is evaluated and acted on 4,096 times per thread either way.

Read the two per-iteration rows against each other for the cost of splitting a warp, and both against the baseline for the cost of asking a question in a loop. That first ratio changed from 1.51x to 1.06x between these runs, so it is a measured property of the compiled kernel, not a portable constant.

Run it yourself

A free Colab T4 or any card you own. The build line is in the repo's README:

nvcc -std=c++17 -O3 -arch=sm_75 -o divergence divergence.cu

There is no Compiler Explorer embed. Size is not the reason: two 4 MiB buffers and a fraction of a second of GPU work sit well inside the 20 second run cap. The reason is that every number here is a ratio between two timings.

The free runner is a shared spot instance with no SLA, so a slow row there can be somebody else's kernel rather than yours. Reading the program on Compiler Explorer is fine; timing it is not.

Exercise

Add a fifth predicate to part 1: threadIdx.x < 16, the one at the top of this page. It needs a mode constant, a case in takesPathA, a label and kNumModes raised by one. Write down how many of the block's 8 warps it splits and how many threads it sends down path A before you build.

Put both numbers in static_asserts and let the compiler mark them. Run it, and say why the row is faster than the two split rows above it without being warp-uniform.

Time: 25 to 40 minutes. Submit: the predicted split count, your row's ratio to the threadIdx.x / 32 & 1 row, one sentence on where the speed came from, and the smallest edit to threadIdx.x < 16 that makes it warp-uniform.

Check: the correctness gate covers your row for free, because the host reference and the kernel share takesPathA, so a predicate you got wrong on one side is reported as a mismatch with the thread index that found it. The static_asserts are the real check: guess the split count wrong and the file does not compile. Then read your ratio against the uniform row, which is the number the sentence has to explain.

Hint 1

Warps are cut from consecutive threadIdx.x values: 0 to 31 is one warp, 32 to 63 the next. Write down, for each of the eight warps, whether its 32 threads agree. Then count the path-A threads in the whole block.

Hint 2

A warp that splits pays for both paths and a warp that agrees pays for one. Add up the eight warps of the block that way, for this predicate and for the 16 and the 8 of the two rows above it, and the ratio is the ratio of those totals.

Solution

One warp of the eight splits, and 16 of the 256 threads take path A. The block therefore pays 9 path-runs: warp 0 runs both paths, the other seven run path B only. Against 16 for threadIdx.x & 1 and 8 for threadIdx.x / 32 & 1, the row should land just above the uniform rows and well under the split ones.

That speed is not the absence of divergence, it is the absence of work. The row that looks nearly as good as the warp-aligned one got there by sending 112 fewer threads down path A, which makes it a different program rather than a faster one.

The smallest edit that makes it warp-uniform is threadIdx.x < 32, one whole warp, and notice what that costs: you cannot warp-align a predicate without changing which threads do which work. When the work assignment is fixed, the move is to change which thread owns which element, which is what day 24 does to a reduction.

Take this away instead of the numbers: compare two kernels only when they do the same work, and count path-runs per warp rather than branches per kernel.

Pitfalls

You made the condition simpler and it still diverges. A threshold is not warp-aligned because it is tidy. Run the predicate through the guide's test. If it depends only on threadIdx.x / 32 and blockIdx.x, no warp splits, and if it cannot, count the warps it splits before you claim a fix.

You compare two kernels that no longer do the same work. The fastest way to make a divergent kernel look fixed is to send fewer threads down the expensive path. Report the work split beside every timing, the way this page's tables do, or the comparison means nothing.

You delete every if and replace it with arithmetic. A uniform branch is one instruction and costs nothing, and the compiler may already be turning a short divergent one into predicated instructions. Branchless code that makes all 32 lanes evaluate both sides can be slower than the branch it replaced.

You expect the warp to reconverge the instant the if ends. Since Volta a warp may stay diverged past the join, so warp-synchronous code that worked on Pascal is not correct now. Use __syncwarp() when you need the lanes back together, and the _sync forms of every shuffle and vote. Day 23 and day 27 cover the rules.

You measure divergence on a memory-bound kernel and find nothing. The memory stall can overlap the extra instruction issue, so the branch adds no measured time in that kernel and expensive in the next one. Find out which kind you have before optimising, which is what day 42 and the roofline on day 30 are for.

You profile as a normal user and get nothing. Nsight Compute reports ERR_NVGPUCTRPERM, in full: The user running <tool_name/application_name> does not have permission to access NVIDIA GPU Performance Counters or the Hardware Event System on the target device. On a machine where you have root, sudo ncu works, because the stock driver default RmProfilingAdminOnly is 1. Colab and Kaggle give you neither. NVIDIA's page on the message, checked 2026-08-30: https://developer.nvidia.com/nvidia-development-tools-solutions-err_nvgpuctrperm-permission-issue-performance-counters

Go deeper

• CUDA C++ Best Practices Guide 13.1, "Branching and Divergence", and 13.2, "Branch Predication", for the warp-alignment formula, the reconvergence caveat and the predication threshold: https://docs.nvidia.com/cuda/cuda-c-best-practices-guide/index.html (checked 2026-08-30)
• CUDA Programming Guide 3.2.2.1, "SIMT Execution Model", for the divergence rule and the definition of an active thread: https://docs.nvidia.com/cuda/cuda-programming-guide/03-advanced/advanced-kernel-programming.html (checked 2026-08-30)
• cuda-samples, cpp/2_Concepts_and_Techniques/reduction, NVIDIA's own sequence of reduction kernels, whose early versions change where the branch falls: https://github.com/NVIDIA/cuda-samples/tree/master/cpp/2_Concepts_and_Techniques/reduction (checked 2026-08-30)
• Programming Massively Parallel Processors, 4th edition, chapter 4, on warps, SIMD hardware and control divergence: https://shop.elsevier.com/books/programming-massively-parallel-processors/hwu/978-0-323-91231-0

Next

Day 23 stops working around the warp and starts using it: shuffles and votes move data between lanes with no shared memory and no barrier, and every one of them carries a mask because of the reconvergence rule above.

Day 24 then compares reduction versions. The first version has a branch that splits every warp, while the second moves that branch onto a warp boundary. The ratio you measured here is the most that change can save; day 24 measures the result on a T4, and the answer is less than zero because the same edit adds bank conflicts.

The rest of module 3 keeps testing predicates this way. Reading a predicate is a working habit rather than a lesson topic, and it is one of the things a kernel engineer is paid for.