最大化指令吞吐量

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

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

  算术指令及其吞吐量之间的关系请参考《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的位操作微观优化。