CUDA优化总结

总体的CUDA优化分以下几个策略：

1. 最大化并行执行来获得最大的利用率
2. 优化内存使用方法来获得最大的内存吞吐量
3. 优化指令使用方式来获得最大的指令吞吐量

具体再可以从以下几个方面入手：

CUDA配置优化：指的是在CUDA核函数执行之前，其参数的设置，根据显卡的参数，结合核函数的特点，需要注意以下几个参数：

1. 占用率：每个多核（个人觉得也可以理解成CUDA流处理）上正在执行的线程数，除以多核上允许执行的最大线程数，对于gtx 1060，其多核上允许同时执行的最大线程数是2048，其块内允许最大线程数是1024，每个多核内运行的最大块数是32（这个还不是很确定），在这些参数下，当线程块的个数是32时，那么每个线程块内的线程数是64。高的线程占用率并不总是意味着高性能：存在一个临界点，当线程占用率超过某个值以后并不能够提升性能。但是，低的线程占用率却总是会导致性能降低。
   a 计算占用率：决定占用率的几个因素之一是可获取的寄存器个数。寄存器运行线程中的局部变量快速访问，但是，在多核中的所有线程都需要共享这些寄存器，并且寄存器的数量是有限的，并且寄存器是一次性分配给整个线程块，所以，如果线程块使用了大量的寄存器，那么多核中的线程块的个数就会降低，这就会导致低的占用率（个人理解，应该是降低线程块内每个线程的寄存器使用数量，这样可以增大占用率）。gtx 1060中每个块内的共享寄存器大小是48k，每个多核中的共享寄存器是64k，每个线程可以使用的最大寄存器个数是255个（每个寄存器是32位），每个线程块中的线程数最好是64的倍数，这样才能使性能最优
2. 核并发执行
3. 多上下文，在一个CUDA应用，中，在同一个GPU上，应该避免使用多个上下文。为了辅助这种情况，CUDA驱动API提供了一种方法去获取和管理每个GPU上的特殊的上下文，也叫首要上下文(primary context)。如果对于一个线程没有现成的上下文可以用，那么CUDA运行时将隐式的使用这个上下文
4. 隐藏寄存器独立性，由于访问寄存器并不消耗额外的时钟周期，但是由于寄存器的读写具有独立性和存储的bank冲突，这会导致24个时钟周期的延迟
5. 线程与线程块选择的启发性：由于有许多因素参与到块的大小选择，因此不可避免的需要一些经验。所以下面的几条应该遵循：
   a 线程块里面的线程数应该是线程warp(一般是32)的倍数
   b 当一个多核上有多个线程块的时候，线程块内的线程数最小应该是64(或者是其倍数)
   c 当你尝试不同的线程块内的线程数的时候，128到256之间的数是一个好的选择
   d 如果寄存器的latency影响性能的时候，应该使用多个线程块(3到4个)，而不是使用一个大的线程块。当核函数频繁调用__syncthread()函数的时候更加有用
   需要注意的是，当线程块分配的寄存器个数超过了多核上现有的寄存器数量的时候，核函数将不能够启动，这对共享寄存器也是一样
6. 共享内存的影响：在前面举的例子中(使用共享存储器和不实用共享存储器计算矩阵相乘)，使用共享存储器能够减少对全局内存的访问，提高系统性能。然而他也同样能够对占用率产生限制，在许多情况下，一个核函数需要的共享存储器是与其线程块的块数相关的，比如：
   一个核函数可能需要一个32X32的共享存储器，但是每个线程块能够达到的最大线程数是512，那么，就不太可能为每个块分配一个32X32的共享存储器，在这种情况下，为每个块分配32X16或32X8的线程数，每个线程处理2个或4个共享存储器元素

CUDA指令优化：
算术指令：单精度提供了最好的性能，高度推荐使用单精度浮点
单个算术操作的吞吐量在 CUDA_C_Programming_Guide中有详细讲解

1. 除法取模操作，使用移位操作，替换昂贵的除法和取模操作：整数除法和取模操作特别耗时，并应该尽量避免或用位操作来替换：如果n是2的指数，那么，
   (i/n)的结果和(i>>log2(n))一样，(i%n)和(i&(n-1))结果一样
2. 倒数平方根：单精度的倒数平方根应该显式的调用rsqrtf()，双精度的应该调用rsqrt(),
3. 其他算术指令：避免自动的双精度到单精度的转换。在以下两种情况下，编译器将时不时的插入转换指令，增加额外的执行周期：
   a 函数操作char或者short时，其值通常要转换到int
   b 双精度浮点常量(没有后缀的浮点数如：1.0是双精度，1.0f是单精度)，作为输入进行单精度浮点计算
   对于b，性能的影响外还有精度的影响，注意到这个不同之处在计算能力为2.x的设备上特别明显
4. 小分数的指数运算：参见CUDA_C_Best_Practices_Guide，11.1.4
5. 数学库：当速度要求大于精度要求时，建议使用快速数学库：有两种类型的运行时数学操作，可以根据其名字来区分，前面有下划线的和前面没有下划线的：
   比如__functionName(),functionaNmae().有下划线的函数将会被直接映射到硬件级别去执行，其速度快，但是精度较低。并且，其吞吐量大。相反，没有下划线的速度慢，但是精度高，吞吐量小。
6. 与编译相关的nvcc的编译选项：
   -use_fast_math会将每一个遇到的没有下划线的函数替换成有下划线的函数
   -ftz 英文愿意：denormalized numbers are flushed to zero
   -prec-div=false(较低精度的除法)
   -prec-sqrt=false(较低精度的平方根)

CUDA内存指令优化：
高度推荐：将全局内存的访问降到最低：转而尽可能的使用共享内存
内存指令包括了任何的读写，共享、局部、全局内存的操作。
推荐使用restrict关键字修饰参数，这样可以减少指令的个数，从而提高性能

CUDA控制指令优化
CUDA控制指令：if switch do for while，高度推荐：避免在一个warp（32线程）内执行不同的执行路径
任何流控制指令能够显著的影响指令吞吐量：因为其（不执行同一个执行路径）会导致同一个warp内的线程不收敛
当不执行同一个执行路径时，不同的执行路径江北串行化，因为线程束中的所有线程共享同一个程序计数器，这将增加这个warp内的总的指令数

1. 分支预测

2. 循环计数器：有符号vs无符号：在循环中，使用有符号的计数器而不是无符号的计数器

3. 循环中发散线程的同步：高度推荐：避免在一个发散代码中使用__syncthreads(),即，要保证循环中的所有线程都会执行到这个函数，要是有某个线程执行不到这个函数，那就将出错