The CUDA Handbook ▸ Chapter 12. Reduction

12.4 Optimizing for Bandwidth

Reduction reads O(N) input and writes an O(1) result, so it is bound by memory bandwidth: the best a reduction can do is read the input as fast as the memory system allows. The kernels so far reach about 316 GB/s on an RTX 3060, roughly 88% of its ~360 GB/s of peak bandwidth. Two changes close most of the gap.

The first is to load the input with vectorized 128-bit transactions rather than one int at a time. Each thread reads an int4 – four ints at once – so the load path issues a quarter as many instructions to cover the same bytes and keeps more memory transactions in flight. The second is to finish in a single pass with an atomic, avoiding the intermediate array. Listing 12-4 gives the resulting kernel, which handles any input length by processing the bulk of the array as int4 and the last few elements individually.

__global__ voidReductionVector_kernel( int *out, const int *in, size_t N ){    cg::thread_block block = cg::this_thread_block();    cg::thread_block_tile<32> warp = cg::tiled_partition<32>( block );    extern __shared__ int sPartials[];     int sum = 0;     // Body: consume the input four ints at a time with 128-bit loads.    const size_t N4 = N / 4;    const int4 *in4 = reinterpret_cast<const int4 *>( in );    for ( size_t i = blockIdx.x*blockDim.x + threadIdx.x;          i < N4;          i += blockDim.x*gridDim.x ) {        int4 v = in4[i];        sum += v.x + v.y + v.z + v.w;    }    // Tail: the up-to-three elements past the last full int4.    for ( size_t i = 4*N4 + blockIdx.x*blockDim.x + threadIdx.x;          i < N;          i += blockDim.x*gridDim.x ) {        sum += in[i];    }     sum = blockReduceCG<int>( block, warp, sum, sPartials );    if ( threadIdx.x == 0 )        atomicAdd( out, sum );} voidReductionVector( int *answer, int *partial,                 const int *in, size_t N,                 int numBlocks, int numThreads ){    int sharedBytes = (numThreads/32) * sizeof(int);    cudaMemset( answer, 0, sizeof(int) );    ReductionVector_kernel<<< numBlocks, numThreads, sharedBytes >>>(        answer, in, N );}

Table 12-1 compares these reductions against cub::DeviceReduce and thrust::reduce on an RTX 3060, summing 64M 32-bit integers. The cg::reduce-based block reduction matches the hand-written shuffle reduction while being markedly simpler, and the vectorized single-pass kernel reaches the same throughput as CUB and Thrust.

Table 12-1. Reduction throughput (RTX 3060, 64M 32-bit integers).

Reaching library-level throughput takes only vectorized loads and a single-pass atomic on top of the cg::reduce block reduction.