技术开发 频道

推荐CUDA程序优化的15个策略

    【IT168 技术】在《CUDA程序优化策略》这篇文章中,我们介绍过CUDA优化的常见策略。今天我们会对CUDA优化策略进行详细讲解。具体策略如下:

  1. memory coalescing,保证内存融合。因为global memory在CC为1.x上是按照half wrap进行访问读写的,而在2.x上是按照wrap进行访问读写的。在显存中,有多个存储器控制器,负责对显存的读写,因此,一定要注意存储器控制器的负载均衡问题。每一个存储器控制器所控制的那片显存中的地址空间称为一个分区。连续的256Byte数据位于同一个分区,相邻的另一组256Byte数据位于另一个分区。访问global memory就是要让所有的分区同时工作。合并访问就是要求同一half-wrap中的thread按照一定byte长度访问对齐的段。在1.0和1.1上,half-wrap中的第k个thread必须访问段里的第k个字,并且half-wrap访问的首地址必须是字长的16倍,这是因为1.0和1.1按照half-wrap进行访问global memory,如果访问的是32bit字,比如说一个float,那么half-wrap总共访问就需要16个float长,因此,每个half-wrap的访问首地址必须是字长的16倍。1.0和1.x只支持对32bit、64bit和128bit的合并访问,如果不能合并访问,就会串行16次。1.2和1.3改进了1.0和1.1的访问要求,引进了断长的概念,与1.0和1.1上的端对齐长度概念不同,支持8bit-段长32Byte、16bit-段长64Byte、32bit-64bit-128bit-段长128Byte的合并访问。对1.2和1.3而言,只要half-wrap访问的数据在同一段中,就是合并访问,不再像1.0和1.1那样,非要按照顺序一次访问才算合并访问。如果访问的数据首地址没有按照段长对齐,那么half-wrap的数据访问会分两次进行访问,多访问的数据会被丢弃掉。所以,下面的情况就很容易理解:对1.0和1.1,如果thread的ID与访问的数据地址不是顺序对应的,而是存在交叉访问,即:没有与段对齐,那么,就会16次串行访问,而对1.2和1.3来讲,会判断这half-wrap所访问的数据是不是在同一个128Byte的段上,如果是,则一次访问即可,否则,如果half-wrap访问地址连续,但横跨两个128Byte,则会产生两次 传输,一个64Byte,一个32Byte。当然,有时还要考虑wrap的ID的奇偶性。1.2和1.3放宽了对合并访问的条件,最快的情况下的带宽是最好的情况下的带宽的1/2,然而,如果half-wrap中的连续thread访问的显存地址相互间有一定的间隔时,性能就会灰常差。比如,half-wrap按列访问矩阵元素,如果thread的id访问2*id的地址空间数据,那么,半个wrap访问的数据刚好是128Byte,一次访问可以搞定,但是,有一半数据会丢失,所以,也表示浪费了带宽,这一点一定要注意。如果不是2倍,而是3倍、4倍,那么,有效带宽继续下降。在程序优化时,可以使用share memory来避免间隔访问显存。

  2. bank conflict,bank冲突。先说一下,share memory在没有bank conflict情况下,访问速度是global和local的100倍呢,你懂的。类似global memory的分区,share memory进行了bank划分。如果half-wrap内的很多thread同时要求访问同一个bank,那么就是bank conflict,这时,硬件就会将这些访问请求划分为独立的请求,然后再执行访问。但是,如果half-wrap内所有thread都访问同一个bank,那么会产生一次broadcast广播,只需要一次就可以相应所有访问的请求。每个bank宽度长度为32bit。对于1.x来讲,一个SM中的share memory被划分为16个bank,而2.x是32个bank。1.x的bank conflict和2.x的bank conflict是不一样的。对1.x来讲,多个thread访问同一个bank,就会出现bank conflict,half-wrap内所有thread访问同一个bank除外。但是,对2.x来讲,多个thread访问同一个bank已经不再是bank conflict了。比如:

  __shared__ char Sdata[32];

  
