CUDA 程序的优化(1) 概述

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

1.精度

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

2.延迟

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

3.计算量

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

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

从相对计算量来说，如果可以并行的部分在整个应用中所占的比例不大，那么GPU对程序整体性能的提高也不会非常明显。如果整个应用中串行部分占用时间较长，而并行部分较短，那么也需要考虑是否值得使用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、同步、原子操作等功能以及volatile关键字。在精度不足或者发生溢出时必须使用双精度浮点或者更长的整数类型。

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

可以采用相同的block和grid维度实现的几个kernel合并为一个，减少对显存的访问。

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

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

为满足合并访问，对数据类型进行对齐(使用_align)o

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

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

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

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

5)优化指令流。在编译过程中，编译器对代码会进行一些优化。但是程序员很难直接控制编译器对代码的优化，所以指令流优化不一定能获得立竿见影的效果。但是，仍然有一些准则可以参考，包括:

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

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

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

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

避免多余的同步。

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

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

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

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

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

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

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

使用zero-copy技术和Write-Combined memory提高可用带宽。

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

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

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

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