【转】CUDA程序优化要点

标签：style   blog   http   color   使用   os   io   strong   

CUDA程序优化应该考虑的点：
精度：只在关键步骤使用双精度，其他部分仍然使用单精度浮点以获得指令吞吐量和精度的平衡；

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

延迟：需要首先缓冲一部分数据，缓冲的大小应该可以保证每个内核程序处理的一批数据能够让GPU慢负荷工作；

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

计算量：计算量太小的程序使用CUDA很不合算；当需要计算的问题的计算密集度很低的时候，执行计算的时间远远比IO花费的时间短，整个程序的瓶颈出现在PCI-E带宽上。

　　　　衡量计算量有绝对和相对两种方式。

　　　　从绝对量来说，如果要优化的程序使用频率比较低，并且每次调用需要的时间也可以接受，那么使用 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 执行。

优秀的CUDA程序特征：

在给定的数据规模下，选用算法的计算复杂度不明显高于最优的算法；
Active warp的数量能够让SM满载，并且active block的数量大于2，能够有效地隐藏访存延迟；
当瓶颈出现在运算指令时，指令流的效率已经过了充分优化；
当瓶颈出现在访问IO时，程序已经选用了恰当的存储器来储存数据，并使用了适当的存储器访问方式，以获得最大带宽；

CUDA的编写与优化需要解决的问题：

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

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

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

4. 优化显存访问，避免显存带宽成为瓶颈。在显存带宽得到完全优化前，其他优化不会产生明显结果。

　　显存访问优化中可以使用的技术包括：合并采用相同block和grid的kernel；尽力避免将线程私有变量分配到local memory；

• 将可以采用相拓扑实现的几个 kernel 合并为一个，减少对显存访问；
• 除非非常必要，应该尽力避免将线程私有变量分配到 local memory ；
• 为满足合并访问，采用 cudaMallocPitch() 或者 cudaMalloc3D() 分配显存；
• 为满足合并访问，对数据类型进行对齐（使用 __align ）；
• 为满足合并访问，保证访问的首地址从 16 的整数倍开始，如果可能，尽量让每个线程一次读 32bit ；
• 在数据只会被访问一次，并且满足合并访问的情况下可以考虑使用 zerocopy ；
• 在某些情况下，考虑存储器控制器负载不均衡造成分区冲突的影响。
• 使用用有缓存的常数存储器和纹理存储器提高某些应用的实际带宽。

5.优化指令流：

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

• 如果只需要少量线程进行操作，那么一定记得要使用类似 (if threadID < N) 的方式，避免多个线程同时运行占用更长时间或者产生错误结果；
• 在不会出现不可接受的误差的前提下采用 CUDA 算术指令集中的快速指令；
• 使用 #unroll ，让编译器能够有效的展开循环；
• 采用原子函数实现更加复杂的算法，并保证结果的正确性；
• 避免多余的同步；
• 如果不产生 bank conflict 的算法不会造成算法效率的下降或者非合并访问，那么就应该避免 bank conflict 。

6.资源均衡：

• 为了使程序能够获得更高的 SM 占用率，调整每个线程处理的数据数量、 shared memroy 和 register 的使用量，这需要在三者间进行调整。当每个线程处理的子任务间有一定的公用部分时，可以考虑让一个线程处理更多的数据来提高指令流和访存效率。为了获得更高的 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-memory技术和Write-Combined memory提高可用带宽；

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

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

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

====================================================================================================

测量程序运行时间：
CUDA内核程序的运行时间：可以在设备端测量，也可以在主机端测量；
CUDA API的运行时间：只能在主机端测量；使用CUDA runtime API时，会在第一次调用runtime API函数时启动CUDA环境，计时的时候应该避免将这一部分计入，因此在正式测试之前应选择一个包含数据输入输出的.....，使得GPU从平时的节能模式进入工作状态，使得测试结果更加可靠；

　　设备端测量时间：

• 调用clock（）函数：返回的是GPU的时钟周期，需要除以GPU的运行频率才能得到以秒为单位的时间；
• 使用CUDA API事件管理功能；

