__shfl_down_sync for Reductions
Warp reductions without barriers.
__shfl_down_sync for Reductions 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.
Shuffle Moves Registers
A shuffle lets one lane read another lane's register value directly. Data hops between threads with no shared memory and no barrier in between.
Meet shfl_down_sync
The workhorse for reductions is shfl_down_sync. Each lane grabs a value from a lane a fixed number of positions higher in the warp.
float v = __shfl_down_sync(mask, val, offset);Reading the Arguments
The call takes an active mask, the value to share, and an offset. Lane i receives the value held by lane i plus offset.
Add What You Receive
A warp sum just adds the shuffled value back in. Lane i takes its neighbor's number and folds it into its own running total.
val += __shfl_down_sync(mask, val, offset);Halve the Offset
Like a tree reduction, the offset starts at 16 and halves each step: 16, 8, 4, 2, 1. After five steps the whole warp is summed.
for (int o = 16; o > 0; o >>= 1)
val += __shfl_down_sync(mask, val, o);Why Five Steps
A warp has 32 lanes, and 2 to the fifth power is 32. So five halving steps are exactly enough to combine all 32 values into one.
Answer Lands in Lane 0
After the loop, the complete warp total sits in lane 0. The other lanes hold partial junk, so only lane 0 should write the result out.
No Barriers Needed
Because exchange happens through registers in lockstep, you skip syncthreads entirely. The sync suffix already coordinates the masked lanes.
Faster Than Shared
A warp shuffle reduction beats the shared-memory version: fewer instructions, no bank conflicts, and no barrier stalls. It is the fast path for the last 32.
Pass the Right Mask
If some lanes already exited, a full mask is wrong. Capture the live set with activemask so each shuffle only touches present lanes.
Two-Level Reductions
Big reductions often reduce each warp with shuffles, then combine warp results in shared memory. Shuffles handle the cheap inner level beautifully.
Quick Check
Think about which lane holds the answer when the loop finishes.
Recap
You summed a warp with shfl_down_sync: halve the offset five times, and lane 0 holds the total, no barriers needed. Next: ballot and vote. ✨
Frequently asked questions
Is the “__shfl_down_sync for Reductions” lesson free?
Yes — the full text of “__shfl_down_sync for Reductions” 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 “__shfl_down_sync for Reductions”?
Warp reductions without barriers. 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 “__shfl_down_sync for Reductions” 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
- Warps, Lanes, and Masks
- __shfl_down_sync for Reductions
- Ballot and Vote Functions
- Cooperative Groups