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

Question

__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

It's a type-pun, not a conversion, using unsigned integer min or max on the FP bit-pattern. (Thanks to the exponent bias, increasing the magnitude of a float always increases the integer value of the bit pattern, e.g. `nextafter(x, +Inf)` can be done for non-negative `x` with an integer increment.) Is that what you're asking about? — Peter Cordes, Oct 17 '22 at 10:37
@PeterCordes so casting does not change the underlying bit pattern? it just reinterprets the floating point bit pattern as an integer. is it the case am I understanding it right? — javesh bhardwaj, Oct 18 '22 at 07:09
Yes, exactly. Note that it's a pointer cast, not `(int)*addr`. — Peter Cordes, Oct 18 '22 at 09:58

talonmies · Accepted Answer · 2022-10-19T13:39:26.063

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.

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

1 Answers1