12.1 概述

因为该二元操作符符合结合律，O（N）个用于计算归约结果的操作可以以任意顺序执行。

图12-1展示了一些处理8元素数组的不同方式。为了方便对比，我们给出它的串行实现。只需要具备一个可以执行 $\oplus$ 操作符的执行单元，但是这种方法的性能比较差，因为它完成计算需要7步。

图12-1 8元素的归约

原书这里的交替方式的公式写错了，应该改成（ $(a_0\oplus a_4)$ ） $\oplus (a_{1}\oplus a_{5})$ ） $\oplus ((a_{2}\oplus a_{6})\oplus (a_{3}\oplus a_{7}))$ 。——译者注

成对的方式是直观的，并且只需要0（1gN）步（这里是3步）来计算结果，但它在CUDA中性能较差。当读取全局内存时，让单个线程访问相邻的内存单元会导致非合并的内存事务。当读取共享内存时，所示的模式会引起存储片冲突（bank conflict）。

不论是对全局内存还是共享内存，基于交替策略效果更好。在图12-1中，交替因子为4。对于全局内存，使用blockDim.x $\times$ gridDim.x的倍数作为交替因子有良好的性能，因为所有的内存事务将被合并。对于共享内存，最好的性能是按照所确定的交错因子来累计部分和，以避免存储片冲突，并保持线程块的相邻线程处于活跃状态。

一旦一个线程块处理完其交替子数组，将结果写入全局内存以备它随后启动的内核进一步处理。启动多个内核可能看起来开销很大，但内核启动是异步的，因此当GPU正在执行第一个内核时，CPU可以请求下一个内核启动。每个内核启动都有机会来指定不同的启动配置。

鉴于内核的性能可随不同的线程和线程块大小而有所不同，因此编写可以正常工作于任何组合形式的线程和线程块大小的内核，是一个好主意。最佳的线程/线程块配置，可以根据经验来确定。

本章的初始归约内核展示一些大家可能熟悉的重要CUDA编程概念：

·合并的内存操作，以最大限度地提高带宽；
·可变大小的共享内存，以方便线程之间的协作；
·避免共享内存的存储片冲突。

优化的归约内核展示更先进的CUDA编程惯例：

·线程束同步代码避免了不必要的线程同步；

• 原子操作和内存栅栏避免了调用多个内核的需要;
  ·洗牌指令支持线程束级别的归约，而无须使用共享内存。