Multi-Block Final Reduction
Combining per-block partial sums.
Multi-Block Final Reduction is a free CUDA Academy lesson on CoddyKit — lesson 4 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.
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. 🏁
Frequently asked questions
Is the “Multi-Block Final Reduction” lesson free?
Yes — the full text of “Multi-Block Final Reduction” 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 “Multi-Block Final Reduction”?
Combining per-block partial sums. 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 4 of 4, so you can start here or from the beginning and move at your own pace.
How long does the “Multi-Block Final Reduction” 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