How do I calculate the square root of a half type in device code?

[jetsonnvidia]

I have sqrt for double and sqrtf for float.

What do I use for half? Must I convert to float first? Seems rather wasteful…

Thanks!

[Robert_Crovella]

https://docs.nvidia.com/cuda/cuda-math-api/group__CUDA__MATH____HALF__FUNCTIONS.html#group__CUDA__MATH____HALF__FUNCTIONS

[jetsonnvidia]

I realise there is hsqrt but this “Calculates half square root in round-to-nearest-even mode.”

I don’t want any rounding especially to the nearest even number (which is incredibly inaccurate for small numbers). I want the result of square root of 5 to be 2.236… not have it rounded to 2. Is this possible?

[Robert_Crovella]

that’s not what “round to nearest even” means. It applies to the least signficant digit of the mantissa/significand. Rounding rules and modes apply when creating a IEEE-style floating point result, even if the function description doesn’t specifically call out one.

Perhaps you should give it a try.

# cat t58.cu
#include <cstdio>
#include <cuda_fp16.h>

__global__ void k(half x, float y){

  printf("%f\n", __half2float(hsqrt(x)));
  printf("%f\n", sqrtf(y));
}

int main(){

  k<<<1,1>>>(5, 5);
  cudaDeviceSynchronize();
}
# nvcc -o t58 t58.cu
# ./t58
2.236328
2.236068
#

[jetsonnvidia]

Thank you so much for explaining that. I have googled the heck out of it and only met with ambiguous explanations. Even the example in the Wikipedia article is confusing: Rounding - Wikipedia (this link followed from the link you gave).

Makes so much more sense!

Muchos grasias.

[jetsonnvidia]

BTW, what I actually get is this:

nano@jetson-nano:/home/nano$ /usr/local/cuda-10.2/bin/nvcc -gencode arch=compute_53,code=sm_53 t58.cu 
nano@jetson-nano:/home/nano$ ./a.out 
nan
2.236068