Killing Warp Divergence
Reindexing to keep warps busy.
Killing Warp Divergence is a free CUDA Academy lesson on CoddyKit — lesson 2 of 4. You can read the complete lesson below for free — then practise it hands-on in the browser with a built-in code editor and a 24/7 AI tutor. It is part of the CUDA Academy learning path, one of 4 lessons in the course, and your progress syncs across the web and the CoddyKit app.
Warps Run in Lockstep
A warp is 32 threads that execute the same instruction together. When their paths agree, the hardware runs at full speed.
What Divergence Costs
If threads in a warp take different branches, that is divergence. The hardware runs each path serially, leaving some lanes idle and wasting cycles.
The Naive Reduction Diverges
The simple version uses tid % (2*s) to pick active threads. Active and idle threads interleave inside every warp, so each warp diverges hard.
if (tid % (2 * s) == 0)
data[tid] += data[tid + s];Idle Lanes Still Cost
Even though half the threads do nothing, they still occupy the warp. The warp cannot finish until both the active and idle paths are handled.
Reindex by Thread ID
The fix is to map active work to the lowest thread IDs instead of scattered ones. Compute an index from tid and the stride.
int index = 2 * s * tid;
if (index < blockDim.x)
data[index] += data[index + s];Why That Helps
Now the busy threads are contiguous: tid 0,1,2,... all work, the rest all rest. Whole warps are either fully active or fully idle.
Fully Idle Warps Are Free
A warp where every lane is idle just retires with no work. There is no per-lane serialization, so the cost of divergence largely disappears.
The Modulo Trap
The hidden villain was the modulo condition. It scattered active threads across each warp, which is exactly what creates divergence.
Same Work, Better Mapping
You did not change the math or the number of additions. You only remapped which thread does each add, and the warps thank you for it.
It Compounds at Scale
Across thousands of blocks and many steps, removing divergence is a real speedup, often a couple of times faster than the naive kernel.
Still One Snag Left
This version reads neighbors that are interleaved in shared memory, which can cause bank conflicts. The next lesson fixes that too.
Quick Check
Think about what causes warp divergence in the naive reduction.
Recap
You killed divergence by giving work to the lowest thread IDs, so warps are all-active or all-idle. Same math, faster reduction. Next: bank conflicts. 🚀
Frequently asked questions
Is the “Killing Warp Divergence” lesson free?
Yes — the full text of “Killing Warp Divergence” is free to read here on the web, and the CUDA Academy course includes 4 lessons in total. To practise it interactively (a built-in code editor and a 24/7 AI tutor) and unlock the rest of the CUDA Academy course, upgrade to CoddyKit PRO.
What will I learn in “Killing Warp Divergence”?
Reindexing to keep warps busy. You practise CUDA Academy with hands-on code you run directly in the browser, and a 24/7 AI tutor answers your questions as you work through the lesson.
Do I need any experience to start CUDA Academy?
No prior experience is required. CUDA Academy on CoddyKit is structured for beginners through advanced learners; this is — lesson 2 of 4, so you can start here or from the beginning and move at your own pace.
How long does the “Killing Warp Divergence” lesson take?
Most CoddyKit lessons take about 5–10 minutes. Each one is bite-sized and interactive, so you make steady progress and pick up exactly where you left off across the web and the app.
Can I write and run code in this CUDA Academy lesson?
Yes. Every CUDA Academy lesson includes a built-in code editor, so you write and run real code right in your browser and get instant AI feedback — no local setup required.
All lessons in this course
- The Reduction Tree Idea
- Killing Warp Divergence
- Sequential Addressing
- Multi-Block Final Reduction