Cuda归约运算（一）——分析cudasample的reduction1-优快云博客

本文链接：https://blog.youkuaiyun.com/qq_32417133/article/details/87923493

何谓归约

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

概述

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

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

串行实现
$({a_0}\bigoplus({a_1}\bigoplus({a_2}\bigoplus({a_3}\bigoplus({a_4}\bigoplus({a_5}\bigoplus({a_6}\bigoplus({a_7})))))))$
对数步长的归约
$((({a_0}\bigoplus{a_1}) \bigoplus({a_2}\bigoplus{a_3}))\bigoplus(({a_4}\bigoplus{a_5})\bigoplus({a_6}\bigoplus{a_7})))$
对数步长的归约
$((({a_0}\bigoplus{a_1}) \bigoplus({a_1}\bigoplus{a_5}))\bigoplus(({a_2}\bigoplus{a_6})\bigoplus({a_3}\bigoplus{a_7})))$

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