[SYCL] Fix nextafter with half on host (intel#10558)

Currently the host-side implementation of sycl::nextafter with
sycl::half uses the float variant of std::nextafter. However, due to the
conversion between half and float, the result may be unexpected.
Likewise, KhronosGroup/SYCL-Docs#440 removes the
reference to single-precision floating point results.

Signed-off-by: Larsen, Steffen <steffen.larsen@intel.com>

sycl/source/detail/builtins_math.cpp

@@ -538,7 +538,23 @@ __SYCL_EXPORT s::cl_double sycl_host_nextafter(s::cl_double x,
 }
 __SYCL_EXPORT s::cl_half sycl_host_nextafter(s::cl_half x,
                                              s::cl_half y) __NOEXC {
-  return std::nextafter(x, y);
+  if (std::isnan(d::cast_if_host_half(x)))
+    return x;
+  if (std::isnan(d::cast_if_host_half(y)) || x == y)
+    return y;
+
+  uint16_t x_bits = s::bit_cast<uint16_t>(x);
+  uint16_t x_sign = x_bits & 0x8000;
+  int16_t movement = (x > y ? -1 : 1) * (x_sign ? -1 : 1);
+  if (x_bits == x_sign && movement == -1) {
+    // Special case where we underflow in the decrement, in which case we turn
+    // it around and flip the sign. The overflow case does not need special
+    // handling.
+    movement = 1;
+    x_bits ^= 0x8000;
+  }
+  x_bits += movement;
+  return s::bit_cast<s::cl_half>(x_bits);
 }
 MAKE_1V_2V(sycl_host_nextafter, s::cl_float, s::cl_float, s::cl_float)
 MAKE_1V_2V(sycl_host_nextafter, s::cl_double, s::cl_double, s::cl_double)

sycl/test/regression/host_half_nextafter.cpp

@@ -0,0 +1,30 @@
+// RUN: %clangxx -fsycl %s -o %t.out
+// RUN: %t.out
+//
+// Checks that sycl::nextafter with sycl::half on host correctly converts based
+// on half-precision.
+
+#include <sycl/sycl.hpp>
+
+void check(uint16_t x, uint16_t y, uint16_t ref) {
+  assert(sycl::nextafter(sycl::bit_cast<sycl::half>(x),
+                         sycl::bit_cast<sycl::half>(y)) ==
+         sycl::bit_cast<sycl::half>(ref));
+}
+
+int main() {
+  check(0x0, 0x0, 0x0);
+  check(0x1, 0x1, 0x1);
+  check(0x8001, 0x8001, 0x8001);
+  check(0x0, 0x1, 0x1);
+  check(0x8000, 0x8001, 0x8001);
+  check(0x0, 0x8001, 0x8001);
+  check(0x8000, 0x1, 0x1);
+  check(0x8001, 0x0, 0x0);
+  check(0x1, 0x8000, 0x8000);
+  check(0x8001, 0x1, 0x0);
+  check(0x1, 0x8001, 0x8000);
+
+  std::cout << "Passed!" << std::endl;
+  return 0;
+}