The CUDA Handbook ▸ Chapter 12. Reduction

12.3 Grid Reduction

Reducing an array in device memory means combining the results of many thread blocks. CUDA provides no general barrier across the blocks of a grid, so a grid-level reduction must either launch a second kernel or coordinate through global memory.

12.3.1 Two-Pass Reduction

The two-pass approach invokes the same kernel twice. On the first pass, each block reduces an interleaved slice of the input with a grid-stride loop and writes its block-level partial to an intermediate array in global memory. On the second pass, a single block reduces that array of partials to the final result. Listing 12-2 gives the kernel and its host function.

__global__ voidReductionCG_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;    for ( size_t i = 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 )        out[blockIdx.x] = sum;} voidReductionCG( int *answer, int *partial,             const int *in, size_t N,             int numBlocks, int numThreads ){    int sharedBytes = (numThreads/32) * sizeof(int);    ReductionCG_kernel<<< numBlocks, numThreads, sharedBytes >>>(        partial, in, N );    ReductionCG_kernel<<< 1, numThreads, sharedBytes >>>(        answer, partial, numBlocks );}

Each thread walks the input with a grid-stride loop, reading elements blockDim.x*gridDim.x apart, so the accesses are interleaved and every thread accumulates a partial before the block reduction begins. The two launches are inexpensive: kernel launches are asynchronous, so the CPU can request the second while the GPU runs the first, and each launch can carry its own configuration.

12.3.2 Single-Pass Reduction with Atomics

When ⊕ is supported by a hardware atomic, the second pass can be eliminated. Each block reduces its slice as before, then adds its block-level partial to the output with a single atomicAdd(). Listing 12-3 gives this kernel, which needs only one launch. The output location must be initialized to 0 before launch, because the kernel cannot safely clear it – there is no way to resolve the race between blocks from within the kernel.

__global__ voidReduction5_kernel( int *out, const int *in, size_t N ){    const int tid = threadIdx.x;    int partialSum = 0;    for ( size_t i = blockIdx.x*blockDim.x + tid;          i < N;          i += blockDim.x*gridDim.x ) {        partialSum += in[i];    }    atomicAdd( out, partialSum );} voidReduction5( int *answer, int *partial,             const int *in, size_t N,             int numBlocks, int numThreads ){    cudaError_t status_cudart;    cuda(Memset( answer, 0, sizeof(int) ));    Reduction5_kernel<<< numBlocks, numThreads>>>( answer, in, N );Error_cudart:;}

An older single-pass formulation, in the threadFenceReduction SDK sample, avoids the atomic on the output by having the last block to finish perform the final reduction; it detects the last block with an atomic counter and orders the global memory writes with __threadfence(). The atomic-accumulate version in Listing 12-3 is simpler and, on current hardware, at least as fast.