What is warp divergence, why does it cost you, and how would you find and fix it in a kernel?

TL;DR: A warp is 32 threads that share one instruction stream. If a branch splits them, the SM executes each taken path in sequence with the non-participating lanes masked, so the warp's time is the sum of the path lengths instead of the length of the longest path the thread needed. The cost is proportional to how much the paths differ and how many warps contain a mix; a branch that splits every warp with a 100-instruction path and a 5-instruction path runs at about 8% lane efficiency. Nsight Compute reports it as threads executed per instruction (ideal 32). The fix is almost always to reorder data so that threads in the same warp take the same path.

How to approach it

Define the unit first: 32 threads, one program counter per warp before Volta, one per thread after but still one issue per warp. Explain what happens at a branch in two sentences. Then take a concrete branch, count the instructions on each path, compute the warp time and the ideal time, and give the efficiency. Name the metric you would look at, then give the fixes in order of how often they apply: data layout, predication, warp-uniform branching.

A strong answer

A typical situation: a kernel with a data-dependent branch runs at half the throughput of the same kernel on sorted input, with identical instruction counts and identical memory traffic. Nothing is wrong with it. The 32 lanes were taking different paths.

The CUDA Programming Model exposes threads, but the hardware schedules warps. An SM sub-partition issues one instruction per cycle to one warp, and all 32 lanes execute it, each on its own registers. A branch whose condition differs across lanes cannot be issued as one instruction to all 32, so the scheduler runs the taken path with the lanes that want it active and the others masked, then the other path with the masks inverted, then reconverges. Both paths cost issue slots; the masked lanes do nothing useful during theirs.

The cost model for one warp:

inputs: path A = a instructions, path B = b instructions, fraction p of lanes take A
warp time with divergence  = a + b            (both paths issued, one after the other)
useful lane-instructions   = 32 × (p × a + (1 - p) × b)
lane efficiency            = useful ÷ (32 × (a + b)) = (p·a + (1-p)·b) ÷ (a + b)

example: a = 100 (a rare slow path), b = 5, p = 1/32 (one lane per warp takes it)
  warp time = 105 issue slots
  useful    = 32 × (100/32 + 31 × 5/32) = 100 + 155 = 255 lane-instructions
  efficiency = 255 ÷ (32 × 105) = 255 ÷ 3,360 = 7.6%
  compare: if the slow lane were in a warp of its own and the rest in clean warps,
           time is 100 for one warp and 5 for the others, efficiency near 100% for the many

sanity: a warp cannot be more than 32 lanes efficient and cannot take less time than its
        longest path, so 7.6% is between the bounds and the loss is 13x, which matches
        the 105 ÷ (average lane work of 8) intuition

The example is the one that bites in practice: a rare special case (a NaN check, a boundary tile, a token that routes to a different expert) that costs 20x the common path. If the rare lanes are scattered so that almost every warp holds one, every warp pays the slow path. The same rare cases packed into a few warps cost almost nothing overall. Divergence is a property of the mapping of data to lanes, not of the branch.

A scenario: a top-k routing kernel for a mixture-of-experts layer checks if (expert_id == my_expert) per token, with tokens laid out in arrival order. Expert IDs are close to random across consecutive tokens, so each warp of 32 tokens contains several experts, and the per-expert path runs several times per warp. Sorting tokens by expert before the kernel (what every production MoE dispatch does) makes each warp uniform, and the same kernel runs at full lane efficiency. The sort costs one pass over the token indices, which is cheap next to the divergent kernel.

Since Volta, each thread has its own program counter and call stack, which lets diverged lanes interleave and makes lock-free code inside a warp correct, but it does not change the cost: still one issue per warp per cycle, still masked lanes. What it changed is where reconvergence happens, so a kernel that relied on implicit reconvergence at the end of an if should call __syncwarp() before warp-level shuffles.

How to see it. In Nsight Compute, smsp__thread_inst_executed_per_inst_executed.ratio is the average active lanes per issued instruction; 32 is the ceiling and anything under about 28 in a hot loop deserves a look. The Source view shows per-line "Predicated-On Thread Instructions Executed" next to instructions executed, so you can find the exact branch. The Warp State Sampling breakdown does not show divergence directly, but a kernel whose issue slots are busy while achieved FLOPS are low, with no memory stall, is the fingerprint.

The fixes, in the order they usually apply:

1. Change the data-to-lane mapping so that a warp's 32 items share a path: sort or bucket by the branch key, pad boundary tiles to a multiple of 32 so the edge check is warp-uniform, and handle the ragged tail in a separate small kernel or a warp-uniform epilogue.
2. Let the compiler predicate short branches. For paths of a few instructions the compiler emits both with predicate masks and no jump, and the cost is a + b either way, which is what the arithmetic above already charged. Writing x = cond ? f(x) : g(x) with cheap f and g is fine; the pathology is a long path taken by few lanes.
3. Make the branch warp-uniform: compute the condition from something all 32 lanes share (block index, warp index, a value read once per warp) so the hardware takes one path with no masking. __ballot_sync() tells you at run time whether any lane wants the slow path, and the warp can skip it entirely when none do.

The reversal condition: a memory-bound kernel, which Memory-Bound vs Compute-Bound Kernels settles with one division. If the kernel is waiting on HBM most of the time, issue slots are not the scarce resource, and a 2x loss in lane efficiency may change nothing measurable. Check the roofline before spending a week on a branch.

What interviewers probe next

• "Does a branch on threadIdx.x < 16 diverge?" Yes: half the lanes each way inside one warp. threadIdx.x < 32 × k for a whole warp does not; that is the warp-uniform pattern.
• "What about loops with data-dependent trip counts?" The warp runs until its longest loop finishes; lanes that exit early are masked for the remaining iterations. Cost is max trip count, not mean, so sort by length or bucket similar lengths together, which is what batched sequence processing does.
• "Independent thread scheduling on Volta fixed divergence, right?" It fixed forward-progress and correctness for intra-warp synchronization. The issue-slot cost is unchanged.
• "How much does a divergent kernel matter in an LLM?" Little in the GEMMs, which are branch-free; a lot in sampling, top-k, MoE dispatch and tokenization-adjacent kernels, which are small but sit on the decode critical path where a 10x slowdown of a 20 µs kernel is visible at batch 1.

Common mistakes

• Describing the cost as "both paths run" without the multiplier: the loss is (a + b) ÷ (what the lanes needed), and it is large only when a long path is taken by few lanes.
• Removing every if from a kernel, including the cheap ones the compiler would have predicated, and producing slower code with more arithmetic.
• Optimizing lane efficiency in a kernel that is memory-bound, where the issue slots were never the limit.
• Forgetting that a loop with variable trip count is a branch, and that its cost is the maximum across the warp.

Key takeaways

• Warp time under divergence = sum of taken path lengths; efficiency = useful lane-instructions ÷ (32 × sum). One slow lane per warp at 100 vs 5 gives 7.6%.
• The metric: smsp__thread_inst_executed_per_inst_executed.ratio, ceiling 32.
• Divergence is a data layout problem: sort or bucket by the branch key so warps are uniform; pad tails to 32.
• Predication is fine for short paths; warp-uniform conditions and __ballot_sync remove the cost for rare long paths. Check memory-bound first.