acosf 实现中的 CUDA __float_as_int

Question

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

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

哪里a和 t1是 floats和 CUDART_PI_F是 float预先设置为接近数学常数 Pi 的数值。 我试图了解条件(if 子句)正在测试什么以及它的 C 等效项或函数/宏 __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()不是 float至 int四舍五入。 (这会产生一个 cvt.s32.f32 。)相反，它分配了 float %f1。作为位复制 ( b32 ) 到 %r1 (注意:%r1 的类型是u32 (unsigned int)!!)然后比较%r1就好像它是一个 s32 (有符号整数，令人困惑!!)与 %r2 (谁的值为 0 )。

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

谁能解释一下发生了什么，特别是解释一下 __float_as_int()是在 if 子句测试是否定的上下文中进行的(<0)？ .. 并提供 if 子句和/或 __float_as_int() 的 C 等价物马可？

Answer 1

__float_as_int重新诠释 float作为int . int是<0当它具有最高有效位时。对于 float这也意味着符号位打开，但并不完全意味着数字是负数(例如，它可以是“负零”)。检查然后检查是否 float 会更快是< 0.0 .

C 函数看起来像:

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

在这个标题的一些其他版本中 __cuda___signbitf而是使用。