Reduction 2: One Atomic Per Value

Move stage 1's sum onto the GPU with the smallest change that is genuinely correct: one thread per element, and one atomicAdd each into a single output. It gives the right answer. Then look at what it costs.

Languages: CUDA C++, Python, Rust, HIP C++

Lesson

Same array, same answer, same N = 4,194,304. What changes is who does the adding.

In stage 1 one core walked the array and one accumulator absorbed every value in order. Here every element gets its own thread, and all of those threads have to reach the same destination. Write three things:

1. This thread's global index, from `blockIdx.x`, `blockDim.x`, and `threadIdx.x`. Identical arithmetic to the elementwise stages.
2. A bounds check, because the grid is rounded up to cover N and the last block has threads with no element to read.
3. The add itself. Not `*out += X[i]`: two threads would read the same old value and one update would vanish. `atomicAdd(out, X[i])` is a CUDA built-in that performs the read, add, and write as one indivisible operation, so no contribution is lost no matter how the threads interleave.

The trusted runner owns everything around your kernel. It fills X, sets `out` to zero before the launch, runs your kernel once per test size, and copies the single result back. It does no combining of its own, so whatever your kernel leaves in `out` is the answer being judged.

This stage has no performance target, and that is deliberate. Write it, run it, and read the evidence. The number to find is **SM throughput**. Note it down; stage 3 is about moving it.

Low SM throughput is a symptom to investigate; it does not uniquely prove atomic contention. Compare the same workload timing and atomic behavior.

What you will learn

Map one CUDA thread to one input element, with a bounds check for the rounded-up grid.
State what atomicAdd guarantees and why a plain += would lose updates here.
Find SM throughput in the run evidence and read it as a contention signal rather than an occupancy problem.

Starter code

    // 1. Which element does this thread own? Same index arithmetic as the
    //    elementwise stages: which block, how big a block, where inside it.
    // TODO: int i = ...

    // 2. The grid is rounded up to cover N, so the last block has spare
    //    threads. They must not read past the end of X.
    // TODO: if (...) {

    // 3. Every thread adds into the SAME address. A plain *out += X[i] would
    //    lose updates; atomicAdd makes the read, add, and write indivisible.
    // TODO:     atomicAdd(out, X[i]);
    // TODO: }