A warp issues one instruction per cycle. When threads in a warp take different if/else paths, the warp cannot split. So all threads run the if block first, with the threads that failed the condition masked off, then all threads run the else block with the other set masked off. This is called warp divergence.

// a kernel with warp divergence
__global__ void divergent(float* A, float* out) {
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    float x = A[tid];

    if (x < 0) {
        // ~100 cycles: only the negative lanes are active,
        // the rest sit masked for all 100
        for (int k = 0; k < 50; k++) x = x * 1.01f - 0.5f;
    } else {
        // ~2 cycles: now the negative lanes are the masked ones
        x = x * 2.0f;
    }

    out[tid] = x;   // reconverged, all 32 lanes active again
}

Lanes not on the current path are masked: still holding their slot on the SM, producing no result. Both paths run in full, so the warp runs for the if block's cycles then the else block's.

Step by step through a divergent warp

• Cycle 1: loads a mix of positive and negative values into all 32 lanes. The threads with loaded values are already shown.
  • The condition splits the warp, no cycle consumed yet.
• Cycle 2: the if path runs first. cycles for this example.
• Cycle 3: the else path runs next. cycles for this example.

// negative lanes execute; non-negative lanes are masked (predicated off)
// they occupy the warp for all 100 cycles and produce no result.
if (x < 0) {
    for (int k = 0; k < 50; k++) x = x * 1.01f - 0.5f;
}

// non-negative lanes execute; negative lanes are masked this time.
else {
    x = x * 2.0f;
}

if (blockIdx.x < 16) is decided once per block, before any of its warps start. Every thread in that block takes the same side, so no warp splits.

Divergence is a per-warp, per-threadIdx cost.

A non-divergent warp

• Issue a uniform LOAD
• Issue the uniform instruction.

float x = A[tid];   // all 32 lanes, one cycle - same mix of values as before

// no branch: every lane runs the same 100 cycles of work
for (int k = 0; k < 50; k++) x = x * 1.01f - 0.5f;