EN
注册
我要评分
文档获取效率
文档正确性
内容完整性
文档易理解
在线提单
论坛求助
鲲鹏小智

最大化指令吞吐量

算术强度本质上是指访问内存所花的时间内做的数学运算次数,所以要最大化算术强度,可参考如下优化思路:

  • 尽量少使用低吞吐量的算术指令。

    算术指令及其吞吐量之间的关系请参考《CUDA C++ Programming Guide》。

    A100(计算能力8.0)典型的算术指令及吞吐量关系表如表1所示。

    表1 A100算术指令及吞吐量关系

    算术指令

    吞吐量

    16bit浮点加法、乘法、乘加运算

    256

    32bit浮点加法、乘法、乘加运算

    64

    64bit浮点加法、乘法、乘加运算

    32

    32bit浮点倒数、倒数平方根、__log2f、exp2f、__sinf、__cosf

    16

    最大最小

    64

    32bit整型位取反

    16

    32-bit AND, OR, XOR

    64

  • 在不影响最后结果的情况下使用低精度类型,使用内部计算接口。

    比如,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位,但速度快很多。

    更多内部函数列表请参考《CUDA C++ Programming Guide》。

  • 在不影响最终结果的情况下用精度换取速度,使用单精度float而不是双精度double。
  • 尽量减少流程控制指令,即减少条件分支。比如:
    1. 控制条件仅取决于(threadIdx / warpSize)。
    2. 使用#pragma unroll展开循环。
  • 尽量使用__fdividef(x,y)单精度浮点除法,其运算效率优于除法运算符。
  • 大部分情况下,rsqrtf()效率比1.0 / sqrtf()高。
  • 使用三角函数时,尽量使用单精度,且x值域较小为宜。
  • 尽量使用位运算代替整数除法和模运算。
  • 半精度计算,使用half2代替half。
  • 调用接口时,尽量保持参数类型一致,否则会有类型转换的消耗。比如函数参数是int类型,输入是char类型或者short类型,需要先将类型转化为int类型。
  • 寄存器优化,比如GPU的位操作微观优化。
搜索结果
找到“0”个结果

当前产品无相关内容

未找到相关内容,请尝试其他搜索词