这一篇我们一起学习一下如何使用CUDA实现并行归约算法。

首先我们要知道什么是并行归约。并行归约（Reduction）是一种很基础的并行算法，简单来说，我们有N个输入数据，使用一个符合结合律的二元操作符作用其上，最终生成1个结果。这个二元操作符可以是求和、取最大、取最小、平方、逻辑与或等等。

我们以求和为例，假设输入如下：

int array[8] =  [3, 1, 7, 0, 4, 1, 6, 3]

在串行的情况下，算法很容易实现，一般我们会使用下面这样的代码。

int 

但当我们试图进行并行计算时，问题就会变得复杂。

由于加法的交换律和结合律，数组可以以任意顺序求和。所以我们会自然而然产生这样的思路：首先把输入数组划分为更小的数据块，之后用一个线程计算一个数据块的部分和，最后把所有部分和再求和得出最终结果。

依照以上的思路，我们可以按下图这样计算。就上面的输入例子而言，首先需要我们开辟一个8个int的存储空间，图中一行的数据代表我们开辟的存储空间。计算时首先将相邻的两数相加（也称相邻配对），结果写入第一个数的存储空间内。第二轮迭代时我们再将第一次的结果两两相加得出下一级结果，一直重复这个过程最后直到我们得到最终的结果，空白方框里面存储的内容是我们不需要的。这个过程相当于，每一轮迭代后，选取被加数的跨度翻倍，后面我们核函数就是这样实现的。

相比与串行计算，我们只用了3轮迭代就得出了8个数的和，时间复杂度由O(N)变为O(logN)。

当然使用归约算法时我们不会只有这么小的输入数组，这时候就要用到线程块了，我们可以把输入数组先划分成很多包含8个int型值的数组，每一个小数组分配给一个线程块，最后再将所有的结果传回主机串行求和。

主函数如下，与我们上一个例程结构类似，先初始化，分配内存，然后运行核函数，最后和CPU对照组对比，检验结果是否正确。我们重点关注分配线程网格和线程块的主干部分。与上面例子里讲的8元素数组不同，实际使用时为了达到高加速比，我们会把输入数组分成更大的块。

这里我们使用一个16M（16777216）大小的输入数组，分成的小数组大小为1K（1024）。所以程序里将block设置为（1024, 1）的大小，每个线程块完成一个小数组的求和。一共需要16K（16384）个线程块，所以我们设置grid为（16384, 1）。每个线程块的求和结果都保存在全局内存里，等所有线程完成后，统一传回主机，在主机里串行求和得到最后的结果。注意这里的block和grid设置并不是最优，只是为了简单，下一篇中我们会进行优化。

int 

核函数具体编程实现如下。我们利用一个stride变量，实现不同轮数时被加数的选择。每当计算一轮后，选取被加数的跨度会乘2。在这里，有心的同学就会发现了，第一轮计算中，其实只有一半的线程是活跃的，而且每进行一轮计算后，活跃的线程数都会减少一半，这是条件表达式的使用造成的。比如在第一轮迭代时，只有偶数ID的线程会为True，其主体才能得到执行。这会导致线程束的分化，也就是说只有一部分线程是活跃的，但是所有的线程仍然都会被调度，因为硬件调度线程是以线程束（连续32个线程）为单位进行调度的。这肯定会影响我们程序的执行效率，不过不用太担心，我们下一篇就会提出解决这个问题的方法。

__global__ 

最后我们看看运行结果，cpu花费89.3ms，而gpu花费2.5ms，运算结果一致。

完整代码和编译生成的可执行文件放在这里：

https://github.com/ZihaoZhao/CUDA_study/tree/master/Reduction

总结一下，这一篇我们使用CUDA完成了一个基础的并行归约算法，不过这算法还有很大的优化空间，接下来我们一起对这个算法进行优化，看看能再加快多少。