---
title: 最大化指令吞吐量
description: "算术强度本质上是指访问内存所花的时间内做的数学运算次数，所以要最大化算术强度，可参考如下优化思路："
url: https://www.hikunpeng.com/document/detail/zh/perftuning/gputuing-wp/kunpenggpu_19_0028.html
sourcePath: /source/zh/perftuning/gputuing-wp/kunpenggpu_19_0028.html
indexId: ea33fed57b4137b9225c49e06761b7eac0f1dab25f162f7f7daa35350549c62957
---
# 最大化指令吞吐量

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

- 尽量少使用低吞吐量的算术指令。
  算术指令及其吞吐量之间的关系请参考《CUDA C++ Programming Guide(https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#arithmetic-instructions)》。

  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(https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#intrinsic-functions)》。


- 在不影响最终结果的情况下用精度换取速度，使用单精度float而不是双精度double。
- 尽量减少流程控制指令，即减少条件分支。比如：

  1. 控制条件仅取决于（threadIdx / warpSize）。
  2. 使用#pragma unroll展开循环。
- 尽量使用__fdividef（x，y）单精度浮点除法，其运算效率优于除法运算符。
- 大部分情况下，rsqrtf()效率比1.0 / sqrtf()高。
- 使用三角函数时，尽量使用单精度，且x值域较小为宜。
- 尽量使用位运算代替整数除法和模运算。
- 半精度计算，使用half2代替half。
- 调用接口时，尽量保持参数类型一致，否则会有类型转换的消耗。比如函数参数是int类型，输入是char类型或者short类型，需要先将类型转化为int类型。
- 寄存器优化，比如GPU的位操作微观优化。