CUDA学习(六十七)

简介:

数学函数:
参考手册列表及其说明列出了设备代码支持的C / C ++标准库数学函数的所有功能,以及所有固有功能(仅在设备代码中支持)
标准函数:
本节中的功能可用于主机和设备代码。
本节规定了每个功能在设备上执行时的错误界限,以及在主机不提供功能的情况下在主机上执行时的错误界限。
加法和乘法符合IEEE标准,因此最大误差为0.5 ulp。
将单精度浮点操作数四舍五入为整数,其结果为单精度浮点数的推荐方法是rintf(),而不是roundf()。 原因是roundf()映射到设备上的8指令序列,而rintf()映射到单个指令。 truncf(),ceilf()和floorf()也分别映射到单个指令。
具有最大ULP误差的单精度数学标准库函数
最大误差表示为正确舍入的单精度结果与CUDA库函数返回的结果之间的差值的绝对值。
1
2
3
4
双精度浮点函数:
将双精度浮点操作数舍入为一个整数,其结果是一个双精度浮点数的推荐方法是rint(),而不是round()。 原因是round()映射到设备上的8个指令序列,而rint()映射到单个指令。 trunc(),ceil()和floor()也分别映射到单个指令。
具有最大ULP误差的双精度数学标准库函数:
最大误差表示为正确舍入的双精度结果与CUDA库函数返回的结果之间的差值的绝对值。
5
6
7
8
内在函数:
在这些函数中,标准函数的某些函数的准确性较低,但速度更快。它们具有以__为前缀的相同名称(如__sinf(x))。 它们映射到更少的原生指令时速度更快。 编译器有一个选项(-use_fast_math),它强制表8中的每个函数编译为其内部对应部分。 除了降低受影响功能的准确性之外,还可能会在特殊情况下处理一些差异。 更稳健的方法是通过调用内部函数来选择性地替换数学函数调用,只有在性能增益的情况下才适用数学函数调用,并且可以容忍更改的属性(如精度降低和不同的特殊情况处理)。
受use_fast_math影响的函数:
9
以_rn结尾的函数使用该轮以最接近的舍入模式运行。
使用_rz后缀的函数使用舍入为零的舍入模式进行操作。
以_ru结尾的函数使用舍入(到正无穷)舍入模式进行操作。
以_rd结尾的函数使用向下舍入(到负向无限)舍入模式进行操作
单精度浮点函数:
_fadd [rn,rz,ru,rd]()和__fmul_ [rn,rz,ru,rd]()映射到编译器不会合并到FMAD中的加法和乘法操作。 相比之下,'*'和'+'运算符产生的加法和乘法运算常常被合并到FMAD中。
浮点除法的精度取决于代码是使用-prec-div = false还是-prec-div = true编译的。 当使用-prec-div = false编译代码时,常规除法/运算符和__fdividef(x,y)具有相同的准确性,但对于2126 __fdividef(x,y)将结果为零 ,而/运算符将正确结果提供给表9中所述的精度。同样,对于2126 __fdividef(x,y)将传递NaN(由于将无穷大乘以零 ),而/运算符返回无穷大。 另一方面,当代码使用-precdiv = true或完全没有任何-prec-div选项编译时,/ operator是符合IEEE的,因为它的默认值为true。
单精度浮点内部函数:
CUDA运行时库支持各自的错误界限
10
11
双精度浮点函数:
__dadd_rn()__dmul_rn()映射到编译器不会合并到FMAD中的加法和乘法操作。 相比之下,'*'和'+'运算符产生的加法和乘法运算常常被合并到FMAD中。
双精度浮点内部函数:
12
timg

目录
相关文章
|
并行计算 C语言 编译器
|
并行计算 API 调度
CUDA学习(八十八)
3.虽然__syncthreads()一直被记录为同步线程块中的所有线程,但Pascal和以前的体系结构只能在warp级别强制执行同步。 在某些情况下,只要每条经线中至少有一条线达到屏障,就可以在不被每条线执行的情况下成功实现屏障。
1697 0
|
存储 并行计算 C语言
|
并行计算 编译器
|
并行计算 API 编译器
CUDA学习(六十五)
很早之前就发现云栖社区的编辑器有一个Bug,往草稿箱存博客,当草稿箱博客数超过十篇时,无法再选择十篇前的博客进行编辑
2366 0
|
并行计算 C语言 编译器
|
并行计算 API 异构计算
|
并行计算 API C语言
|
并行计算 API

热门文章

最新文章