char data = Sdata[BaseIndex+tid];

  在1.x上属于bank conflict,因为,0~3thread访问同一个bank,4~7访问同一个bank,类推,这种情况属于4-way bank conflict。但是,对于2.x来讲,这种情况已经不是bank conflict了,以为2.x采用了broadcast机制,牛吧,哈哈。 这里要多看看矩阵乘积和矩阵转置例子中的share memory的使用,如何保证memory coalescing和避免bank conflict的。

  3. texture memory是有cache的,但是,如果同一个wrap内的thread的访问地址很近的话,那么性能更高。

  4.以下是要注意的:

  (1)在2.x的CC上,L1 cache比texture cache具有更高的数据带宽。所以,看着使用哈。

  (2)对global memory的访问,1.0和1.1的设备,容易造成memory uncoalescing,而1.2和1.3的设备,容易造成bandwidth waste。 而对2.x的设备而言,相比1.2和1.3,除了多了L1 cache,没有其他的特别之处。

  (3)采用-maxrregcount=N阻止complier分配过多的register。

  (4)occupancy是每个multiprocessor中active wrap的数目与可能active wrap的最大数目的比值。higher occupancy并不意味着higher performance,因为毕竟有一个点,超过这个点,再高的occupancy也不再提高性能了。

   5.影响occupancy的一个因素,就是register的使用量。比如,对于1.0和1.1的device来讲,每个multiprocessor最多有8192个register,而最多的simultaneous thread个数为768个,那么对于一个multiprocessor,如果occupancy达到100%的话,每个thread最多可以分配10个register。另外,如果在1.0和1.1上,一个kernel里面的一个block有128个thread,每个thread使用register个数为12,那么,occupancy为83%,这是因为一个block有128个thread,则,由于multiprocessor里面最大的simultaneous thread为768,根据这个数目计算,最多同时有6个active block,但是6个active block,就会导致总共thread个数为128*6*12个,严重超过了8192,所以不能为6,得为5,因为128*5<768, and 128*5*12<8192, 5是满足要求的最大的整数。如果一个kernel里面的一个block有256个thread,同样一个thread用12个register,那么occupancy为66%,因为active block为2。可以在编译选项里面加入--ptxas-options=-v查看kernel中每个thread使用register的数量。同时,NV提供了CUDA_Ocuppancy_calculator.xls作为occupancy计算的辅助工具。顺便说一下,对于1.2和1.3的device来讲,每个multiprocessor最多的simultaneous thread个数为1024个。

  6. 为了隐藏由于register dependent寄存器依赖造成的访问延迟latency,最小要保证25%的occupancy,也就是说,对于1.x的device来讲,一个multiprocessor最少得发起192个thread。对于1.0和1.1来讲, occupancy为192/768=25%,达到要求,但是对于1.2和1.3而言,192/1024=18.75%,不过,也只能这样。对于2.x系列的device来讲,由于是dual-issue,一个multiprocessor最多发起simultaneous thread个数为1536个,所以,一个multiprocessor最少同时发起384个thread时,occupancy为384/1536=25%,又达到了25%。

  7. 对于block和thread的分配问题,有这么一个技巧,每个block里面的thread个数最好是32的倍数,因为,这样可以让计算效率更高,促进memory coalescing。其实,每个grid里面block的dimension维度和size数量,以及每个block里面的thread的dimension维度和size数量,都是很重要的。维度呢,采用合适的维度,可以更方便的将并行问题映射到CUDA架构上,但是,对性能不会有太大改进。所以,size才是最重要的,记住叻! 其实,访问延迟latency和occupancy占有率,都依赖于每个multiprocessor中的active wrap的数量,而active wrap的数量,又依赖于register和share memory的使用情况。首先,grid中block的数目要大于multiprocessor的数目,以保证每个multiprocessor里面最少有一个block在执行,而且,最好有几个active block,使得blocks不要等着__syncthreads(),而是占用了hardware。其次,block里面的thread的数目也很重要。对于1.0和1.1的设备来讲,如果一个kernel里面block的大小为512个thread,那么,occupancy为512/768=66%,并且一个multiprocessor中只有一个active block,然而,如果block里面的thread为256个thread,那么,768/256=3,是整数,因此,occupancy为100%,一个multiprocessor里面有3个active block。但是,记住了,higher occupancy don't mean better performance更高的占有率并不意味着更好的性能。还是刚才那个例子,100%的occupancy并不比66%的occupancy的性能高很多,因为,更低的occupancy使得thread可以有更多的register可以使用,而不至于不够用的register分配到local memory中,降低了变量存取访问速度。一般来讲啊,只要occupancy达到了50%,再通过提高occupancy来提高性能的可能性不是很大,不如去考虑如何register和share memory的使用。保证memory coalescing和防止bank conflict。记住如下几点:

  (1)block里面thread个数最好为wrap大小的倍数,即:32的倍数。使得计算效率更高,保证memory coalescing。

  (2)如果multiprocessor中有多个active block时,每个block里面的thread个数最好为64的倍数。

  (3)当选择不同的block大小时,可以先确定block里面thread个数为128到256之间,然后再调整grid中block大小。

  (4)如果是让问延迟latency造成程序性能下降时,考虑在一个block里面采用小block划分,不要在一个multiprocessor中分配一个很大的block,尽量分配好几个比较小的block,特别是程序中使用了__syncthreads(),这个函数是保证block里面所有wrap到这里集合,所以,block里面的thread越少越好,最好是一个wrap或者两个wrap,这样就可以减少__syncthreads()造成的访问延迟。

  (5)如果如果一个block里面分配的register超过了multiprocessor的最大极限时,kernel的launch就会fail。

  8. share memory的使用量也是影响occupancy的一个重要因子。thread与share memory的元素之间,没有必要是一对一的。一个线程可以一起负责处理share memory数组中的第一个、第二个以及第三个元素,都ok的。第一个thread处理share memory中的第一个元素,第二个thread负责处理第二个元素,类推如下,这种情况不是必须的,有时也没必要这么做。在代码里面,采用一个thread负责处理share memory数组中的多个元素的方法,是非常好的策略。这是因为如果share memory里面各个元素要进行相同的操作的话,比如乘以2,那么,这些操作可以被负责处理多个元素的一个thread一次搞定,分摊了thread处理share memory元素数量的成本费用。

  9. 当上面那些high level级别的优化策略都检查使用过以后,就可以考虑low level级别的优化:instruction optimization指令集优化。这个也可以很好的提高性能的。指令集的优化,可以稍微总结如下:

  (1)尽量使用shift operation位移运算来取代expensive昂贵的division除法和modulo取余运算,这里说的都是integer运算,float不行的。如果n是2幂数,(i/n)=(i>>log2(n)), (i%n)=(i&(n-1)). 其实,这只是一个量的问题,对于1.x的设备而言,如果一个kernel里面使用了十多个tens of这样的指令,就要考虑用位移运算来取代了;对于2.x的设备而言,如果一个kernel里面使用了20个这样的指令,也要考虑使用位移运算来取代除法和取余运算。其实,compiler有时会自动做这些转换的,如果n是2的幂数。

  (2)reciprocal square root,对于平方根倒数1.0f/sqrtf(x),编译器会采用rsqrtf(x)来取代,因为硬件做了很多优化。当然,对于double型的平方根倒数,就采用rsqrt(x)啦。呵呵,记住了。

  (3)编译器有时会做一些指令转化。在要做计算的单精度浮点型常数后面,一定要加入f,否则,会被当做双精度浮点运算来计算,对于2.x以上的设备来讲,这一点很重要,记好了。

  (4)如果追求速度speed,而不是精度precision,那么尽量使用fast math library。比如,__sinf(x)、__expf(x)比sinf(x)和expf(x)有更快的速度,但是,精度却差一些。如果是__sinf(x-2)则比sinf(x-2)的速度要快一个数量级,因为x-2运算用到了local memory,造成太高的访问延迟。当然,在compiler option中使用-use_fast_math可以让compiler强制将sinf(x)和expf(x)转化为__sinf(x)和__expf(x)进行计算。对于transcendental function超越函数,作用对象是单精度浮点型数据时,经常这么用,其他类型的数据,性能提升不大。

  (5)对于2和10为底做指数运算,一定要采用exp2()或者expf2()以及exp10()或者expf10(),不要采用pow()和powf(),因为后者会消耗更多的register和instruction指令。 另外,exp2()、expf2()、exp10()、expf10()的性能和exp()以及expf()性能差不太多,当然比pow()和powf()要快10多倍呢。加好了哈。

  (6)减少global memory的使用,尽量将global memory数据加载到share memory,再做访问。因为访问uncached的显存数据,需要400~600个clock cycle的内存延迟。

  10. 下一个就是control flow了。一定要避免在同一个wrap里面发生different execution path。尽量减少if、swith、do、for、while等造成同一个wrap里面的thread产生diverge。因为,一旦有divergence,不同的execution path将会顺序的串行的执行一遍,严重影响了并行性。但是:

