数学库:
中等优先级:每当速度超过精度时使用快速数学库。
支持两种类型的运行时数学运算。 它们可以用它们的名字加以区分:一些名称带有前置下划线,而另一些则没有(例如,__functionName()
和functionName()
)。 遵循__functionName()
命名约定的函数直接映射到硬件级别。 它们速度更快,但准确度稍低(如__sinf(x)
和__expf(x)
)。 functionName()命名约定的函数较慢但具有较高的准确性(例如,sinf(x)和expf(x))。 __sinf(x)
,__cosf(x)
和__expf(x)
的吞吐量远远大于sinf(x),cosf(x)和expf(x)的吞吐量。 如果参数x的大小需要减小,后者变得更加昂贵(大约慢一个数量级)。 此外,在这种情况下,自变量缩减代码使用本地内存,由于本地内存的高延迟,这会更多地影响性能。 更多详细信息可在CUDA C编程指南中找到.
还要注意,无论何时计算相同参数的正弦和余弦,sincos系列指令都应该用于优化性能:
- __sincosf()用于单精度快速数学(参见下一段)
- sincosf()用于常规的单精度
- sincos()用于双精度
nvcc的-use_fast_math编译器选项强制每个functionName()调用等效的__functionName()
调用。 只要准确度优先于性能,就应该使用此开关。 超验函数通常是这种情况。 请注意,此开关仅在单精度浮点上有效。
中等优先:如果可能,优先选择速度更快,更专业化的数学函数,而不是更慢,更一般的数学函数。
对于小整数幂(例如$x^2$或$x^3$),显式乘法几乎肯定比使用pow()等一般指数运算例程更快。 虽然编译器优化改进不断寻求缩小这种差距,但明确的乘法(或使用等效的专用内联函数或宏)可以具有显着的优势。 当需要几个相同基数的幂(例如,其中$x^2$和$x^5$都靠近计算)时,这个优点增加,因为这有助于编译器进行其常见的子表达式消除(CSE)优化。
对于使用基数2或10的指数,使用函数exp2()或expf2()和exp10()或expf10()而不是函数pow()或powf()。 pow()和powf()在寄存器压力和指令计数方面都是重量级的函数,这是由于在大多数幂运算中产生了许多特殊情况,并且难以在整个基数和指数的整个范围内实现良好的精度。 另一方面,函数exp2(),exp2f(),exp10()和exp10f()在性能方面与exp()和expf()类似,并且可以比它们快十倍 pow()/ powf()等价物。
对于指数为1/3的指数,使用cbrt()或cbrtf()函数而不是泛指数函数pow()或powf(),因为前者明显快于后者。 同样,对于指数为-1/3的指数,使用rcbrt()或rcbrtf()。
用sincospi()替换sinpi(),cos(π )与cospi()和sincos(π )的sin(π )。 就精度和性能而言这是有利的。 作为一个特殊的例子,为了评估正弦函数而不是弧度,使用sinpi(x / 180.0)。 类似地,当函数参数的形式为π 时,单精度函数sinpif(),cospif()和sincospif()应该替换对sinf(),cosf()和sincosf()的调用。 (性能优势sinpi()超过sin()是由于简化了参数约简;精确度优势是因为sinpi()仅隐式乘以π,有效地使用无限精确的数学π而不是单精度或双精度逼近。)
精度相关的编译器标志:
默认情况下,nvcc编译器为计算能力2.x的设备生成符合IEEE标准的代码,但它也提供了生成代码的选项,这些代码的精确度稍低但速度更快,并且更接近为早期设备生成的代码:
- -ftz = true(非规格化数字被刷新为零)
- -prec-div = false(不太精确的划分)
- -prec-sqrt = false(不精确的平方根)
另一个更积极的选项是-use_fast_math,它强制每个functionName()调用等效的__functionName()
调用。 这使得代码运行速度更快,代价是精度和准确性降低。
内存指令:
高优先级:尽量减少全局内存的使用。 尽可能优先使用共享内存访问。
内存指令包括从共享,本地或全局内存读取或写入的任何指令。 访问未缓存的本地或全局内存时,存在400至600个时钟周期的内存延迟。 例如,以下示例代码中的赋值运算符具有较高的吞吐量,但关键的是,从全局内存中读取数据的延迟时间为400至600个时钟周期:
__shared__ float shared[32];
__device__ float device[32];
shared[threadIdx.x] = device[threadIdx.x];
如果在等待全局内存访问完成时可以发出足够的独立算术指令,那么线程调度程序可以隐藏大部分全局内存延迟。 但是,最好尽可能避免访问全局内存。