多线程块最终归约
合并各线程块的部分和
多线程块最终归约 是 CoddyKit 上的免费 CUDA Academy 课时。 这是第 4 节课，共 4 节。 你可以在下方免费阅读本课时的完整内容 — 然后在浏览器中使用内置代码编辑器和全天候 AI 导师进行实践。 这是 CUDA Academy 学习路径的一部分，你的进度在网页和 CoddyKit 应用中同步。 CUDA Academy 课程共包含 4 节课。
本课时的部分内容尚未翻译，以英文显示。
Blocks Cannot Talk
A reduction within a block is easy, but blocks run independently and cannot synchronize with each other mid-kernel. So one launch cannot sum everything.
Each Block Produces a Partial
So every block reduces its own chunk to one number, a partial sum, and writes it to a small output array indexed by blockIdx.
if (tid == 0)
out[blockIdx.x] = data[0];Now You Have Fewer Values
With 1000 blocks you go from a million inputs to 1000 partials. The hard part is done; only a tiny array remains to combine.
Strategy One: Launch Again
The simplest finish is a second launch of the same kernel on the partials. Repeat until only one value is left.
Recursive Until One
Each pass shrinks the array by the block size. A few recursive launches reduce millions down to a single final sum.
Strategy Two: Atomics
Alternatively, thread 0 of each block can add its partial straight into one global total with atomicAdd, avoiding a second kernel.
if (tid == 0)
atomicAdd(total, data[0]);Atomics Trade Off
Atomics are simple and need only one launch, but many blocks contending on the same address can serialize. With few partials it is usually fine.
Strategy Three: Grid-Stride
A grid-stride loop lets each thread first sum many elements into a register, so far fewer blocks are needed before the final step.
for (int i = gid; i < n; i += gridDim.x * blockDim.x)
sum += in[i];Fewer Blocks, Less Overhead
Doing more work per thread up front means fewer partials and fewer launches. This often beats spawning one thread per element.
Zero the Total First
If you use atomics, remember to zero the global total before launching, or your sum starts from garbage left in that memory.
Pick by Problem Size
Small inputs love atomics for their simplicity; huge inputs favor a two-pass or grid-stride design. Measure on your data to choose.
Quick Check
Think about why a single kernel launch cannot sum the whole array directly.
Recap
Blocks each emit a partial sum, then you combine them with a second launch, atomics, or grid-stride. You can now reduce arrays of any size. 🏁
常见问题解答
「多线程块最终归约」课时是免费的吗？
是的 — 「多线程块最终归约」的完整文本可在网页上免费阅读。要进行交互式练习（内置代码编辑器和全天候 AI 导师）并解锁 CUDA Academy 课程的其余内容，请升级到 CoddyKit PRO。 CUDA Academy 课程共包含 4 节课。
「多线程块最终归约」这节课中我会学到什么？
合并各线程块的部分和 你通过在浏览器中直接运行的动手代码来练习 CUDA Academy，全天候 AI 导师会在你学习这节课的过程中回答你的问题。
学习 CUDA Academy 需要有经验吗？
无需任何先前经验。CoddyKit 上的 CUDA Academy 课程适合初学者到高级学习者，你可以从这里开始或从头开始，按照自己的节奏学习。 这是第 4 节课，共 4 节。
「多线程块最终归约」课时需要多长时间？
大多数 CoddyKit 课程大约需要 5–10 分钟。每节课都很精短且互动，所以你能稳步进步，并在网页和应用中从离开的地方继续。
我能在这节 CUDA Academy 课中编写并运行代码吗？
能。每节 CUDA Academy 课都包含内置代码编辑器，你可以在浏览器中直接编写并运行真实代码，并获得即时 AI 反馈 — 无需本地设置。