Reduction 4: Warp-Shuffle Block Sum

Lesson

**What your kernel produces here is different, so read this before you write anything.** Stages 2 and 3 wrote one global answer. This kernel writes **one partial sum per block** into `partial[blockIdx.x]`, and nothing on the device combines those values. The driver tab shows the host summing `partial[]` afterwards, purely to check your answer.

That means:

- You are optimizing the inside of **one block**. Getting `partial[blockIdx.x]` right for every block is the whole job.
- `partial[blockIdx.x]` is not supposed to hold the total. It holds this block's share.
- The final cross-block reduction is deliberately out of scope. Writing a second kernel or your own `main()` comes later.

Now the optimization. The starter is correct and serial: one thread per block adds blockDim elements one at a time while 31 of every 32 lanes do nothing.

`__shfl_down_sync(mask, value, offset)` returns the `value` register belonging to the lane `offset` above the caller. No memory is involved, in either direction. So a warp folds itself by halves, offsets 16, 8, 4, 2, 1, and turns 32 values into one in five instructions with no shared memory and no `__syncthreads()` inside the warp.

A 256-thread block is eight warps, so after that fold you have eight values living in eight different warps' registers, and registers are private to a warp. **That** is what shared memory is for here: lane 0 of each warp writes its total to a small shared array, one `__syncthreads()`, and warp 0 folds those eight the same way. Shared memory between warps, registers within one.

A shuffle addresses a fixed lane, so it assumes a full warp. Keep the block size a multiple of 32.

There are five shuffle rounds for a 32-lane warp, not five total machine instructions. Every lane named by the participation mask must execute the matching collective.

What you will learn

• State the partial-output contract: one value per block, combined by the host, not by your kernel.
• Use __shfl_down_sync to reduce a warp's 32 lanes in registers, without shared memory or __syncthreads.
• Explain why the per-warp totals still need shared memory: registers are private to a warp.
• Combine per-warp partials into one block result with a small shared array and a final warp reduction.
• Understand that warp-shuffle assumes a full warp, so the block size should be a multiple of 32.

Starter code

    // Serial and wasteful: only thread 0 works while the rest of the block idles.
    int blockStart = blockIdx.x * blockDim.x;
    if (threadIdx.x == 0) {
        float sum = 0.0f;
        for (int i = 0; i < blockDim.x; ++i) {
            int idx = blockStart + i;
            if (idx < N) sum += X[idx];
        }
        partial[blockIdx.x] = sum;
    }