CUDA 性能优化指北（四）- Maximize Instruction Throughput

为了最大化指令吞吐量，应用程序应该:

• 尽量减少使用低吞吐量的算术指令；这包括在不影响最终结果的情况下以精度换取速度，例如使用 intrinsic 函数而不是regular functions，单精度而不是双精度，或将非规格化的数字刷新为零;
• 最小化由控制流指令引起的 divergent warps
• 减少指令的数量，例如，尽可能优化同步点，或使用限制性指针，如__limitt__ 所述。

在本节中，吞吐量以每个多处理器每个时钟周期的操作数表示。对于32的warp大小，一条指令对应32个操作，因此如果N是每个时钟周期的操作数，则指令吞吐量为每个时钟周期的N/32条指令。

所有吞吐量都是针对一个多处理器的。它们必须乘以设备中的多处理器数量，才能获得整个设备的吞吐量。

# Arithmetic Instructions

Table 3 给出各种计算能力设备的硬件中native支持的算术指令的吞吐量。

其他指令和函数是在native指令之上实现的。对于具有不同计算能力的设备，实现可能不同，编译后的本机指令数量可能随每个编译器版本而变化。对于复杂的函数，根据输入可以有多个代码路径。cuobjdump可用于检查cubin 对象中的特定实现。

CUDA头文件(math_functions.h, device_functions.h，…)中可以很容易地一些函数的实现。

通常，使用-ftz=true编译的代码(反规范化的数字被刷新为零)往往比使用-ftz=false编译的代码具有更高的性能。类似地，使用-precc-div=false(较不精确的除法)编译的代码，其性能往往高于使用-precc-div=true编译的代码，而使用-precc-sqrt=false(较不精确的平方根)编译的代码，其性能往往高于使用-precc-sqrt=true编译的代码。nvcc用户手册更详细地描述了这些编译标志。

# Single-Precision Floating-Point Division

__fdividef(x, y)(参见内在函数)提供了比除法运算符更快的单精度浮点除法。

# Single-Precision Floating-Point Reciprocal Square Root

为了保持IEEE-754语义，编译器只能在倒数和平方根都是近似的情况下，将1.0/sqrtf()优化为rsqrtf()(即-prec-div=false和-prec-sqrt=false)。因此，建议在需要的地方直接调用 rsqrtf()。

# Single-Precision Floating-Point Square Root

单精度浮点平方根被实现为倒数平方根和倒数平方根，而不是倒数平方根和乘法，因此它对0和无穷大给出正确的结果。

# Sine and Cosine

sinf (x)， cosf(x)， tanf(x)， sincosf(x)，以及相应的双精度指令都要昂贵得多，如果参数x的幅度很大，就更贵了。

更准确地说，参数缩减代码(参见实现的数学函数)包含两个代码路径，分别称为快速路径和慢路径。

快速路径用于大小足够小的参数，基本上由一些乘法运算组成。慢路径用于大小较大的参数，包括在整个参数范围内获得正确结果所需的冗长计算。

目前三角函数的参数约简代码对于单精度函数的参数大小小于105615.0f，对于双精度函数的参数大小小于2147483648.0，选择快速路径。

由于慢路径比快路径需要更多的寄存器，尝试通过在本地内存中存储一些中间变量来减少慢路径中的寄存器压力，这可能会影响性能，因为本地内存的高延迟和带宽(参见设备内存访问)。目前，单精度函数使用28字节的本地内存，双精度函数使用44字节的本地内存。不过，具体金额可能会有所变化。

由于长时间的计算和在慢路径中使用本地内存，当需要慢路径约简时，这些三角函数的吞吐量比快速路径约简低一个数量级。

# Integer Arithmetic

整数除法和取模运算的成本很高，因为它们最多可编译20条指令。在某些情况下，它们可以用位运算替换:如果n是2的幂，(i/n)等价于(i>>log2(n))和(i%n)等价于(i&(n-1));如果n是字面量，编译器将执行这些转换。

__brev和__popc映射到单个指令，而__brevll和__popcll映射到几个指令。

# Half Precision Arithmetic

