[SYCL] Fix float-to-half conversion for minimum subnormal on host (#7401)

Currently when converting a floating point value representing the
minimum subnormal value for half to a half on host will result in 0
rather than the expected minimum value. This is due to a faulty check
for the minimum supported single-precision exponent when converting.
This commit fixes this check.

Signed-off-by: Larsen, Steffen <[email protected]>

Author: steffenlarsen

sycl/include/sycl/half_type.hpp

@@ -78,7 +78,7 @@ inline __SYCL_CONSTEXPR_HALF uint16_t float2Half(const float &Val) {
     // Tie to even.
     else if (roundBits == halfway)
       Frac16 += Frac16 & 1;
-  } else if (__builtin_expect(Exp32Diff > -24, 0)) {
+  } else if (__builtin_expect(Exp32Diff > -25, 0)) {
     // subnormals
     Frac16 = (Frac32 | (uint32_t(1) << 23)) >> (-Exp32Diff - 1);
   }

sycl/test/regression/half_host_subnormal_min.cpp

@@ -0,0 +1,25 @@
+// RUN: %clangxx -fsycl %s -o %t.out
+// RUN: %t.out
+//
+// Checks that sycl::half on host can correctly cast its minimum subnormal value
+// to and from a floating point value.
+
+#include <sycl/sycl.hpp>
+
+int main() {
+  sycl::half SubnormalMin =
+      sycl::bit_cast<sycl::half>((uint16_t)0b0000000000000001u);
+  sycl::half ConvertedSubnormalMin =
+      static_cast<sycl::half>(static_cast<float>(SubnormalMin));
+
+  if (SubnormalMin != ConvertedSubnormalMin) {
+    std::cout << "Failed! (0x" << std::hex
+              << sycl::bit_cast<uint16_t>(SubnormalMin) << " != 0x"
+              << sycl::bit_cast<uint16_t>(ConvertedSubnormalMin) << ")"
+              << std::endl;
+    return 1;
+  }
+
+  std::cout << "Passed!" << std::endl;
+  return 0;
+}