The CUDA Handbook ▸ Chapter 12. Reduction

12.2 Block Reduction

To reduce a thread block, each warp reduces its own 32 values, writes that partial to shared memory, and then the first warp reduces the per-warp partials. Listing 12-1 gives blockReduceCG(), which does exactly this with two calls to cg::reduce() and one array of shared memory holding a single partial per warp.

//// Reduce v across the whole thread block. The result is returned to// thread 0; other threads' return values are undefined. sPartials must// have room for one T per warp (blockDim.x/32 elements).//template<class T>__device__ TblockReduceCG( cg::thread_block block,               cg::thread_block_tile<32> warp,               T v, T *sPartials ){    v = cg::reduce( warp, v, cg::plus<T>() );   // one value per warp    if ( warp.thread_rank() == 0 )        sPartials[warp.meta_group_rank()] = v;    block.sync();     if ( warp.meta_group_rank() == 0 ) {        v = ( warp.thread_rank() < warp.meta_group_size() ) ?                sPartials[warp.thread_rank()] : (T) 0;        v = cg::reduce( warp, v, cg::plus<T>() );    }    return v;}

Because cg::reduce() works on any number of lanes and the cross-warp step needs only one slot per warp, this formulation places no constraint on the block size – it need not be a power of 2 – and contains no hand-written log-step loop.

The classic formulation, which predates the shuffle instruction, instead stages every thread’s value through shared memory and performs the log-step reduction there, halving the number of active threads at each step. That approach still works, but it moves more data through shared memory and, done naively, incurs the bank conflicts described earlier; avoiding them takes the interleaved addressing that the shuffle-based reduction gets for free. Section 7.5 covers the “warp-synchronous” optimization that the shared memory reduction once used for its final 32 elements – a technique NVIDIA has since deprecated – along with its correct replacement, the shuffle sequence of Section 12.1.