Ivan's wiki

This page is read only. You can view the source, but not change it. Ask your administrator if you think this is wrong.
# Warp divergence (CUDA)
All 32 threads in a [[warps-cuda|warp]] share one instruction stream. **Warp divergence** happens when threads in the same warp take different branches, and the hardware has to run both sides one after the other, masking off the threads that are not on the current path.

```c
if (threadIdx.x % 2 == 0)
    a[i] = expensive_f(i);   // 16 lanes active, 16 masked off
else
    a[i] = expensive_g(i);   // then the other 16
```

Both branches execute, so the cost is the sum of the two paths rather than the maximum. In the worst case, a 32-way switch inside a warp runs 32 times slower than a uniform one.

The fix is to make the branch uniform across the warp, so every thread in a warp agrees. Branching on data is fine as long as the data is arranged so that neighbouring threads agree:

```c
if (blockIdx.x % 2 == 0)     // whole blocks branch together, no divergence
    a[i] = expensive_f(i);
```

Some things that look like divergence are not. A branch where every thread in the warp takes the same side costs nothing beyond the test itself. The `if (i < n)` bounds guard found in almost every kernel only diverges in the final warp of the grid, which is negligible.

Short divergent regions are also often cheaper than they look, because the compiler converts them into predicated instructions rather than real branches. Divergence is worth restructuring code for when the two paths are long or when the branch is inside a hot loop, not for a two-line `if`.