The warp is the natural place to begin, because its 32 threads can exchange values directly through shuffle instructions (Section 8.6.1) without touching shared memory. The cooperative groups library expresses a warp reduction in one call:
namespace cg = cooperative_groups;
cg::thread_block_tile<32> warp =
cg::tiled_partition<32>( cg::this_thread_block() );
int sum = cg::reduce( warp, myValue, cg::plus<int>() );After the call, every lane in the warp holds the sum of the warp’s 32
input values. The compiler lowers cg::reduce() to a
sequence of shuffle instructions – the same sequence one would otherwise
write by hand. Using the “butterfly” variant
__shfl_xor_sync(), the five steps are:
int mySum = myValue;
mySum += __shfl_xor_sync( 0xffffffff, mySum, 16 );
mySum += __shfl_xor_sync( 0xffffffff, mySum, 8 );
mySum += __shfl_xor_sync( 0xffffffff, mySum, 4 );
mySum += __shfl_xor_sync( 0xffffffff, mySum, 2 );
mySum += __shfl_xor_sync( 0xffffffff, mySum, 1 );The first argument to __shfl_xor_sync() is a mask of the
participating lanes; 0xffffffff names all 32. Figure 12-2
shows how the value propagates: each thread’s partial sum is drawn as a
4W×8H rectangle, with a dark square marking which threads have
contributed to it. With each step, the number of contributions doubles,
until after lg(32) = 5 steps every thread holds the full reduction.
Figure 12-2. Reduction using the shuffle instruction.
Because the butterfly makes the result available to all 32 lanes, no
separate broadcast is needed. The shuffle-up and shuffle-down variants
take the same number of steps but leave the result in only one lane. For
floating-point reductions, note that the butterfly can leave different
lanes with slightly different values, since they sum the contributions
in different orders; if a single canonical value is required, one lane’s
result can be broadcast with __shfl_sync().