[CUDA KERNEL] Some tips for write CUDA kernel function

1 CUDA device端 内建函数

__exp(),__log(),__power(),...这些函数降低了精度，但是提高了运算速度

 

2 线程同步 __syncthreads() 这个函数是“block线程内”的线程同步

以前有一次发现使用到__syncthreads()的程序，发现经常运行的结果不一样，就是因为__syncthreads() 这个函数是“block线程内”的线程同步，改用一个线程块就好了

例如：squareKernel<<<(dim + 511) / 512, 512, 0, stream>>>(dim, input, buffer);

参数1（blocksPerGrid）：(dim + 511) / 512, ---使用“block线程内”的线程同步函数时，这个参数必须为1

 

3 线程块内 Wrap的分割 对线程优化的作用

视频中举得例子是规约求和，方法分别是按照 1，2，4...和 ...4，2，1的方式，看上去程序一样，但是，运行完成的线程，可以以wrap为单位，释放GPU运算资源

使相邻线程一致性提高，减少线程发散

概念：block内划分以32个线程为1个wrap，wrap是最基本的运行和调度单元，所以写代码时，最好要让同wrap内的程序逻辑一致性提高

 

4 使用预处理指令，使代码循环展开：同时也可以优化，少量增加运行速度

例如：

#pragma unroll (16)

for (int i=0;i<BLOCK_DIM;i++)

{

pValue+=Mds[ty][i]*Nds[i][tx];

}

 

[C++11]

constexpr

constexpr是C++11中新增的关键字，其语义是“常量表达式”，也就是在编译期可求值的表达式。

最基础的常量表达式就是字面值或全局变量/函数的地址或sizeof等关键字返回的结果，而其它常量表达式都是由基础表达式通过各种确定的运算得到的。

constexpr值可用于enum、switch、数组长度等场合。

constexpr的好处：

是一种很强的约束，更好地保证程序的正确语义不被破坏。

编译器可以在编译期对constexpr的代码进行非常大的优化，比如将用到的constexpr表达式都直接替换成最终结果等。

相比宏来说，没有额外的开销，但更安全可靠。