最大化指令吞吐量:
为了最大化指令吞吐量,应用程序应
- 尽量减少低吞吐量的算术指令的使用; 这包括在不影响最终结果的情况下交易的速度精度,例如使用内部函数而不是常规函数(内部函数在内部函数中列出),单精度而不是双精度,或者将非正则化数字清零。
- 最小化由控制流程指令引起的分歧warp
- 减少指令的数量,例如,如同步指令中所述,尽可能优化出同步点,或使用限制指针(如__restrict__中所述)。
在本节中,吞吐量以每个多处理器每个时钟周期的操作次数给出。 对于32的变形大小,一条指令对应于32个操作,所以如果N是每个时钟周期的操作数,则指令吞吐量是每个时钟周期N / 32条指令。
所有的吞吐量都是针对一个多处理器的。 它们必须乘以设备中的多处理器数量才能获得整个设备的吞吐量。
算术指令:
表2给出了各种计算能力的设备硬件本地支持的算术指令的吞吐量
表2本地算术指令的吞吐量(每个多处理器每个时钟周期的结果数量)
其他指令和功能在本机指令的基础上实现。 对于具有不同计算能力的设备,实现可能不同,并且编译后的本机指令的数量可能随着每个编译器版本而波动。 对于复杂的函数,根据输入可能有多个代码路径。 cuobjdump可用于检查cubin对象中的特定实现。
CUDA头文件(math_functions.h,device_functions.h,...)中提供了一些函数的实现。
通常,使用-ftz = true编译的代码(非规格化数字被刷新为零)与使用-ftz = false编译的代码相比具有更高的性能。 类似地,用-prec div = false编译的代码(不太精确的分割)往往比用-prec div = true编译的代码具有更高的性能代码,而用-prec-sqrt = false编译的代码(不太精确的平方根)往往具有 比使用-prec-sqrt = true编译的代码更高的性能。 nvcc用户手册更详细地描述了这些编译标志。
单精度浮点除法:
__fdividef(x,y)(请参阅内部函数)提供比除法运算符更快的单精度浮点除法。
单精度浮点平方倒数:
为了保留IEEE-754语义,只有当倒数和平方根都近似时(即,-precdiv = false和-prec sqrt = false),编译器才能优化1.0 / sqrtf()到rsqrtf()。 因此建议直接在需要的地方调用rsqrtf()
单精度浮点平方根:
单精度浮点平方根实现为倒数平方根,后跟倒数,而不是倒数平方根,然后乘以,从而得到0和无穷大的正确结果。
正弦和余弦:
sinf(x),cosf(x),tanf(x),sincosf(x)以及相应的双精度指令会更加昂贵,而且如果参数x的幅度很大,则更是如此。
更确切地说,参数缩减代码(参见用于实现的数学函数)包括分别被称为快速路径和慢速路径的两个代码路径。
快速路径用于幅度足够小的参数,基本上由几个乘加操作组成。 慢速路径用于幅值较大的参数,并且包含为在整个参数范围内获得正确结果所需的冗长计算。
目前,三角函数的参数缩减代码为单精度函数选择大小小于105615.0f的参数的快速路径,并且为双精度函数选择小于2147483648.0的参数。
由于慢速路径比快速路径需要更多寄存器,因此尝试通过在本地存储器中存储一些中间变量来降低慢速路径中的寄存器压力,由于本地存储器的高延迟和带宽,这可能会影响性能(请参阅设备内存访问)。 目前,单精度函数使用28个字节的本地存储器,双精度函数使用44个字节。 但是,具体金额可能会发生变化。
由于慢速路径中的本地存储器的冗长计算和使用,当需要慢速路径减少时,这些三角函数的吞吐量降低一个数量级,而不是快速路径减少。
整数运算:
整数除法和模运算代价高达20个指令。 在某些情况下,它们可以用逐位运算代替:如果n是2的幂,则(i / n)等于(i >> log2(n))并且(i%n)等于(i&1)); 如果n是文字,编译器将执行这些转换。
__brev和__popc映射到单个指令,而__brevll和__popcll映射到几条指令。
__ [u] mul24是传统的内在函数,不再有任何理由被使用。
半精度算法:
为了达到很好的半精度浮点加法,乘法或乘加吞吐量,建议使用half2数据类型。 然后可以使用向量内部函数(例如__hadd2,__hsub2,__hmul2,__hfma2)在单个指令中执行两个操作。 使用half2来代替使用一半的两个调用也可以帮助其他内部函数的性能,例如warp shuffles。
提供内在的__halves2half2来将两个半精度值转换为half2数据类型。
类型转换:
有时,编译器必须插入转换指令,引入额外的执行周期。 这是:
- 函数操作类型为char或short的操作数通常需要转换为int的变量,
- 用作单精度浮点计算的输入的双精度浮点常量(即那些没有任何类型后缀定义的常量)(按照C / C ++标准的规定)
最后一种情况可以通过使用单精度浮点常量来避免,用f后缀定义,如3.141592653589793f,1.0f,0.5f
控制流程指令:
任何流量控制指令(if,switch,do,for,while)都可以通过使同一个warp的线程发散(即遵循不同的执行路径)而显着影响有效的指令吞吐量。 如果发生这种情况,则必须对不同的执行路径进行序列化,从而增加为这个warp执行的指令总数。
为了在控制流程依赖于线程ID的情况下获得最佳性能,应该写入控制条件以使发散的warp数量最小化。 这是可能的,因为在SIMT体系结构中所提到的在整个块上的warp分布是确定性的。 一个简单的例子是当控制条件只取决于(threadIdx / warpSize)warpSize是warp的大小。 在这种情况下,由于控制条件与warp完美对齐,因此没有warp分叉。
有时,编译器可能会展开循环,或者可能会使用分支预测来优化short或if块,如下所述。 在这些情况下,没有任何扭曲可以分歧。 程序员还可以使用#pragma unroll指令控制循环展开(请参阅#pragma unroll)。
当使用分支预测时,跳过执行取决于控制条件的指令。 相反,它们中的每一个都与根据控制条件被设置为真或假的预读条件代码或谓词相关联,并且尽管这些指令中的每一个被计划执行,但是只有具有真谓词的指令才被实际执行。 带有假谓词的指令不会写结果,也不会计算地址或读操作数。
同步指令:
对于计算能力3.x,每个时钟周期的__syncthreads()的吞吐量为128个操作;对于计算能力为6.0和7.0的设备,每个时钟周期为32个操作;对于计算能力为5.x的设备,每个时钟周期为64个操作6.2.
请注意,通过强制多处理器空闲,__syncthreads()可以影响性能,详见设备内存访问。