Warp Divergence and Why Branches Are Expensive on GPUs

SIMT, in one paragraph

CUDA GPUs execute threads in groups of 32 called warps, and every thread in a warp runs the same instruction at the same time on its own data, a model NVIDIA calls SIMT (single instruction, multiple thread). This is close to SIMD, but with an important difference: each thread has its own program counter conceptually, and the hardware allows individual threads within a warp to take different control flow paths. The catch is that the warp scheduler does not actually run those paths concurrently. It serializes them.

What divergence actually does

When threads in a warp hit a conditional branch and disagree on which way to go, the warp executes both paths sequentially, masking off the threads that are not on the active path for each pass. Consider:

__global__ void divergent_kernel(float* data, int n) {
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    if (tid >= n) return;

    if (data[tid] > 0.0f) {
        data[tid] = expf(data[tid]);   // path A
    } else {
        data[tid] = -expf(-data[tid]); // path B
    }
}

If half the threads in a warp take path A and half take path B, the warp first executes path A with the path-B threads masked off (idle but still occupying a slot), then executes path B with the path-A threads masked off. Total time for that warp is roughly the sum of both paths’ cost, not the max, even though logically each thread only ever does one of them. With full divergence across a 32-wide warp, you can lose most of the parallelism you paid for.

The reconvergence point (where control flow merges back together) determines how long the divergence lasts. Modern architectures (Volta and later) use independent thread scheduling, which allows more flexible interleaving and can reduce some pathological cases like divergent locks, but it does not eliminate the fundamental serialization cost of a warp with threads on different paths.

How costly this actually is

where $P$ is the number of distinct paths threads in the warp take and $T_p$ is the cost of path $p$. In the best case ($P=1$, no divergence) the warp pays only that path’s cost. In the worst case, with as many distinct paths as threads, a warp can pay up to 32x the cost of a single path, though this extreme is rare in practice since most divergent kernels have only two or three branches.

A useful mental model: divergence does not make a warp’s threads run wrong, it just makes the warp behave like a serial machine for however many distinct paths exist, which quietly erodes the SIMT throughput advantage that justified using a GPU in the first place.

Patterns that reduce divergence

The most reliable fix is to restructure data or work assignment so that threads within a warp tend to agree on control flow, since divergence is a property of the warp, not of the algorithm in the abstract.

• Sort or bucket data by predicate before the kernel launch. If you know some threads will take a slow branch, group those elements contiguously so entire warps go one way, not individual threads scattered across warps.
• Replace branches with predication or arithmetic where the branch body is cheap. data[tid] = (data[tid] > 0) ? a : b can sometimes compile to predicated instructions instead of a real branch, especially for short bodies, avoiding divergence entirely for cheap operations.
• Move the branch to block or warp granularity, not thread granularity. If an entire warp (or block) can be routed to the same code path based on a coarser decision, the SIMT model has nothing to diverge on.
• Avoid divergent loop trip counts within a warp. A loop whose iteration count depends on data[tid] forces the warp to wait for its slowest thread every iteration; padding to a common trip count with masked no-ops is sometimes faster than the divergence it avoids.

Practical note

Warp divergence is invisible in a naive correctness check and only shows up in profiling (nsight compute reports warp execution efficiency directly), which is why it is worth checking early rather than after a kernel is already deep in a larger pipeline. A kernel that is correct but divergent will still pass every unit test while quietly running at a fraction of achievable throughput.