Block Reduction: Scaling Reductions with Shared Memory and Warp Handoff

In warp reduction, we saw how 32 threads can collapse their values down to a single number in just five steps using register shuffles. Because threads in a warp execute synchronously, they pass values directly without writing to memory or waiting on barriers.

However, CUDA programs rarely launch only 32 threads. Real-world kernels launch thread blocks of 128, 256, 512, or 1024 threads to achieve high occupancy and saturate the GPU's compute units.

Because warp shuffle instructions cannot reach across warp boundaries, threads in Warp 1 cannot directly read registers from Warp 0. To reduce an entire block, we need a two-tier strategy:

1. Shared memory tree reduction: Pool data into on-chip SRAM and repeatedly halve the active participants until only 32 values remain.
2. Warp-level handoff: Let the first warp pick up the final 32 values and sprint across the finish line with register shuffles.

---

The Mental Model: An Auditorium and a Council Table

Imagine an auditorium filled with 256 researchers, each holding an observation score. You need to find the overall maximum score.

Instead of all 256 people shouting across the room or queuing up at a single microphone:

1. Write on the whiteboard: Everyone walks up to a large whiteboard at the front of the room (shared memory) and writes their score in their assigned box (Index 0 to 255). Everyone waits until the whiteboard is completely filled (__syncthreads()).
2. First round of cuts (Halving to 128): Researchers 0 to 127 walk up to the board, compare their score with researcher (i + 128), and update their box with the higher number. Half the scores are eliminated in one pass.
3. Second round of cuts (Halving to 64): Researchers 0 to 63 compare with researcher (i + 64).
4. Third round of cuts (Halving to 32): Researchers 0 to 31 compare with researcher (i + 32).
5. The council handoff: Now only 32 scores remain, sitting in boxes 0 to 31. The remaining 224 researchers sit down. The first 32 researchers (Warp 0) do not need the whiteboard or slow barriers anymore. They sit in a circle and play the 5-step note-passing game (warp_reduce).
6. Publishing the winner: Researcher 0 writes the final champion score into box 0. Everyone is told the game is finished (__syncthreads()), and every researcher in the room can read box 0 as they leave.

---

The Implementation

Here is the complete implementation of block_reduce:

template <typename Op>
__device__ float block_reduce(float val, float *shared, int tid, Op op) {
    // 1. Stage each thread's local value into shared memory
    shared[tid] = val;
    __syncthreads();

    // 2. Tree reduction down to 32 elements
    for (int stride = blockDim.x / 2; stride >= 32; stride >>= 1) {
        if (tid < stride) {
            shared[tid] = op(shared[tid], shared[tid + stride]);
        }
        __syncthreads();
    }

    // 3. Final 32 elements reduced by Warp 0
    if (tid < 32) {
        val = warp_reduce(shared[tid], op);
        if (tid == 0)
            shared[0] = val;
    }

    // 4. Synchronize so all threads can read the broadcast result
    __syncthreads();
    return shared[0];
}

---

Why Stop at Stride 32?

A common beginner mistake when writing CUDA tree reductions is running the shared memory loop all the way down to stride = 1:

// Slower approach: running shared memory loop to the very end
for (int stride = blockDim.x / 2; stride > 0; stride >>= 1) {
    if (tid < stride) {
        shared[tid] = op(shared[tid], shared[tid + stride]);
    }
    __syncthreads();
}

While correct, this approach is noticeably slower. Every single iteration requires a block-wide barrier (__syncthreads()). Barriers force every warp in the block to wait, stalling execution units and flushing pipeline stages.

By stopping the loop when stride >= 32, we eliminate five expensive block-wide synchronization barriers (strides 16, 8, 4, 2, 1).

Once the data is narrowed down to 32 elements, all 32 values reside in shared[0] through shared[31]. That maps perfectly to the 32 threads of Warp 0 (tid < 32). Warp 0 loads those 32 values into registers and reduces them using warp_reduce at full register-to-register speed without a single barrier.

---

Step-by-Step Trace for a 256-Thread Block

Let us trace block_reduce for a block with blockDim.x = 256:

Phase	Active Threads	Action	Synchronization
Initial Staging	256 threads	shared[tid] = val	__syncthreads()
Stride 128	Threads 0 to 127	shared[tid] = op(shared[tid], shared[tid + 128])	__syncthreads()
Stride 64	Threads 0 to 63	shared[tid] = op(shared[tid], shared[tid + 64])	__syncthreads()
Stride 32	Threads 0 to 31	shared[tid] = op(shared[tid], shared[tid + 32])	Loop exits
Warp Handoff	Threads 0 to 31	val = warp_reduce(shared[tid], op)	Warp shuffle (no barriers)
Store Champion	Thread 0	shared[0] = val	-
Final Broadcast	All 256 threads	Every thread returns shared[0]	__syncthreads()

Only 3 block-wide barriers are spent in the tree loop (for strides 128, 64, and 32), plus one initial staging barrier and one final broadcast barrier.

---

Memory Bank Conflicts: Contiguous vs Interleaved Addressing

How you index into shared memory during a tree reduction significantly impacts memory bandwidth.

Notice that our loop uses contiguous half-block indexing:

// Good: Contiguous addressing (zero bank conflicts)
if (tid < stride) {
    shared[tid] = op(shared[tid], shared[tid + stride]);
}

Compare this to the older interleaved pattern often found in legacy tutorials:

// Bad: Interleaved addressing (causes severe bank conflicts)
for (int s = 1; s < blockDim.x; s *= 2) {
    if (tid % (2 * s) == 0) {
        shared[tid] = op(shared[tid], shared[tid + s]);
    }
    __syncthreads();
}

The interleaved pattern causes two major performance problems:

1. Thread Divergence: In the first step, only even threads work (tid % 2 == 0). In the next step, only every fourth thread works (tid % 4 == 0). Warps become severely diverged.
2. Bank Conflicts: Because threads read at powers-of-two multiples, requests map to the same shared memory banks, causing multi-way serialization.

The contiguous stride pattern keeps active threads bundled in low-numbered warps (tid < stride), ensuring full warp utilization and zero bank conflicts as threads access sequential memory locations.

---

Summary

Block reduction combines the strengths of both GPU memory layers:

• Shared memory provides high-capacity, on-chip communication to gather data across hundreds of threads in different warps.
• Warp shuffles take over in the final lap to finish the reduction with five single-cycle register operations, skipping five unnecessary synchronization barriers.