CUDA程序性能调优

1、kernel function parameters

2、 local variables

3、shared memory with __syncthreads__ call

4、device function calls

5、loop/if

参考文献：CUDA程序调优指南（二）：性能调优 - 知乎

介绍了GPU的结构以及资源的控制要素（GPU硬件结构和程序具体参数设置_yu132563的专栏-CSDN博客），下面就可以对CUDA进行程序的调优，从而在不同的GPU上面运行同一个CUDA程序的参数设置方法。

对于一个CUDA kernel function而言，其通常由如下几个部分组成：

kernel function paras
local variables
shared memory with __syncthreads__ call
device function call
loop/if
<<<BlocksNum, ThreadsNumPerBlock>>>

我们分别考虑如何对这些部分进行优化。

1、kernel function parameters

__global__ function的参数是存在constant mem里的，并且其大小被限制为4KB。通常来说，我们传入的参数都比较少，因此这些参数大部分是直接缓存在SM上的register上的，因此读取起来最快。但如果register不够用，那么就会被放到constant mem/cache上，速度就会变慢。

2、 local variables

对于单独的local var，其会被放在register里面，因此读写极快。对于数组类型的local var，根据访问pattern的不同，速度也不同（见per-thread array的访问）。总结而言是这样的：

【Static Indexing】：会被放到register里，读写极快
【Dynamic indexing with Uniform Access】：会被放到Local Mem里，略慢。但 更高的math/load比 配合 缓存到L1/L2cache里，也可以很快
【Dynamic indexing with Non-uniform Access】：由于会让SM产生多个访存指令，所以最慢。但可以通过把其放到shared mem里加速

3、shared memory with `syncthreads` call

shared memory需要尽量避免bank confilicts，这样读取最快，一个cycle clock就可以读取128byte。除此之外，因为shared memory由同一block内的thread共享，所以在初始化shared mem之后，需要调用__syncthreads__来对同一block内的所有threads进行同步。

4、device function calls

编译器会自行决定device function是否会被inlined（多数情况下会inline）。通常来说，inline会更好，因为少了函数调用的开销。

5、loop/if

在书写kernel function时，应尽量避免loop/if，因为这两个在代码里引入了分支结构。

如果代码有分支，那么

SM首先执行那些进入first branch的threads，此时其余threads在等待
然后执行进入second branch的threads，前面的线程等待。

因此在可能的情况下，我们应该尽量使用模板参数来替换掉loop/if：

如果if所判断的条件在kernel launch时就能确定（即作为kernel的一个parameter传入的），那就可以用模板参数来代替该parameter，这样编译器会自动优化掉另一分支。
同理，如果for-loop的边界也在kernel launch时就能确定，那也可用模板参数代替。

针对一些边界是常量（如0->5）的循环，在循环体足够简单的情况下，可以使用#pragma unroll来告诉编译器展开该循环。

参考文献：CUDA程序调优指南（二）：性能调优 - 知乎