　　主机端测量时间：

• 使用c标准库中的clock_t()函数测试，由于其精度很低，因此应该运行多次然后求平均运行时间；注意异步函数（比如内核函数和带有asyn后缀的存储器拷贝函数），在GPU上执行完成之前，CPU线程已经得到了它的返回值；从主机测量一系列CUDA调用需要的时间的时候，要首先调用cudaThreadSynchronize()函数等，使得GPU线程执行完毕后，进入CPU线程，从而得到正确的执行效果；在一串流中的第一个流（ID为0的流）的行为总是同步的，因此使用这些函数对0号流进行测时，得到的记过是可靠的。

任务划分原则：

• 在两次主机—设备通信之间进行尽量多的计算；
• 考虑使用流运算隐藏主机—设备通信时间，通过Pinned memory、zero—copy、write—combined memory等手段提高实际传输带宽；
• 尽量使得每个block中线程数量是32的整数倍，最好保持在64~256之间，并根据任务的具体情况确定每个维度上的大小，以减少计算访存地址时的整数除法和求模运算；
• 对一个block的任务进行划分后，再按照block的维度和尺寸要求对grid进行划分：每个block 的访存均匀分布在显存的各个分区中；block 间的负载可以存在一定程度的不均衡；

Grid和Block的维度设计：

• 首先考虑block的尺寸，grid的尺寸一般越大越好；
• 每个SM中至少要有6个active warp用于隐藏流水线延迟，并且拥有至少2个active block；

　　计算每个SM上active warp和active block的数量：

1. 确定每个SM使用的资源数量：使用nvcc的—keep编译选项，或者在.cu编译规则（cuda build rule）中选择保留中间文件，得到.cubin文件，用写字板打开后可以看到imem 和reg分别代表内核函数中每个线程使用的local memory和register数量；
2. 根据硬件确定SM上的可用资源：可以用SDK中的deviceQuery 获得每个SM中的资源；根据内核不同，SM上的warp总数上限，block总数上限，寄存器数量，shared memory数量都不同；
3. 每个block中的线程数量不能超过512；（？？？）
4. 计算每个block使用的资源，并确定active block和active warp数量：

     e.g. 每个block中有64个线程，每个block使用256 Byte shared memory，8个寄存器文件，
　　　　那么：每个人block使用的shared memory： 256 Byte；每个block使用的寄存器文件数量： 8*64 = 512；

      　　每个block中使用的warp数量：64/32 = 2;

           如果在G80/92 GPU中运行这个内核程序：
           由shared memory数量限制的active block数量： 16384、256  = 64；
           由寄存器数量限制active block数量：8192/512 = 16;

           由warp数量限制的active block数量 24/2 = 12;

           每个SM中的最大active block数量：8；
　　　　这些计算可以由NVIDIA在CUDA SDK中提供的 CUDA occupancy calculator完成；

• Block 的维度和每个维度上的尺寸的主要作用是避免做整数除法和求模运算，对执行单元效率没有什么显著影响；

计算grid中各个维度上block的数量：grid在x轴上的block数量 = （问题在x轴上的尺寸+每个block在x轴上的尺寸-1）/每个block在x轴上的尺寸；
 

存储器访问优化：    

 主机—设备通信优化：
      目前一条PCI—E 2.0*16通道的理论带宽是每向8GB/s, 远小于显存和GPU片内存储器带宽；
      Pinned memory：强制让操作系统在物理内存中完成内存申请和释放工作，不用参会页交换，因此速度比pageable memory快；
                         声明这些内存会占用操作系统的可用内存，可能会影响到操作系统运行需要的物理内存；
                         需要合理规划CPU和GPU各自使用的内存，使整个系统达到最优；
异步执行：
      内核启动和显存内的数据拷贝（Device to Device）总是异步的；
      内存和显存间的数据拷贝函数有异步和同步两个版本：
　　同步(顺序执行)： cudaMemcpy(a_d,a_h,size,cudaMemcpyHostToDevice);

                               cpuFunction();

