Tesla架构下的CUDA程序优化

CUDA优化的最终目的是：在最短的时间内，在允许的误差范围内完成给定的计算任务。在这里，“最短的时间”是指整个程序运行的时间，更侧重于计算的吞吐量，而不是单个数据的延迟。在开始考虑使用GPU和CPU协同计算之前，应该先粗略的评估使用CUDA是否能达到预想的效果，包括以下几个方面：

精度：目前GPU的单精度性能要远远超过双精度性能，整数乘法、求模、求余等运算的指令吞吐量也较为有限。在科学计算中，由于需要处理的数据量巨大，往往采用双精度或者四精度才能获得可靠的结果，目前的Tesla架构还不能很好的满足高精度计算的需要。如果你的计算需要很高的精度，或者需要进行很多轮的迭代，最好考虑在关键的步骤中使用双精度，而在其他部分仍然使用单精度浮点以获得指令吞吐量和精度的平衡。而如果你对精度有更高的要求，那么现在的架构还不能太高的加速比。不过，在2010年将会普及的下一代架构中，双精度浮点和整数处理能力将有很大的提升，这种情况会有根本性的改变。

延迟：目前CUDA还不能单独为某个处理核心分配任务，因此必须先缓冲一定量的数据，再交给GPU进行计算。这样的方式可以获得很高的数据吞吐量，不过单个数据经过缓冲、传输到GPU计算、再拷贝回内存的延迟就比直接由CPU进行串行处理要长很多。如果对应用实时性要求很高，比如必须在数毫秒内完成对一个输入的处理，那么使用CUDA可能会影响系统的整体性能。对于要求人机能够实时交互的系统，应该将延迟控制在数十毫秒，以响应用户的输入。通过减小缓冲，可以减小延迟，但至少要保证每个内核程序处理的一批数据能够让GPU满负荷工作。不过在大多数情况下，在计算吞吐量较大，需要由GPU才能实时实现的系统，投入相同成本使用CPU很难做到接近实时。如果确实对实时性和吞吐量都有很高要求，应该考虑ASIC、FPGA或者DSP实现，这需要更多的投入，更长的开发时间和硬件开发经验。

计算量：如果计算量太小，那么使用CUDA是不合算的。衡量计算量有绝对和相对两种方式。

从绝对量来说，如果要优化的程序使用频率比较低，并且每次调用需要的时间也可以接受，那么使用CUDA优化并不会显著改善使用体验。对于一些计算量非常小（整个程序在CPU上可以在几十毫秒内完成）的应用来说，使用CUDA计算时在GPU上的执行时间无法隐藏访存和数据传输的延迟，此时整个应用程序需要的时间反而会比CPU更长。此外，虽然GPU的单精度浮点处理能力和显存带宽都远远超过了CPU，但由于GPU使用PCI-E总线与主机连接，因此它的输入和输出的吞吐量受到了IO带宽的限制。当要计算的问题的计算密集度很低时，执行计算的时间远远比IO花费的时间短，那么整个程序的瓶颈就会出现在PCI-E带宽上。此时无论如何提高浮点处理能力和显存带宽，都无法提高系统性能。

相对的计算量指可以并行的部分在整个应用中所占的时间。如果整个应用中串行部分占用时间较长，而并行部分较短，那么也需要考虑是否值得使用GPU进行并行计算。例如，假设一个程序总的执行时间为1.0，其中串行部分占0.8，而并行部分只占0.2，那么使用GPU将并行部分加速10倍，总的执行时间也只能从1.0降低到0.82。即使是在CPU和GPU可以同时并行计算的应用中，执行时间也至少是CPU串行计算需要的0.8。只有在并行计算占用了绝大多数计算时间的应用中，使用CUDA加速才能获得很高的加速比。不过，随着GPU+CPU并行计算的普及和GPU架构的进一步改进，即使获得的加速比较小，也可能会由GPU执行。

完成对GPU加速能够达到的效果的粗略评估后，就可以开始着手编写程序了。为了在最短的时间内完成计算，需要考虑算法、并行划分、指令流吞吐量、存储器带宽等多方面问题。总的来说，优秀的CUDA程序应该同时具有以下几个特征：

在给定的数据规模下，选用算法的计算复杂度不明显高于最优的算法；

active warp的数量能够让SM满载，并且active block数量大于2，能够有效的隐藏访存延迟；

当瓶颈出现在指令流时，指令流效率已经经过了充分优化；

当瓶颈出现在访存或者IO时，已经选用了恰当的存储器和和存储器访问方式以获得最大带宽。

按照开发流程的先后顺序，CUDA程序的编写与优化需要解决以下问题：

1．确定任务中的串行部分和并行部分，选择合适的算法。首先，需要将问题分为几个步骤，并确定哪些步骤可以用并行算法实现，并确定要使用的算法。

2．按照算法确定数据和任务的划分方式，将每个需要并行实现的步骤映射为一个满足CUDA两层并行模型的内核函数。在这里就要尽量让每个SM上拥有至少6个活动warp和至少2个活动线程块。

