使用 __shfl_down_sync 进行归约
无需屏障即可完成线程束归约
使用 __shfl_down_sync 进行归约 是 CoddyKit 上的免费 CUDA Academy 课时。 这是第 2 节课，共 4 节。 你可以在下方免费阅读本课时的完整内容 — 然后在浏览器中使用内置代码编辑器和全天候 AI 导师进行实践。 这是 CUDA Academy 学习路径的一部分，你的进度在网页和 CoddyKit 应用中同步。 CUDA Academy 课程共包含 4 节课。
本课时的部分内容尚未翻译，以英文显示。
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. ✨
常见问题解答
「使用 __shfl_down_sync 进行归约」课时是免费的吗？
是的 — 「使用 __shfl_down_sync 进行归约」的完整文本可在网页上免费阅读。要进行交互式练习（内置代码编辑器和全天候 AI 导师）并解锁 CUDA Academy 课程的其余内容，请升级到 CoddyKit PRO。 CUDA Academy 课程共包含 4 节课。
「使用 __shfl_down_sync 进行归约」这节课中我会学到什么？
无需屏障即可完成线程束归约 你通过在浏览器中直接运行的动手代码来练习 CUDA Academy，全天候 AI 导师会在你学习这节课的过程中回答你的问题。
学习 CUDA Academy 需要有经验吗？
无需任何先前经验。CoddyKit 上的 CUDA Academy 课程适合初学者到高级学习者，你可以从这里开始或从头开始，按照自己的节奏学习。 这是第 2 节课，共 4 节。
「使用 __shfl_down_sync 进行归约」课时需要多长时间？
大多数 CoddyKit 课程大约需要 5–10 分钟。每节课都很精短且互动，所以你能稳步进步，并在网页和应用中从离开的地方继续。
我能在这节 CUDA Academy 课中编写并运行代码吗？
能。每节 CUDA Academy 课都包含内置代码编辑器，你可以在浏览器中直接编写并运行真实代码，并获得即时 AI 反馈 — 无需本地设置。