最大化指令吞吐量
算术强度本质上是指访问内存所花的时间内做的数学运算次数,所以要最大化算术强度,可参考如下优化思路:
- 尽量少使用低吞吐量的算术指令。
算术指令及其吞吐量之间的关系请参考《CUDA C++ Programming Guide》。
A100(计算能力8.0)典型的算术指令及吞吐量关系表如表1所示。
- 在不影响最后结果的情况下使用低精度类型,使用内部计算接口。
比如,Shuffle指令是一组针对warp的指令,Shuffle指令最重要的特性就是warp内的寄存器可以相互访问。在没有shuffle指令的时候,各个线程在进行通信时只能通过shared memory来访问彼此的寄存器。而采用了shuffle指令之后,warp内的线程可以直接对其他线程的寄存器进行访存,通过这种方式可以减少访存的延时。目前绝大多数访存类算子,如softmax,batch_norm,reduce等,都是用Shuffle实现的,对一个warp里面的32个线程进行reduce操作代码如下所示:
__device__ __forceinline__ float warpReduceSum(float sum){ if(blockSize >= 32)sum += __shfl_down_sync(0xffffffff,sum,16); if(blockSize >= 16)sum += __shfl_down_sync(0xffffffff,sum,8); if(blockSize >= 8)sum += __shfl_down_sync(0xffffffff,sum,4); if(blockSize >= 4)sum += __shfl_down_sync(0xffffffff,sum,2); if(blockSize >= 2)sum += __shfl_down_sync(0xffffffff,sum,1); return sum; }
另外CUDA支持许多常见的数学运算,比如sin/cos/exp,这些函数的计算精度比C语言的内置数学库少2~3位,但速度快很多。
- 在不影响最终结果的情况下用精度换取速度,使用单精度float而不是双精度double。
- 尽量减少流程控制指令,即减少条件分支。比如:
- 控制条件仅取决于(threadIdx / warpSize)。
- 使用#pragma unroll展开循环。
- 尽量使用__fdividef(x,y)单精度浮点除法,其运算效率优于除法运算符。
- 大部分情况下,rsqrtf()效率比1.0 / sqrtf()高。
- 使用三角函数时,尽量使用单精度,且x值域较小为宜。
- 尽量使用位运算代替整数除法和模运算。
- 半精度计算,使用half2代替half。
- 调用接口时,尽量保持参数类型一致,否则会有类型转换的消耗。比如函数参数是int类型,输入是char类型或者short类型,需要先将类型转化为int类型。
- 寄存器优化,比如GPU的位操作微观优化。
父主题: kernel代码优化