[翻译]CUDA-C-Programming-Guide Maximize InstructionThroughput

5.4 最大化指令吞吐量

为了达到最大的指令吞吐量，程序应该：

• 最小化使用低吞吐量的计算指令；有以下方法：在不影响结果的情况下以精度换取执行速度，比如使用指令来代替内置函数，用单精度代替双浮点精度，或者将非归一化数据刷新为0.
• 最小化分支
• 减少指令的数量，

本章节中，吞吐量通过每个SM每个时钟周期进行的运算次数来表示。对于一个warp=32，一条指令对应了32次运算，所有如果一个时钟周期进行了N次运算，那么指令吞吐量就是每个时钟周期N/32个指令。

吞吐量是针对一个SM的，将其乘以SM的数量就得到整个设备的吞吐量。

5.4.1 计算指令

下表给出了各种不同设备天然支持的计算吞吐量

   

两外还有些指令和函数在天然指令的顶端实现。对于不同计算能力设备函数实现可能都不相同，不同的编译器版本，生成的指令个数也不相同。对于一些复杂点的函数，根据输入不同会有不同的分支代码。可以用cuobjdump来检查cubin 对象更具体的实现。

在头文件中可以查看这些函数的实现(math_functions.h,device_function.h,…)

一般的，带了-ftz=true（非正规化数据刷新为0）比带-ftz=false 编译选项编译出来的代码有更高的性能。加了-prec div=false（更低的精度）比-prec div=true编译选项编出来的代码有更高的性能。加了-prec-sqrt=false（更低精度的平方根）比-prec-sqrt=true编译选项生成的代码有更高的性能。更多详细的信息可以查看nvcc的用户手册。

单精度浮点除法

__fdividef(x,y)提供了比一般除法运算更高性能的单精度浮点数除法。

单精度浮点对数平方根

在符合IEEE-754的条件下，当对数和平方根近似时，编译器将1.0/sqrtf()优化成rsqrtf()，所以在需要的地方推荐直接调用rsqrtf()。

单精度浮点数平方根

单精度平方根被实现为平方根的倒数的倒数，而不是平方根的倒数后接乘法，因为这样可以得到0或者无穷大的结果。

正玄和余玄函数

Sinf(x),cosf(x),tanf(x),sinconsf(x)和对应的双精度浮点数函数在参数x很大时性能都比较低。

更准确的说，输入的参数决定了函数内部代码的走向：快速的实现还是慢速的实现。

当输入参数小时，运算只是一些简单的乘加运算，函数会执行快速的实现。当输入参数大，就会包括一些很冗长的运算，函数为了得到准确的结果就会执行慢速的实现。

目前来讲，当参数小于105615.0f的单精度时函数执行快速实现；小于2147483648.0的双精度时，函数执行慢速的实现。

因为慢速的实现会比快速的实现使用更多的寄存器，并且尝试减少寄存器的压力，在慢速实现代码中，一些中间临时变量都存在局部存储器，而局部存储器延迟高，带宽低。当前的实现时局部存储器的28字节被单精度函数使用，双精度函数使用局部存储器的44字节。

整数运算

整数运算和模运算是非常耗时的，因为他们需要多达20条指令。在有些情况下，他们可以用位运算代替：如果n是2的次幂，那么1/n等于i>>log2(n)，I % n 等于(i&(n-1));

半精度运算

为了获得高的半精度单浮点数加，乘法或者乘加的吞吐量，推荐使用half2数据类型。向量指令如__hadd2,__hsub2,__hmul2,__hfma2 能够在单条指令下执行两次运算。使用half2代替half可以提高性能。

__halves2half2指令可以将两个单精度值转换成half2数据类型。

类型转换

有些时候，编译器必须插入一些转换指令，从而引入额外的执行时钟周期。这些情况包括：

• char或short 运算，转成int
• 双精度浮点数常量(没有使用类型后缀)传入参数为单精度浮点数的函数。这种情况可以在定义的时候避免，定义时增加f作为后缀，比如1.0f,0.5f

5.4.2 流程控制指令

由于流程控制指令(if,switch,do,for,while)会导致一个warp中的线程执行不同的路径，所有流程控制指令会严重的影响指令吞吐量。如果这种情况发生，不同的执行路径必须顺序执行，这样就增加了执行的指令数量。

对于流程控制依赖于线程ID的情况，可以通过设计控制条件使得有分叉的warp数量最小，从而获得好的性能。举个简单的例子：如果将控制条件设计成只依赖于threadIdx/warpSize，这样的化，没有分叉的warp，因为控制条件完美的和warp对齐。

有些时候，编译器会展开循环，或者通过分支预判来优化短的if或者switch语句块。编程人员可以使用#pragma unroll语句来控制展开循环。

当使用分支预判时，所有分支的所有指令都不会被跳过，相反的，他们每一个都会附带一个根据控制条件设置成true或false的条件代码，虽然他们都会执行，但是只有设置成true的指令会真正执行，而那些设置成false的指令不会写结果，并且不会计算地址或者读取操作数。

5.4.3 同步指令

__syncthreads()在计算能力为3.x的设备上是每时钟周期128次，6.0上是32每时钟周期32次，7.x上是每个时钟周期16次，5.x，6.1，6.2上每个时钟周期64次。