100 DAYS OF GPU Course preview
← Glossary
CUDA glossaryExecution model
CC 7.5

What is warp divergence in CUDA?

What happens when lanes of one warp take different branches, so the hardware runs both paths with some lanes masked off.

Count path-runs, not branches. One instruction goes to all 32 lanes of a warp at once, so when the lanes disagree about a condition the hardware runs the if, then runs the else, switching off the lanes that are not on the path it is currently executing. The warp is charged the sum of the paths its lanes took, never the average and never the longer of the two, which means one lane going the other way costs exactly what sixteen would. Take a block of 256 threads with half the work on each side: split inside every warp and the block issues 16 path-runs, split on warp boundaries and it issues 8, for identical arithmetic.

So the predicate decides everything, and two predicates that look nothing alike can be the same branch. threadIdx.x & 1 puts odd and even lanes on opposite sides, interleaved. (threadIdx.x & 31) < 16 cuts each warp down the middle. Different pictures, the same 8 splits per block, and day 22 timed them at 2.02x and 2.00x against a version that splits nothing. NVIDIA states the fix as a formula rather than as advice: "A trivial example is when the controlling condition depends only on (threadIdx / WSIZE) where WSIZE is the warp size. In this case, no warp diverges because the controlling condition is perfectly aligned with the warps." Divide the index by 32 before you test it and the condition cannot change inside a warp.

This course's own first draft got that wrong, which is why the entry exists. It claimed if (threadIdx.x < 16) was the cheap alternative to if (threadIdx.x % 2). That predicate is wrong twice over. It splits the one warp holding thread 16, so it diverges. And in a 256-thread block it sends 16 threads down path A where the other sends 128, so any timing that shows it winning is reporting the work it skipped. Publish the work split beside every divergence number or the comparison says nothing.

The habit worth avoiding is deleting ifs. A branch every lane agrees on is an ordinary branch: blockIdx.x & 1 splits no warp and timed at 1.00x, and the if (i < n) bounds check that ends nearly every kernel in this course splits at most one warp in a grid of thousands. What costs is how often the lanes disagree, not how much code sits inside the branch, and moving the same one-line branch inside the loop took the lane-parity kernel from 2.02x to 2.29x. Since Volta the warp is not guaranteed to come back together at the join either, which is independent thread scheduling, and it is where divergence stops being a throughput problem and starts being a correctness one, as soon as a warp shuffle is in the branch.

Measured

Day 22 passes the predicate to one kernel as a runtime argument, so every row runs the same instructions in the same order and only the identity of the threads on each path moves. All four send half the grid down each side and both sides are the same length. Tesla T4, driver 595.84, CUDA 12.6 (V12.6.85), built with nvcc -O3 -arch=sm_75, 4096 blocks of 256 threads, 4096 multiply-adds per thread.

Predicate Warps split, of 8 Time vs uniform
threadIdx.x & 1 8 3.272 ms 2.02
(threadIdx.x & 31) < 16 8 3.235 ms 2.00
threadIdx.x / 32 & 1 0 1.616 ms 1.00
blockIdx.x & 1 0 1.617 ms 1.00

Rows one and two are the finding. They are the same branch in different clothes, and a page that contrasts them is teaching a difference the hardware does not have. Rows three and four both branch and neither pays, because a warp whose lanes agree runs one path.

Moving that branch inside the loop, so the split happens on every iteration instead of once:

Kernel Time vs no branch
no branch 1.615 ms 1.00
per-iteration, lane parity 3.703 ms 2.29
per-iteration, warp parity 2.457 ms 1.52

Every row above moved 8,388,608 bytes and ran 4096 multiply-adds per thread. There is no branch-efficiency counter here: day 22 ran under no profiler, the split counts came from a constexpr predicate the compiler checks, and the times came from CUDA events after a warm-up per kernel. Nsight Compute would add the counter and needs root on a stock driver, which Colab does not give you.

Diagram

occupancy-stepper, set to the divergence axis. One 256-thread block drawn as 8 warp rows of 32 lane boxes, path A filled and path B hollow, with a path-run tally beside each row. threadIdx.x & 1 alternates within every row, (threadIdx.x & 31) < 16 fills the left half of every row, and both tally 16. threadIdx.x / 32 & 1 fills whole rows alternately and tallies 8.

Alt text: "The same half-and-half work split, arranged three ways across eight warps. Splitting inside each warp costs sixteen path-runs per block whether the halves interleave or sit side by side, and splitting on the warp boundary costs eight."

Code

From code/day22-divergence/divergence.cu. splitWarpsPerBlock walks the block's 8 warps at compile time and counts the ones whose lanes disagree, so the claim this page rests on turns the build red rather than sitting in prose.

static_assert(splitWarpsPerBlock(kLaneParity) == kWarpsPerBlock,
              "threadIdx.x & 1 must split every warp in the block");
static_assert(splitWarpsPerBlock(kLaneHalf) == kWarpsPerBlock,
              "a threshold inside the warp splits every warp too");
static_assert(splitWarpsPerBlock(kWarpParity) == 0,
              "threadIdx.x / 32 is constant across a warp");

The same constexpr predicate feeds the kernels and the CPU reference, so the page, the check and the hardware cannot drift into describing different programs.

Related terms

Where you meet this

Sources

Byline

Written by: pending. Reviewed by: pending. Written on: pending. Last checked on: pending. The numbers came off the verification node on 2026-08-30, and this entry stays a draft until a named author and a different named reviewer sign it.