Search code examples
cudaatomicmin

Is there no precision loss because of casting, in below code?


__device__ __forceinline__ float atomicMinFloat (float * addr, float value) {
        float old;
        old = (value >= 0) ? __int_as_float(atomicMin((int *)addr, __float_as_int(value))) :
             __uint_as_float(atomicMax((unsigned int *)addr, __float_as_uint(value)));

        return old;
}

Original Answer: https://stackoverflow.com/a/51549250/5858405


Solution

  • There is no loss of precision in that code because there is no type conversion involved, only a pointer cast. The atomic comparison operation is performed on the bit pattern of the floating point numbers stored in IEEE-754 binary32 format.

    Consider the following trivial example of three floating point numbers:

     Decimal      Nearest IEEE 754
     8.7654f  -->  0x410c3f14
     4.3210f  -->  0x408a45a2
    -1.2345f  -->  0xbf9e0419
    

    and follow the code:

    floatmin(8.7654f, 4.3210f)  --> intmin(0x410c3f14, 0x408a45a2) = 0x408a45a2 --> 4.32100009918
    floatmin(8.7654f, -1.2345f) --> intmax(0x410c3f14, 0xbf9e0419) = 0xbf9e0419 --> -1.23450005054
    floatmin(4.3210f, -1.2345f) --> intmax(0x408a45a2, 0xbf9e0419) = 0xbf9e0419 --> -1.23450005054
    

    Note that the sign in IEEE-754 format is stored in the most significant bit, so the comparison used needs to change depending on the sign of the arguments. But please read all the answers and all the comments in the question you linked to carefully, because there are some corner cases which that code does not handle correctly and you must understand those clearly before using that function in your own program.