Will the float loses its precision when a number is small enough (GPU Kernel, RTX 2060)?
Say,
#include <stdio.h>
__global__ void calc()
{
float m1 = 0.1490116119;
float m2 = -0.000000007450580;
float res = m1 + m2;
printf("M1: %.15f M2: %.15f Res: %.15f\n", m1, m2, res);
}
int main()
{
calc<<<1,1>>>();
cudaDeviceSynchronize();
return 0;
}
Results as:
M1: 0.149011611938477 M2: -0.000000007450580 Res: 0.149011611938477
M2 does not matter in this case, which confuses me.
I know that float number may not be as accurate as expected, but I wonder if there is something wrong that float in the order of 1e-9 is ignored.
Assume that if I want to add a very small number to a given variable (called M) gradually, this benchmark would imply that M will not change any more finally?
Thank you very much!
Your observations are correct and they don't have anything to do with CUDA:
$ cat t1981.cu
#include <stdio.h>
#ifndef DONT_USE_CUDA
__global__ void calc()
#else
int main()
#endif
{
float m1 = 0.1490116119;
float m2 = -0.000000007450580;
float res = m1 + m2;
printf("M1: %.15f M2: %.15f Res: %.15f\n", m1, m2, res);
}
#ifndef DONT_USE_CUDA
int main()
{
calc<<<1,1>>>();
cudaDeviceSynchronize();
return 0;
}
#endif
$ nvcc -o t1981 t1981.cu
$ ./t1981
M1: 0.149011611938477 M2: -0.000000007450580 Res: 0.149011611938477
$ cp t1981.cu t1981.cpp
$ g++ -DDONT_USE_CUDA t1981.cpp -o t1981
$ ./t1981
M1: 0.149011611938477 M2: -0.000000007450580 Res: 0.149011611938477
$
M2 does not matter in this case, which confuses me.
A float quantity corresponds to IEEE 754 32-bit floating point. This number representation has ~23 bits for mantissa representation. 23 bits corresponds to roughly 6 or 7 decimal digits. The way I like to think about this is that there can be about 6 or 7 decimal digits between the most significant digit you want to represent and the least significant digit you want to represent (inclusive).
I know that float number may not be as accurate as expected, but I wonder if there is something wrong that float in the order of 1e-9 is ignored.
It's not that categorically, a number of the order of 1e-9 is ignored. Instead, it's better to conclude that if there are 9 significant digits between the digits in the number(s) that you care about, float may be an insufficient representation. You can easily work with numbers that are in the range of 1e-9, it's just that you may not have good results combining those with numbers in the range of 1e-1, which is what you are trying to do.
Assume that if I want to add a very small number to a given variable (called M) gradually, this benchmark would imply that M will not change any more finally?
Yes, if you want to handle calculations like this, one possible solution would be to switch from float to double. double can handle typically 12-15 or more decimal digits of significance.
There are also other alternatives, such as kahan summations, to deal with problems like this. You could also sort a set of numbers to be added. In a serial fashion, or blockwise in CUDA, adding from smallest to largest may give better results than an unsorted or naive summation.
Also note that a typical parallel reduction effectively performs the pairwise summation approach discussed here and so may perform better than a naive serial running-sum approach.