c - ACOSF实现中的CUDA __float_as_int

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

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

其中acosf和a是t1并且floats是先前设置为接近数学常数Pi的数值的CUDART_PI_F。
我试图理解条件（if子句）测试的是什么，它的C等价物是什么，或者函数/宏是什么。我寻找float的实现，但没有成功。似乎__float_as_int(a)是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_as_int()到__float_as_int()的舍入。（这会产生一个float）而不是将int作为位拷贝（cvt.s32.f32）分配给float %f1（注意：b32是%r1类型（无符号int）！！）然后比较%r1，就好像它是一个u32（带符号的int，令人困惑！！）使用%r1（who's value iss32）。
对我来说这看起来有点奇怪。但显然这是正确的。
有人能解释发生了什么，特别是在if子句测试为否定（%r2）的上下文中，解释0在做什么吗？.. 并提供if子句和/或__float_as_int()marco的C等价物？

最佳答案

__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。

关于c - ACOSF实现中的CUDA __float_as_int，我们在Stack Overflow上找到一个类似的问题：https://stackoverflow.com/questions/13801808/