12.4 使用原子操作的归约

对于那些为硬件的本地原子操作符支持的操作符 $\oplus$ 的归约，编写更简单的归约算法是可能的：只要循环遍历输入数据并采取“用后即弃”的方式把输入添加到结果所在的内存位置来接收输出值。代码清单12-7中给出的Reduction5_kernel比以前的方法要简单得多。每个线程计算输入的部分和，并在尾部对输出执行一个atomicAdd。

需要注意的是，Reduction5_kernel不能正常工作，除非out指向的存储位置初始化为0。[1]像threadFenceReduction示例一样，这个内核的优点是只需要一个内核调用就可完成归约操作。

代码清单12-7 使用全局内存原子操作的归约

(reduction5Atoms.cuh)

global void Reduction5_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  $<  \overline{\mathbf{N}}$  . i  $+ =$  blockDim.x\*gridDim.x）{ partialSum  $+ =$  in[i]; } atomicAdd( out, partialSum);   
}   
void Reduction5( int \*answer, int \*partial, const int \*in, size_t N, int numBlocks, int numThreads)   
{ CUDAMemset( answer, 0, sizeof(int)); Reduction5_kernel<< numBlocks, numThreads>>>(answer, in, N);

[1] 内核本身不能执行这个初始化，因为CUDA执行模型不支持解决不同线程块之间资源竞争的机制。请参见7.3.1节。