cuda编程_CUDA编程入门（五）更高效的并行归约算法

上一篇我们写了一个简单的并行归约算法，虽然相比与CPU已经达到了较快的速度，但其实程序还有很大的优化空间。

比较明显的是有两个需要优化的点：

• 线程束分化
• 内存访问

下面我们来一一解决。

线程束分化

首先我们来回忆一下上个程序的实现思路，我们是利用stride变量来实现每轮迭代时的被加数选择，这造成线程束的分化，也就是每个线程束中只有部分线程是活跃的，但由于硬件设计，调度会以一整个线程束为单位进行，所以影响了程序的效率。

这种情况下，我们可以通过重新组织线程索引来解决线程束分化的问题。我们把核函数的代码修改成这样：

__global__ 

相比于之前的初版代码，我们现在使用线程ID来生成一个数组访问索引，这种方式有效地避免了线程束的分化。可能会有点迷惑，我们更具体说说，在每个线程块有1024个线程（32个线程束）时，在第一轮迭代，前16个线程束执行计算，后16个线程束什么都不做；而原来的代码中，32个线程束都执行计算，但只有偶数标号的线程活跃。第二轮迭代时，前8个线程束执行计算，后24个线程束执行计算；而原来的代码中，32个线程束都执行计算，但只有标号是4的倍数的线程活跃。这样重新组织线程ID后，线程束分化就被避免掉了。我们来实际运行一下，看看效果。

终端截图

运行时间从2.5ms变成了1.5ms，加速了1.6倍，非常可观，可见尽量避免线程束的分化十分重要。这给了我们一点的启发，看似不起眼的细节，在CUDA编程中却会产生不小的影响，这也是我们需要了解底层硬件运行机制的一个重要原因。

内存组织

在CPU编程中，我们学过空间局部性与时间局部性，这在CUDA编程中对我们也是很有启发的。之前我们采用相邻配对进行求和，这种方法最为直观，但会造成在第二轮之后内存访问的不连续（因为使用了stride变量作为跨度）。为了缓解这种现象，我们重新组织一下配对方法，让对内存的访问更加集中，如下图，这种新的配对方法我们称为交错配对。

重新组织后的配对方法

新编写的核函数如下：

__global__ 

我们来看看运行结果：

终端截图

在改进线程束分化后，我们重新组织了内存的访问方式，运行时间从1.54ms提升到了1.36ms，获得了大约1.13倍的加速。

这给我们的启发是，对全局内存的访问要尽量进行合并访问与存储，这样才能达到最大的带宽。

完整代码放在这里：

ZihaoZhao/CUDA_study

小结一下，这一篇我们使用两种思路对程序进行了简单的优化，程序从2.48ms加速到1.36ms，累计获得了1.8倍的加速比。而这两点，恰恰对应了系列开篇中我们提到的CUDA编程的特点，线程组织和内存组织，这也是我们在CUDA编程中需要随时注意的。