TL;DR: Within one warp of 32 threads, mixed branch outcomes force the hardware to run both paths serially with the inactive lanes masked off, so a 2-way split can halve throughput and a 32-way split costs 32x. The fix is to make warps see uniform outcomes (reorder data, predicate cheap branches), and the boundary case interviewers chase is that warp-aligned predicates like
threadIdx.x < 64do not diverge at all.
How to approach it
Define it mechanically (SIMT execution, masked serialization), quantify the worst case, then show you know the boundary condition: divergence only costs you within a warp. The follow-up about if (threadIdx.x < 64) is almost always coming, so set yourself up for it.
A strong answer
A warp is 32 threads executing in SIMT lockstep, one instruction issued for all 32 lanes. When a branch makes threads within the same warp take different paths, the hardware cannot issue two instructions at once, so it serializes: it runs the if-side with the else-threads masked off, then the else-side with the if-threads masked off. Both paths consume issue slots; the masked lanes do dead work. With a 2-way split you can lose up to half your throughput; in the pathological case, a 32-way switch on threadIdx.x % 32, each lane executes alone and you get 1/32 of peak.
The key boundary: divergence is a within-warp phenomenon. If warp 0 takes the if-branch and warp 1 takes the else-branch, there is zero penalty, because warps are independently scheduled. So if (threadIdx.x < 64) does not diverge at all (the boundary at 64 is warp-aligned: warps 0-1 all go one way, warp 2+ the other), while if (threadIdx.x % 2 == 0) is maximally painful. Branching on data values (if (x[i] > 0)) diverges whenever a warp's 32 elements are mixed, common in sparse data, ray tracing, and tokenizer-style code.
Counting issue slots on one concrete warp makes the cost exact instead of hand-wavy. Say if (x[i] > 0) splits a warp 20 lanes true, 12 lanes false, with 10 instructions on the if-side and 6 on the else-side. The warp issues all 16 instructions serially: 10 with 12 lanes masked, then 6 with 20 lanes masked. Useful work is 20 x 10 + 12 x 6 = 272 lane-instructions out of 32 x 16 = 512 issued, 53% efficiency (computed, and Nsight's "average active threads per warp" is measuring exactly this ratio). The same count also shows why sorting the data first wins: bucket the positives and negatives so each warp is all-true or all-false, and every warp runs only its own path at 100% lane efficiency, except the single warp straddling the boundary. One partially divergent warp per boundary, instead of mixed outcomes scattered through all of them, is the entire trade, and it is why the fix reorders data rather than code.
Fixes, in the order I would try them: restructure the data so warps see uniform values (sort or bucket elements by branch outcome before the kernel, paying an O(n log n) sort once and winning on every subsequent pass); make the code branchless where the work is cheap (predication via min/max/ternaries, which the compiler does automatically for short branches, since executing both sides costs less than serialization anyway); or split into separate kernel launches per case when the branches are heavyweight. Worth knowing: since Volta, threads have independent program counters ("independent thread scheduling"), which fixes old deadlock patterns and allows finer-grained reconvergence, but the throughput cost of divergence is unchanged. There is still one issue path per warp.
A practical note for ML infra roles: this is part of why GPUs eat dense linear algebra but struggle on branchy preprocessing, and why you keep tokenization and ragged control flow on the CPU or batch it carefully.
What interviewers probe next
- "Is
if (threadIdx.x < 32) { ... }divergent?" No: that is exactly warp 0 vs everyone else; warp-aligned predicates do not diverge. This is the classic filter question. - "How would you detect divergence in a real kernel?" Nsight Compute's branch efficiency and "average active threads per warp" metrics; below ~80% active threads, start hunting.
- "Both sides of my branch are one instruction each, should I care?" Usually no; the compiler predicates short branches and the cost is trivial. Optimize the long, hot divergent paths.
- "Does this exist on other hardware?" Yes; any SIMD/SIMT machine (AMD wavefronts of 32/64, CPU vector predication) pays for mixed lanes somehow.
Common mistakes
- Claiming any
ifin a kernel is bad. Uniform branches (all 32 lanes agree) are essentially free; blanket "avoid branches" advice signals you have never profiled. - Missing the warp-alignment subtlety and calling
threadIdx.x < 64divergent. This is the exact discriminator interviewers use. - Saying Volta's independent thread scheduling "fixed" divergence. It fixed correctness pitfalls (intra-warp lock patterns), not the serialization cost.
- Proposing fixes that move the branch instead of removing the mixed outcomes per warp. Reordering code does not help if each warp still sees both cases; reordering data does.
Key takeaways
- Divergence serializes paths within a warp only; across warps it is free.
- Warp-aligned predicates (multiples of 32) never diverge; data-dependent ones often do.
- The real fix reorders data so warps see uniform outcomes, not just moving the branch.