3．编写一个能够正确运行的程序，作为优化的起点。程序必须能够稳定运行，不能发生存储器泄漏的情况。为了保证结果正确，在必要的时候必须使用memory fence、同步、原子操作等功能。在精度不足或者发生溢出时必须使用双精度浮点或者更长的整数类型。

3．优化显存访问，避免显存带宽成为瓶颈。在显存带宽得到完全优化前，其他优化不会产生明显结果。显存访问优化中可以使用的技术包括：

将可以采用相拓扑实现的几个kernel合并为一个，减少对显存访问；

除非非常必要，应该尽力避免将线程私有变量分配到local memory；

为满足合并访问，采用cudaMallocPitch()或者cudaMalloc3D()分配显存；

为满足合并访问，对数据类型进行对齐（使用__align）；

为满足合并访问，保证访问的首地址从16的整数倍开始，如果可能，尽量让每个线程一次读32bit；

在数据只会被访问一次，并且满足合并访问的情况下可以考虑使用zerocopy；

在某些情况下，考虑存储器控制器负载不均衡造成分区冲突的影响。

使用用有缓存的常数存储器和纹理存储器提高某些应用的实际带宽。

4．优化指令流。由于编译器会进行一些优化，而编译过程基本无法控制，所以指令流优化不一定能获得立竿见影的效果。但是，仍然有一些准则可以参考，包括：

如果只需要少量线程进行操作，那么一定记得要使用类似(if threadID < N)的方式，避免多个线程同时运行占用更长时间或者产生错误结果；

在不会出现不可接受的误差的前提下采用CUDA算术指令集中的快速指令；

使用#unroll，让编译器能够有效的展开循环；

采用原子函数实现更加复杂的算法，并保证结果的正确性；

避免多余的同步；

如果不产生bank conflict的算法不会造成算法效率的下降或者非合并访问，那么就应该避免bank conflict。

5．资源均衡。调整shared memroy和register的使用量。为了使程序能够获得更高的SM占用率，应该调整每个线程处理的数据数量、shared memory和register的使用量。这需要在三者间进行调整。当每个线程处理的子任务间有一定的公用部分时，可以考虑让一个线程处理更多的数据来提高指令流和访存效率。为了获得更高的SM占用率，必须控制每个线程的shared memory和register的使用量。通过调整block大小，修改算法和指令，以及动态分配shared memory，都可以提高shared的使用效率。而减小register的使用则相对困难，因为register的 使用量并不是由内核程序中声明的变量大小决定，而是由内核程序中使用寄存器最多的时刻的用量决定的。由于编译器会尽量减小寄存器的用量，因此实际使用的寄 存器有可能会小于在程序中声明的量。但是在通常情况下，由于需要暂存中间结果并且一些指令也需要更多的寄存器，一般寄存器用量都大于内核程序中声明的私有 变量的总数量。使用以下方法可能可以节约一些寄存器的使用：使用shared memory存储变量；使用括号更加明确的表示每个变量的生存周期；用对[u]long型或的处理代替对两个相邻的[u]short型或者四个相邻的[u]char型的处理；使用占用寄存器较小的等效指令代替原有指令，如用__sin函数代替sin函数。不过，由于不能对编译器的优化过程进行控制，即使使用了这些手段也不一定能减小寄存器的用量。值得注意的是，采用--maxrregcount编译选项只是让编译器将超出限制的私有寄存器分配在local memory中，造成较大的访存延迟。

6. 与主机通信优化。由于PCI-E带宽相对较小，应该尽量减少CPU与GPU间传输的数据量，并通过一些手段提高可用带宽。可用的技术包括：

使用cudaMallocHost分配主机端存储器，可以获得更大的带宽；

一次缓存较多的数据，再一并传输，可以获得较高的实际带宽；

需要将结果显示到屏幕时，直接使用与图形学API互操作功能完成，避免将数据传回；

       使用流式处理隐藏与主机的通信时间。

       对于用CUDA C语 言编写的程序，按照上述流程进行优化，是比较适合的。不过在优化中，各种因素往往相互制约，很难同时达到最优。读者需要按照要处理问题的类型，瓶颈出现的 部位和原因具体分析。按照预想进行优化也不是总能达到预想中的效果，有时优化手段反而会降低性能。在实践中，仍然需要不断实验各种优化方法，在不断试验与 迭代中一步步排除不可行的方案，最后得到一个比较理想的方案。

使用CUDA C并不总是能够编译到最优的指令。如果确实必要，可以用PTX优化程序中最关键的步骤。

除此以外，还要灵活采用宏和模版，动态分配内存和显存以及动态划分数据等手段提高程序的通用性，并在处理不同规模、不同数据类型的问题时选用不同的优化策略。

在2010年将要推出的新处理器中，各种存储器的带宽和延迟会有一定的调整，大部分指令的吞吐量也会有非常显著的提升。随着GPU架构的进一步改进和编译器性能的不断提高，下一代GPU上的CUDA程序优化工作会变得更加简单。