cuda - -use-fast-math 选项是否将 SP 乘法转换为内在函数?

标签 cuda nvcc fast-math

我快速浏览了 CUDA 编程指南 w.r.t -use-fast-math optimizations,虽然附录 C 提到了要转换为内在的除法,但没有提到乘法。我问这个问题的原因是,我的内核有很多乘法。我知道 NVCC 会尝试融合乘法和加法(当使用常规的“*”和“+”运算符时,并且内在函数永远不会合并到 FMAD 操作中)。但是,如果我的代码是大量的乘法,那么如果舍入像 __fmul_rn 这样的 SP 内在函数,会有好处吗?用来?

所以有两个问题:

  • -use-fast-math 选项是否将带有“*”运算符的乘法转换为像 __fmul_rn 这样的 SP 内在函数?
  • 手动编码乘法以显式使用 __fmul_rn 是否有性能优势?一个例子或一些数字会帮助我理解。
  • 最佳答案

    “独立”单精度乘法始终编译为硬件指令(“内部指令”)。没有其他类型的浮点乘法指令。 nvcc 中的 -use_fast_math 选项对为计算能力 1.x 目标发出的浮点乘法指令没有影响。在计算 2.x 和 3.x 目标上,它将编译器置于兼容模式并且所有单精度乘法指令将为 mul.ftz.f32 (清零)。

    您提到的浮点内部函数( __fmul_{rm,rn,rp,rz,ftz,sat} )仅提供对 IEEE 舍入行为的显式控制。我不相信它们在 Fermi 或 Kepler GPU 上的吞吐量存在差异。

    关于cuda - -use-fast-math 选项是否将 SP 乘法转换为内在函数?,我们在Stack Overflow上找到一个类似的问题: https://stackoverflow.com/questions/11507440/

    相关文章:

    c++ - CUDA 写入其他 warp 看不到的全局内存

    c++ - 在编译错误时制作 nvcc 输出跟踪

    cuda - 我可以在编译时确定是否设置了 --use_fast_math 吗?

    c++ - 您可以在第二个声明中使用 `= delete` 模板函数吗?

    visual-studio-2008 - 使用仿真模式在nvcc中出现错误

    c - 是否有任何浮点密集型代码在任何基于 x86 的架构中产生位精确的结果?

    cuda - 如何使用 CUDA 在 GPU 上运行 "host"函数?

    使用 cos 时,Cuda 从 __device__ 函数返回错误值

    c# - 返回 CUDA 中数组的最小和最大元素