switch(threadIdx.x)

{

case 0

break;

case 1:

break;

...

case 31:

break;

}

  上面这个例子,则不会发生divergence,因为控制条件刚好和wrap里面的thread相对应。

  其实,有时,compiler会采用branch predication分支预测来打开loop循环或者优化if和switch语句, 这时,wrap就不会出现divergence了。在写code时,我们也可以自己采用#pragma uroll来打开loop循环。在使用branch predication时,所有指令都将会执行,其实,只有预测正确的真正的执行了,而预测错误的,其实就是thread,不会去读取该instruction的地址和数据,也根本不会写结果。其实,编译器做分制预测,是有条件的,只有分支条件下的指令instruction的个数小于等于某个阈值的时候,才会做分支预测branch predication。如果编译器觉得可能会产生多个divergent wrap,那么阈值为7,否则为4。(这里很不理解7和4是怎么来的)。

  11. 在loop循环的counter,尽量用signed integer,不要用unsigned integer。比如:for(i = 0; i < n; i++) {out[i] = in[offset+stride*i];} 这里呢,stride*i可以会超过32位integer的范围,如果i被声明为unsigned,那么stride*i这个溢出语句就会阻止编译器做一些优化,比如strength reduction。相反,如果声明为signed,也没有溢出语句时,编译器会对很多地方做优化。所以,loop counter尽量设置为int,而不是unsigned int。

  12. 在1.3及其以上的device上,才支持double-precision floating-point values,即:64位双精度浮点运算。当使用double时,在编译器选项里面添加:-arch=sm_13

  13. 还有一点需要注意,如果A、B、C都是float,那么A+(B+C)并不一定等于(A+B)+C。

  14. 先看下面两个语句:float a; a = a * 1.02;

  对于1.2及其以下的device来讲,或者1.3及其以上device,但是没有打开支持double运算的选项,那么,由于不支持double,所以,1.02*a这个乘积是一个float;

  对于1.3及其以上的device来讲,如果打开了支持double运算的选项,那么,a*1.02是一个double,而将乘积赋值给a,这个结果是float,所以,是先做了从float到double的promotion扩展,然后做了从double到float的truncation截取。

  15. 多GPU编程。如果有p个GPU同时并行,那么,程序中就需要p个CPU threads。这些thread可以用OpenMP(小规模)或者MPI(大规模)进行管理。GPU之间的数据拷贝,必须通过CPU实现。对于OpenMP,是这样的:一个CPU thread将数据从对应的GPU中拷贝到host端的share memory region中,然后另一个CPU thread将数据从host端的share memory region拷贝到对应的GPU中。也就是说:OpenMP是通过share memory进行数据拷贝的。而对于MPI而言,数据是通过message passing进行传递的。一个CPU thread使用cudaMemcpy将数据从device拷贝到host,然后通过MPI_Sendrecv(),另一个CPU thread就使用cudaMemcpy将数据从host端拷贝到呃device端。编译选项,记着采用nvcc -Xcompiler /openmp或者nvcc -Xcompiler mpicc。

        更多内容请点击:

        CUDA专区:http://cuda.it168.com/

        CUDA论坛:http://cudabbs.it168.com/

2
相关文章