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