23、CUDA 中的归约与扫描算法详解

CUDA 中的归约与扫描算法详解

1. 单遍归约

1.1 原理

传统的两遍归约方法，是由于 CUDA 块之间无法相互同步而采用的一种变通方案。在没有块间同步机制来确定何时可以开始处理最终输出时，就需要第二次内核调用。而通过结合原子操作和共享内存，可以避免第二次内核调用。

1.2 代码实现

1.2.1 Reduction4_LogStepShared 函数

template<unsigned int numThreads>
__device__ void
Reduction4_LogStepShared( int *out, volatile int *partials )
{
    const int tid = threadIdx.x;
    if (numThreads >= 1024) {
        if (tid < 512) { 
            partials[tid] += partials[tid + 512];
        }
        __syncthreads();
    }
    if (numThreads >= 512) {
        if (tid < 256) { 
            partials[tid] += partials[tid + 256];
        }
        __syncthreads();
    }
    if (numThreads >= 256) {
        if (tid < 128) {