rtn in ternary way (#632)

* rtn in ternary way

* Check both flags to preserve NaN

* Format

* Rearrange flag1

* Apply suggestions from code review
Co-authored-by: Ronan Keryell <ronan@keryell.fr>

---------
Co-authored-by: Rosty Geyyer <rosty.geyyer@amd.com>
Co-authored-by: Rostyslav Geyyer <46627076+geyyer@users.noreply.github.com>
Co-authored-by: Ronan Keryell <ronan@keryell.fr>

include/ck/utility/data_type.hpp

@@ -974,8 +974,6 @@ inline __host__ __device__ constexpr bhalf_t type_convert<bhalf_t, float>(float
         uint32_t int32;
     } u = {x};
 
-    if(~u.int32 & 0x7f800000)
-    {
     // When the exponent bits are not all 1s, then the value is zero, normal,
     // or subnormal. We round the bfloat16 mantissa up by adding 0x7FFF, plus
     // 1 if the least significant bit of the bfloat16 mantissa is 1 (odd).
@@ -987,15 +985,13 @@ inline __host__ __device__ constexpr bhalf_t type_convert<bhalf_t, float>(float
     // has the value 0x7f, then incrementing it causes it to become 0x00 and
     // the exponent is incremented by one, which is the next higher FP value
     // to the unrounded bfloat16 value. When the bfloat16 value is subnormal
-        // with an exponent of 0x00 and a mantissa of 0x7F, it may be rounded up
+    // with an exponent of 0x00 and a mantissa of 0x7f, it may be rounded up
     // to a normal value with an exponent of 0x01 and a mantissa of 0x00.
     // When the bfloat16 value has an exponent of 0xFE and a mantissa of 0x7F,
     // incrementing it causes it to become an exponent of 0xFF and a mantissa
     // of 0x00, which is Inf, the next higher value to the unrounded value.
-        u.int32 += 0x7fff + ((u.int32 >> 16) & 1); // Round to nearest, round to even
-    }
-    else if(u.int32 & 0xffff)
-    {
+    bool flag0 = ~u.int32 & 0x7f800000;
+
     // When all of the exponent bits are 1, the value is Inf or NaN.
     // Inf is indicated by a zero mantissa. NaN is indicated by any nonzero
     // mantissa bit. Quiet NaN is indicated by the most significant mantissa
@@ -1003,9 +999,11 @@ inline __host__ __device__ constexpr bhalf_t type_convert<bhalf_t, float>(float
     // mantissa bit being 0 but some other bit(s) being 1. If any of the
     // lower 16 bits of the mantissa are 1, we set the least significant bit
     // of the bfloat16 mantissa, in order to preserve signaling NaN in case
-        // the bloat16's mantissa bits are all 0.
-        u.int32 |= 0x10000; // Preserve signaling NaN
-    }
+    // the bfloat16's mantissa bits are all 0.
+    bool flag1 = !flag0 && (u.int32 & 0xffff);
+
+    u.int32 += flag0 ? 0x7fff + ((u.int32 >> 16) & 1) : 0; // Round to nearest, round to even
+    u.int32 |= flag1 ? 0x10000 : 0x0;                      // Preserve signaling NaN
 
     return uint16_t(u.int32 >> 16);
 }