GPU Glossary
/
Performance·2 min read

Warp Divergence

When threads in the same warp take different execution paths because of control flow, so some must be masked off.

Warp divergence happens when threads in the same warp take different execution paths through control flow statements. Because a warp issues one instruction at a time for all its threads, divergent paths cost extra work.

An example#

Consider this kernel:

__global__ void divergent_kernel(float* data, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        if (data[idx] > 0.5f) {
            // A
            data[idx] = data[idx] * 4.0f;
        } else {
            // B
            data[idx] = data[idx] + 2.0f;
        }
        data[idx] = data[idx] * data[idx];
    }
}

At the data-dependent conditional, some threads in a warp need to run block A and others block B, depending on data[idx]. Given that data dependency and the structure of the CUDA programming model and its implementation in the PTX machine model, neither the programmer nor the compiler can avoid this split inside the warp.

Instead, the warp scheduler handles the divergent paths by "masking" some threads so they don't execute a given instruction. The masks are held in predicate registers.

What the hardware runs#

Here is the generated SASS (Godbolt link):

LDG.E.SYS R4, [R2]                       // L1 load data[idx]
FSETP.GT.AND P0, PT, R4.reuse, 0.5, PT   // L2 set P0 to data[idx] > 0.5
FADD R0, R4, 2                           // L3 store 2 + data[idx] in R0
@P0 FMUL R0, R4, 4                       // L4 in some threads, store 4 * data[idx] in R0
FMUL R5, R0, R0                          // L5 store R0 * R0 in R5
STG.E.SYS [R2], R5                       // L6 store R5 in data[idx]

Step by step:

  1. L1: the data is loaded into R4.
  2. L2: all 32 threads run FSETP.GT.AND together, and each gets its own P0 based on its value in R4.
  3. L3: a bit of compiler cleverness: all threads run the + 2.0f branch, writing to R0.
  4. L4: only threads where P0 is true run the * 4.0f branch, overwriting R0. On this instruction the warp is "divergent".
  5. L5: all threads run the same code again. Once the warp scheduler realigns them by issuing the same instruction on the same clock cycle, the warp has "converged".

Why the compiler does it this way#

The naïve encoding would predicate both L3 and L4. The compiler's version is presumably more efficient: we can trust the compiler, and heuristically it trades cheap, plentiful CUDA Core computation for more expensive flow control.

Key idea In GPU programming it's often better to waste compute (here, an unnecessary FADD for every thread that takes L4) than to add control complexity, even simple predication.

One reason compilers work hard to avoid divergence: on early (pre-Volta) GPUs, divergent warps were always fully serialized. Divergence still reduces efficiency, but modern GPUs with independent thread scheduling don't necessarily pay the full serialization penalty.