为了实现16位精度浮点加法、乘法或乘法加法的良好性能，建议使用half2数据类型实现半精度，使用__nv_bfloat162实现__nv_bfloat16精度。Vector intrinsic(例如，__hadd2， __hsub2， __hmul2， __hfma2)可以用来在一条指令中执行两个操作。使用half2或__nv_bfloat162来代替使用half或__nv_bfloat16的两个调用也可能有助于其他intrinsic的性能，例如warp shuffle。

提供了固有的__halves2half2来将两个半精度值转换为half2数据类型。

提供了固有的__halves2bfloat162来将两个__nv_bfloat精度值转换为__nv_bfloat162数据类型。

# Type Conversion

有时，编译器必须插入转换指令，引入额外的执行周期。这是以下情况的情况:

• 操作于char或short类型变量的函数，其操作数通常需要转换为int，
• 用作单精度浮点计算的输入(由C/ c++标准强制要求)的双精度浮点常量(即不带任何类型后缀定义的那些常量)。

最后一种情况可以通过使用单精度浮点常量来避免，该常量使用f后缀定义，如3.141592653589793f, 1.0f, 0.5f。

# Control Flow Instructions

任何流控制指令(if、switch、do、for、while)都可以通过导致warp中的线程发散(即遵循不同的执行路径)而显著地影响有效指令吞吐量。如果发生了这种情况，就必须序列化不同的执行路径，从而增加为这种扭曲而执行的指令总数。

为了在控制流依赖于线程ID的情况下获得最佳性能，控制条件应该被写下来，以便最小化divergent warps数。这是可能的，因为正如SIMT体系结构中提到的，整个块的warps分布是确定的。一个简单的例子是，当控制条件只依赖于(threadIdx / warpSize)时，其中warpSize是warp大小。在这种情况下，由于控制条件与warp完全对齐，所以没有warp diverges。

有时，编译器可能展开循环，或者通过使用分支预测来优化短if或switch块。在这种情况下，任何warp都不可能发散。 程序员还可以使用#pragma unroll指令控制循环展开(参见#pragma unroll)。

# Synchronization Instruction

对于计算能力为3.x的设备，__syncthreads()的吞吐量是每个时钟周期128个操作。计算能力为6.0的设备每个时钟周期32次操作，计算能力为7.x和8.x的设备每个时钟周期16次操作。对于具有计算能力5.x、6.1和6.2的设备，每个时钟周期64个操作。

注意，__syncthreads()可以通过强制多处理器空闲来影响性能，详见设备内存访问。

# Minimize Memory Thrashing

经常分配和释放内存的应用程序可能会发现，分配调用往往会随着时间的推移而变慢，直到达到一个限制。这通常是预期的，因为将内存释放回操作系统供其自己使用。为了在这方面取得最佳表现，我们建议如下:

• 试着对手头的问题进行分配。不要尝试使用cudaMalloc / cudaMallocHost / cuMemCreate 分配所有可用内存，因为这将强制内存立即驻留，并阻止其他应用程序能够使用该内存。这可能会给操作系统调度程序带来更大的压力，或者只是阻止使用相同GPU的其他应用程序完全运行。
• 尝试在应用程序早期以适当大小的分配方式分配内存，并仅在应用程序不需要内存时才分配内存。减少应用程序中cudaMalloc+cudaFree调用的数量，特别是在性能关键区域。
• 如果应用程序不能分配足够的设备内存，可以考虑使用其他内存类型，如cudaMallocHost或cudaMallocManaged，这可能不太好，但会使应用程序取得进展。
• 对于支持该特性的平台，cudaMallocManaged允许超订阅，并且启用了正确的cudaMemAdvise策略，将允许应用程序保留cudaMalloc的大部分(如果不是全部)性能。cudaMallocManaged也不会强制一个分配驻留在需要它或预取它之前，这减少了操作系统调度程序的总体压力，并更好地支持多原则用例。

# 参考文献

• https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.htm
• https://on-demand.gputechconf.com/gtc/2018/presentation/s81006-volta-architecture-and-performance-optimization.pdf