Reduction 3: Replace Global Atomic Contention

Stage 2 left you with a correct kernel spending almost all of its time waiting, because every one of the 4,194,304 threads queued to update the same 4 bytes. Keep every thread and every value. Change only where the combining happens: fold inside each block first, then issue one atomicAdd per block instead of one per element.

Languages: CUDA C++, Python, Rust

Lesson

The starter is the kernel you wrote in stage 2, and nothing about it is wrong. That is what makes this stage different from a bug hunt: you are not repairing a wrong answer, you are removing contention from a right one.

At the moment each of a block's 256 threads makes its own trip to the same global address. Let them combine locally first:

1. Each thread stores its value into a shared-memory slot. Shared memory is on-chip and private to the block, so resident blocks fold independently; the 16,384 blocks are scheduled in resource-limited waves.
2. The block folds pairwise: 256 values become 128, then 64, and on down to 1. That is log2(256) = 8 steps, each one done in parallel by the threads that still have work.
3. Thread 0 issues a single atomicAdd with the block's total.

The atomic count drops from N = 4,194,304 to one per block, 16,384 of them. That 256x is what moves the throughput number.

Two things about the tail and the barrier, because they are where correct-looking code goes wrong:

- **Guard the value, not the thread.** Writing `if (i >= N) return;` looks equivalent to the bounds check you wrote in stage 2, and it is not. A lane that leaves early never writes its shared slot, and the fold reads that slot anyway. Uninitialized shared memory can easily happen to be zero, in which case the answer comes out right and tells you nothing. Store the additive identity explicitly and let every lane keep participating: `s[tid] = (i < N) ? X[i] : 0.0f;`.
- **A barrier is block-wide.** `__syncthreads()` is only well defined when the whole block converges on it, which is the second reason not to let some lanes leave early. You need one after the initial store and one after each fold step, so that what you read next is what the previous step actually wrote.

The launch already reserves 1024 shared bytes, which is 256 threads times 4 bytes, so `extern __shared__ float s[]` has room without your touching the launch config.

After Run, the number to put next to your stage 2 result is SM throughput. The gate is 50%; a correct block reduction lands around 91%.

Blocks execute in scheduled waves subject to resource limits. A throughput percentage is supporting evidence; compare elapsed time as well.

What you will learn

Explain why a correct atomicAdd kernel can still be slow: contention on one address serializes updates that the hardware must apply one at a time.
See why that cost shows up as collapsed SM throughput rather than a wrong answer or a lower occupancy.
Fold a block's values in shared memory with a pairwise tree in log2(blockDim.x) steps.
Guard the value rather than the thread: an early return leaves a shared slot unwritten, which the fold reads anyway.
Go from one atomic per element to one atomic per block.

Starter code

    // Your stage 2 kernel: correct, and serialized on one address.
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < N) {
        atomicAdd(out, X[i]);
    }