1

CUDA C 的数学函数实现 ( cuda/math_function.h)acosf包含以下段落：

if (__float_as_int(a) < 0) {
  t1 = CUDART_PI_F - t1;
}

其中a和t1是floats和CUDART_PI_F是float先前设置为接近数学常数 Pi 的数值。我试图了解条件（if-clause）正在测试什么以及它的 C 等价物或 function/macro 是什么__float_as_int(a)。我搜索了__float_as_int()但没有成功的实现。这似乎__float_as_int()是 NVIDIA NVCC 的内置宏或函数。查看 NVCC 从上述段落中产生的 PTX：

    .reg .u32 %r<4>;
    .reg .f32 %f<46>;
    .reg .pred %p<4>;
    // ...
    mov.b32         %r1, %f1;
    mov.s32         %r2, 0;
    setp.lt.s32     %p2, %r1, %r2;
    selp.f32        %f44, %f43, %f41, %p2;

很明显，这不是__float_as_int()四舍五入。（这将产生一个。）相反，它将作为位副本（）分配给（注意：是类型（无符号整数）！！），然后将其作为（有符号整数，混淆！！）与（谁的价值是）。floatintcvt.s32.f32float %f1b32%r1%r1u32%r1s32%r20

对我来说，这看起来有点奇怪。但显然它是正确的。

有人可以解释发生了什么，特别是解释__float_as_int()在 if 子句测试是否定的（<0）的上下文中正在做的事情吗？.. 并提供与 if 子句和/或__float_as_int()marco 等效的 C 语言？

4

1 回答 1

4

__float_as_int重新解释float为int. int是<0它具有最高有效位的时候。因为float它也意味着符号位打开，但并不完全意味着数字是负数（例如，它可以是“负零”）。float检查然后检查 if is会更快< 0.0。

C 函数可能如下所示：

int __float_as_int(float in) {
     union fi { int i; float f; } conv;
     conv.f = in;
     return conv.i;
}

在其他版本中__cuda___signbitf使用此标头。

于 2012-12-10T13:48:34.167 回答