　　异步(同时执行)： cudaMemcpyAsync（…………）；
                                               cpuFunction()；
      属于同一个流中的内核启动总是同步的；
      如果几次内核启动属于不同的流，那么他们的执行可能是乱序的；
利用异步提高计算效率：
      使用流和异步是CPU和GPU同事进行运算；
      利用不同流之间的异步执行，使流之间的传输和运算能够同时执行，更好地利用GPU资源；
全局存储器访问优化：
       需要考虑half-warp访问的对齐问题，不同的硬件要求不同；（存疑？？？？？？？？）
       采用合并访问；
       尽量避免间隔访问：比如按列访问矩阵，可以借助shared memory来实现这一点；
Shared memory访问优化：
       共享存储器被组织为16个可以被同时访问的存贮器模块，称为bank；
Bank组织方式：宽度32bit，相邻的32bit字被组织在相邻的bank中，每个bank在每个时钟周期可以提供32bit的带宽；
一个warp被分为两个half-warp进行访问；
避免bank conflict：在SDK中，使用宽度为17或则会threadDim.x+1的行来避免bank conflict；（存疑？？？？？？？？）
Shared memory采用了广播机制：在相应一个对同一个地址的读请求时，一个32bit字可以被读取并同时广播给不同的线程；
当一个half-warp中有多个线程读取同一个32bit字地址中的数据时，可以减少bank conflict的数量；
如果half-warp中的线程全都读取同一地址中的数据时，此时完全不会发生bank conflict；
如果half-warp内有多个线程要对同一地址进行读写操作，此时会产生不确定结果，这种情况应该使用shared memory的原子操作；
共享存储器保存着加载kernel时传递过来的参数，以及kernel执行配置参数，如果参数列表很长，应该将其中一部分参数放入constant memory；
使用纹理存储器：
        主要用于存放图像和查找表：不用严格遵守合并访问条件，就能达到较高带宽；
                                  对于少量数据的随机访问，效率不会太差；
                                  可以使用线性滤波和自动类型转换等功能调用硬件的不可编程计算资源，不必占用可编程计算单元；
使用常数存储器：
        主要用于存放指令中的常数；速度低于shared memory；
指令流优化：
增大吞吐量手段：
避免使用地吞吐量指令；
       优化每种类型的存储器，有效利用带宽；
       允许线程调度单元精良用多的数学计算来覆盖访存延迟，需要有教导的算术密度；
吞吐量：每个多处理器在一个时钟周期下执行的操作数目；
算术指令：尽量使用单精度浮点单元进行运算，在计算能力小于等于1.2的设备中，每个双精度的变量将会转换成单精度格式，双精度运算也会转为单精度算术运算；
          单精度浮点基本算术运算：加，乘，乘加运算的吞吐量是每个时钟周期8个操作；
          求导数运算：每个时钟周期2个操作；
          单精度除法：每个时钟周期0.88个操作；
          单精度浮点倒数平方根：2；
          平方根：1；
          对数：2；
          正弦余弦：参数较大的时候，采用归约操作将x的绝对值减小；有快路径和慢路径（大参数）；
          整数算术运算：整数加法（8），乘（2）；除法和取模开销特别大，尽量地避免或者用位运算代替；
          比较，min，max：（8）；
          位运算（8）；
          类型转换（8）；
控制流指令： If, switch, do, for, while 可能引起一个warp线程跳转到不同的分支，严重影响指令吞吐量；
访存指令：包括任何读写memory的指令；
对于local memory只有在register不够用或者编译器无法解析的时候才会发生；
将较大的数据（float，double）拆分成每个线程32bit，或者将多个[u]char,[u]short合并成每个线程32bit的形式访问；
在访问local/global memory时候，会有额外的400~600个时钟周期的访问延迟；
同步指令：_syncthreads()的吞吐量是每时钟周期8个操作；

http://blog.csdn.net/ouczoe/article/details/5137063

【转】CUDA程序优化要点,布布扣,bubuko.com

【转】CUDA程序优化要点

标签：style   blog   http   color   使用   os   io   strong   

原文地址：http://www.cnblogs.com/huangshan/p/3918368.html