基于CUDA的GPU优化方法

初学CUDA，往往拿到代码无从下手，也没有什么明确的思路。我想有必要把前人的经验总结拿出来，便于后来者更快掌握这门技术。

对于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。

但是有一点切记，高的占有率并不代表高性能。伯克利还专门发过一篇文章讲为什么会出现这样的情况。

下面给大家一些无脑记忆的，所谓常识吧。

block里面thread个数最好为wrap大小的倍数，即：32的倍数。使得计算效率更高，保证memory coalescing。
如果multiprocessor中有多个active block时，每个block里面的thread个数最好为64的倍数。
当选择不同的block大小时，可以先确定block里面thread个数为128到256之间，然后再调整grid中block大小。
如果是让问延迟latency造成程序性能下降时，考虑在一个block里面采用小block划分，不要在一个multiprocessor中分配一个很大的block，尽量分配好几个比较小的block，特别是程序中使用了__syncthreads()，这个函数是保证block里面所有wrap到这里集合，所以，block里面的thread越少越好，最好是一个wrap或者两个wrap，这样就可以减少__syncthreads()造成的访问延迟。
如果如果一个block里面分配的register超过了multiprocessor的最大极限时，kernel的launch就会fail。

下面给出一些比较低级的，但是实际应用中非常有效的优化细节（有时是通用的，在其他优化加速里也会用到）：

1.位运算：可以考虑用位运算来代替除法和取余，如果n是2幂数，(i/n)=(i>>log2(n)), (i%n)=(i&(n-1))。这种方法只适合interger类型的运算，float类型千万不要用。

2.利用fast math library。例如__sinf(x)、__expf(x)就比sinf(x)和expf(x)有更快的速度。

3.global memory和share memory的使用，尽量多利用share memory

定期更新