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表达式都直接替换成最终结果等。
相比宏来说,没有额外的开销,但更安全可靠。