Cuda归约运算（一）——分析cudasample的reduction1

何谓归约

归约（redution）是一类并行算法，对传入的O(N)个输入数据，使用一个二元的复合结合律的操作符，生成O(1)的结果。这类操作包括取最小、最大、求和、平方求和、逻辑与、逻辑或、向量点积。归约也是其他高级运算中要用的基础算法。
除非操作符的求解代价极高，否则归约倾向于带宽受限型任务。下面就从SDK提供的reduction例子入手，详细理解该归约算法。

概述

因为该二元操作符复合结合律，O(N)个用于计算归约结果的操作可以以任意顺序执行。
$$\sum _i^8{a}_i=a_0⨁a_1⨁a_2⨁a_3⨁a_4⨁a_5⨁a_6⨁a_7$$

下面展示了一些处理8元素数组的不同方式。为了方便对比。下面先给出它的串型实现。只需要具备一个可以执行的$⨁$操作符的执行单元，但是这种方法的性能就比较差了，因为完成它需要7步运算。

• 串行实现
  $$(a_0⨁(a_1⨁(a_2⨁(a_3⨁(a_4⨁(a_5⨁(a_6⨁(a_7)))))))$$

• 对数步长的归约
  $$(((a_0⨁a_1)⨁(a_2⨁a_3))⨁((a_4⨁a_5)⨁(a_6⨁a_7)))$$

• 对数步长的归约
  $$(((a_0⨁a_1)⨁(a_1⨁a_5))⨁((a_2⨁a_6)⨁(a_3⨁a_7)))$$

成对的方式是最直观的，并且只需要*O(lgN)*步来计算结果，但它在cuda中的性能较差。当读取全局内存时，让单个线程访问相邻的内存单元会导致非合并的内存事物。当读取共享内存时，所示的模式会引起存储片的冲突。
不论对全局内存还是共享内存，基于交替的策略效果更好。对于全局内存，使用blockDim.x * gridDim.x的倍数作为交替因子有良好的性能，因为所有的内存事物将被合并，最好的性能是按照所确定的交替因子来累计部分和，以避免存储片冲突，并保持线程块的相邻线程处于活跃状态。
一旦一个线程块处理完其交替子数组，将结果写入全局内存以备它随后启动的内核进一步处理。启动多个内核可能看起来开销很大，但内核启动时异步的，因此当GPU正在执行一个第一个内核时，CPU可以请求下一个内核启动，每个内核启动都有机会来指定不同的启动配置。