Warp divergence: when 32 threads disagree

A warp is 32 GPU threads that must run the same instruction in lockstep. So what happens when they hit an if/else and want to go different ways? Pick a branch pattern and step through it.

__global__ void kernel(...) {
    if ( condition(threadIdx.x) )
        doA();     // some threads want this
    else
        doB();     // others want this
}

With 32 threads per warp: if they all take the same branch, the warp runs that path once and stays fully busy. If they diverge, the hardware must run doA() with one group while the other sits idle, then run doB() with the second group while the first sits idle — two passes, half the threads wasted each time. That's warp divergence, and it's why branchy code belongs on the CPU.