__saturatef() 内在函数没有等效的双精度

Posted

技术标签:

【中文标题】__saturatef() 内在函数没有等效的双精度【英文标题】:__saturatef() intrinsic has no double-precision equivalent 【发布时间】:2017-10-14 15:43:40 【问题描述】:

Cuda 支持intrinsic functions。一些映射到设备指令,如融合乘加,不能用正常语法表示。其他的近似值应该比“标准”函数更快(尽管可能不太准确)。

似乎后一种类型的内在函数不支持双精度参数,仅支持单精度浮点数。有道理:如果您使用双打,那么您肯定对准确性感兴趣而不是速度。

然而,令我惊讶的是,__saturatef() 内部函数将其参数限制在 0 和 1 之间,并没有双精度版本。

如果我在幼稚的实现上使用 __saturatef() 内在函数,是否有任何潜在的数据丢失?如果是这样,有人知道这个内在函数是如何工作的吗?如果不是,为什么 nvidia 会忽略双精度版本?

【问题讨论】:

“潜在的数据丢失”是什么意思?能给我举个例子吗? CUDA 的设备内在函数的主要目的是公开一些其他方式无法访问的硬件功能。由于相关用例,硬件为单精度算术提供饱和,它为双精度算术提供饱和。原样的问题是边界题外话。我建议改写这个问题,这样它显然是关于编程的。例如:假设没有双精度内在 __saturate(double),那么模拟这个操作的最快方法是什么? 【参考方案1】:

CUDA 的设备函数内在函数的目的是公开以其他方式无法访问的特定硬件功能,例如一些代数和超越函数的快速逼近,或者在 __saturatef() 的情况下钳制到区间 [0,1],它映射到 GPU 的机器指令 F2F.FTZ.F32.F32.SAT 用于 5.0 之前的计算能力(sm_50)FADD.SAT适用于大于或等于计算能力 5.0 的架构。

查看PTX documentation 表明浮点饱和支持半精度 (.f16) 和单精度 (.f32) 运算,但不支持双精度 (.f64) 运算.因此,提供的内在函数的非正交性是由 GPU 硬件中的非正交性引起的。由于相关用例(包括图形)通常不使用双精度,因此为较低精度提供了饱和度。

快速详尽的测试表明__saturatef (float) 的行为与fminf (fmaxf (float, 0.0f), 1.0f) 相同。查看反汇编代码,CUDA 编译器似乎也将这个 min / max 习惯用法与 __saturatef() 的使用相同,作为优化。因此,可以很容易地创建一个双精度等价物

double my_saturate (double a)

    return fmin (fmax (a, 0.0), 1.0);

【讨论】:

以上是关于__saturatef() 内在函数没有等效的双精度的主要内容,如果未能解决你的问题,请参考以下文章

ARM64 固有的 x86_64 点向量积

SSE 比较内在 - 如何从比较中获得 1 或 0?

_mm256_loadu2_m128i 内在函数在 g++ 下不可用?

C 内在函数、SSE2 点积和 gcc -O3 生成的程序集

Python 中的函数的双下划线“__”

等效于嵌入式 C 中的 NOP?