CUDA实践指南（二十六）

数学库：
中等优先级：每当速度超过精度时使用快速数学库。
支持两种类型的运行时数学运算。 它们可以用它们的名字加以区分：一些名称带有前置下划线，而另一些则没有（例如，__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];

如果在等待全局内存访问完成时可以发出足够的独立算术指令，那么线程调度程序可以隐藏大部分全局内存延迟。 但是，最好尽可能避免访问全局内存。