cuda c权威编程指南_CUDA编程入门（六）展开循环继续优化

这一节我们还是要继续优化并行归约的程序。（惊不惊喜意不意外.jpg）

不过不用担心，这一篇我们会从另外几个角度来进行优化，与上一篇并不重复。而且反复优化程序对我们理解CUDA编程是很有帮助的，还请大家坚持看下去。

展开循环以隐藏延时

现在使用高级语言编程时，我们已经不会刻意去进行循环展开了，因为这件事会由编译器帮我们完成。但在CUDA编程中，循环展开具有很重要的意义，它能给线程束调度器提供更多可用的线程束，以帮助我们有效地隐藏延时。

于是我们可以使用循环展开的方法，对并行归约的程序再来一波优化。

我们之前只是用一个线程块来处理一个小数组，我们称其为一个数据块。如果我们使用一个线程块手动展开两个数据块的处理，那会怎样呢？先给个结论，这样通过减少指令消耗和增加更多的独立调度指令，更多的并发操作被添加到流水线上，以产生了更高的指令和内存带宽。反映在宏观上就是程序执行的总体时间变少了。

改进后的核函数代码在这里：

__global__ 

从概念上来讲，可以把它作为归约循环的一个迭代，此循环可在数据块间进行归约。

如果每个线程处理两个数据块，那么我们需要的线程块总量会变为原来的一半，因此主函数也要对应修改。

看上去，这样处理后线程块减少了，与我们之前要使用尽量多线程块的理论不符。但实际我们通过这种方式，让一个线程中有更多的独立内存加载/存储操作，这样可以更好地隐藏内存延时，更好地使用设备内存读取吞吐量的指标，以产生更好的性能。所以我们编程时，各个策略要针对实际情况结合使用。

执行结果如下图。

展开2块

相比于上篇最后的优化结果，程序执行速度再次获得了不小的提升，从1.37ms提升到0.8ms，1.7倍加速比。

既然一个线程块处理2个数据块能获得这么高的加速比，那么处理4个，8个呢？代码不赘述了，我们直接来看结果。

展开8块

随着处理数据块数量的增多，处理时间不断降低。不过随着设备内存吞吐量逐渐到达极限，这个时间就不会继续降低了。

总结一下，为了隐藏延时，我们需要合理地增加一个线程块中需要处理的数据量，以便线程束调度器进行调度。

展开线程

我们可以再进一步想想，当执行到最后几次迭代时，当只需要32个或更少线程时，每次迭代后还需要进行线程束同步。为了加速，我们可以把这最后6次迭代进行展开，使用下面的语句：

if

这样就避免了执行控制循环逻辑和线程同步逻辑的时间。

把这段添加到上面的代码中，然后修改一下之前的迭代次数，运行得到如下结果，可以看到又加速了不少。

展开线程

完全展开

到目前位置，我们的代码里只剩一个循环了。不过理论上如果编译时已知一个循环中的迭代次数，就可以把循环完全展开。因为我们的核函数中的循环迭代次数是基于一个线程块维度的，而且一个线程块支持最大1024个线程，所以我们可以将这个循环进行完全展开。核函数代码如下：

__global__ 

运行一下看看结果：

完全展开

不过结果令人意外，完全展开后反而时间增加了，这和书上的结果不太一样，看来nvcc编译器后来又做了很多新的优化。

优化成果小结

画个表格总结一下从最初版本到现在我们的加速情况。

优化效果总结

优化效果还是很明显的，最终我们在初版代码的基础上获得了8.21倍的加速。再次感谢《CUDA C编程权威指南》的作者能为我们提供这么精彩的例子。

最后，本篇的代码放在这里：

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

看到这里，不用说大家应该也有了一定体会，对于并行编程而言，优化是无止境的。可以说，不考虑编程时间的优化都是耍流氓。在优化中我们往往会发现有些改变是“唾手可得”的，这些改变会很容易，而且带来可观的加速。当这些被改进掉后，就必须采用更复杂的策略来进一步优化，这会花费我们更多的时间，代码可读性会变差，也会引入更多潜在的bug。总体来说，如果以我们的编程时间为x轴，加速比为y轴，我们能获得的往往是一个对数曲线，这一点是我们做实际项目时需要铭记的。

长呼一口气，我们终于把并行归约的程序地优化地差不多了，能看到这里的你真的不容易，来评论留个言，我们交个朋友。顺便感慨一下，学并行编程算法的人比学深度学习算法的人少多了，我觉得是好事，至少证明这里没那么大泡沫，大家共勉。