float(0.0) vs 0.0f in CUDA

float(0.0) vs 0.0f in CUDA

本文关键字:in CUDA 0f float vs      更新时间:2023-10-16

我想编写浮点和双精度有效的代码。我正在做这样的事情:

typedef real float;
//typedef real double;
__global__ void foo(real a, real *b){
  b[0] = real(0.5)*a;
}
int main(){
  real a = 1.0f;
  real *b;
  cudaMalloc(&f, sizeof(real));
  foo<<<1,1>>>(a,b);
  return 0;
}

这让我想到,我不想在做双精度时在常量中失去精度0.5f,但我不想在做单精度时将0.5提升到双精度!

因此,我最终使用了操作符real(),如示例中所示。在单精度模式下,如果我使用real(0.5)反汇编函数'foo',我得到没有提升到double,而不是只使用0.5,那里会发生提升。

你可以使用:

$nvcc test.cu -arch=sm_52 -lineinfo --source-in-ptx -g -G -O0 ; cuobjdump -sass  a.out | grep "foo" -A 35

I see

/*0078*/                   FMUL R0, R0, 0.5;               /* 0x3868004000070000 */

使用real(0.5)或0.5f时和:

/*0078*/                   F2F.F64.F32 R4, R0;           /* 0x5ca8000000070b04 */
/*0088*/                   DMUL R4, R4, 0.5;               /* 0x3880004000070404 */
/*0090*/                   F2F.F32.F64 R0, R4;           /* 0x5ca8000000470e00 */

当只写0.5。

这听起来可能太明显了。但是由于我不知道"real(0.5)"在做什么,所以我不知道这是否只是编译器在这个非常特殊的情况下的作用。反汇编代码在真实版本(0.5)和0.5f中似乎是相同的!

所以问题仍然存在:

real(0.5)(又名float(0.5))到底在做什么?

float(0.5)和0.5f之间有什么区别吗?(OR double(0.5) and 0.5)

real(0.5) (AKA float(0.5))到底在做什么?

real(0.5)函数样式强制转换,在本例中它降低为static_cast

real(0.5)
static_cast<real>(0.5) //exactly the same thing

这意味着areal变量(在本例中为float)相乘,这意味着不需要像double * float乘法那样执行到double的提升。

float(0.5)和0.5f有什么区别吗?(OR double(0.5) and 0.5)

有人可能认为float0.5的初始化可以在运行时发生,但这对于任何现代编译器来说都是不现实的。它应该是一个no-op,它已经是一个op了。

除此之外,使用float(0.5f)与仅使用0.5f没有任何区别,double(0.5)0.5也是如此。