Réduction finale sur plusieurs blocs

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.

So every block reduces its own chunk to one number, a partial sum, and writes it to a small output array indexed by blockIdx.

With 1000 blocks you go from a million inputs to 1000 partials. The hard part is done; only a tiny array remains to combine.

The simplest finish is a second launch of the same kernel on the partials. Repeat until only one value is left.

Each pass shrinks the array by the block size. A few recursive launches reduce millions down to a single final sum.

Alternatively, thread 0 of each block can add its partial straight into one global total with atomicAdd, avoiding a second kernel.

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.

A grid-stride loop lets each thread first sum many elements into a register, so far fewer blocks are needed before the final step.

Doing more work per thread up front means fewer partials and fewer launches. This often beats spawning one thread per element.

If you use atomics, remember to zero the global total before launching, or your sum starts from garbage left in that memory.

Small inputs love atomics for their simplicity; huge inputs favor a two-pass or grid-stride design. Measure on your data to choose.

Think about why a single kernel launch cannot sum the whole array directly.

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. 🏁