[SYCL] Fix numeric_limits<half> after #1089 (#1138)

Error was reproducible in two cases:
- using something like `numeric_limits<half>::min()` in within another
  `constexpr`
- not treating SYCL headers as system ones with `-Winvalid-constexpr`
  treated as error

Signed-off-by: Alexey Sachkov <[email protected]>

Author: AlexeySachkov

sycl/include/CL/sycl/half_type.hpp

@@ -18,6 +18,14 @@
 #include <iostream>
 #include <limits>
 
+#ifdef __SYCL_DEVICE_ONLY__
+// `constexpr` could work because the implicit conversion from `float` to
+// `_Float16` can be `constexpr`.
+#define __SYCL_CONSTEXPR_ON_DEVICE constexpr
+#else
+#define __SYCL_CONSTEXPR_ON_DEVICE
+#endif
+
 __SYCL_INLINE_NAMESPACE(cl) {
 namespace sycl {
 namespace detail {
@@ -130,7 +138,7 @@ class half {
   half(const half &) = default;
   half(half &&) = default;
 
-  half(const float &rhs) : Data(rhs) {}
+  __SYCL_CONSTEXPR_ON_DEVICE half(const float &rhs) : Data(rhs) {}
 
   half &operator=(const half &rhs) = default;
 
@@ -216,15 +224,6 @@ using half = cl::sycl::detail::half_impl::half;
 // Partial specialization of some functions in namespace `std`
 namespace std {
 
-#ifdef __SYCL_DEVICE_ONLY__
-// `constexpr` could work because the implicit conversion from `float` to
-// `_Float16` can be `constexpr`.
-#define CONSTEXPR_QUALIFIER constexpr
-#else
-// The qualifier is `const` instead of `constexpr` that is original to be
-// because the constructor is not `constexpr` function.
-#define CONSTEXPR_QUALIFIER const
-#endif
 
 // Partial specialization of `std::hash<cl::sycl::half>`
 template <> struct hash<half> {
@@ -307,35 +306,43 @@ template <> struct numeric_limits<half> {
 
   static constexpr const float_round_style round_style = round_to_nearest;
 
-  static CONSTEXPR_QUALIFIER half min() noexcept { return SYCL_HLF_MIN; }
+  static __SYCL_CONSTEXPR_ON_DEVICE const half min() noexcept {
+    return SYCL_HLF_MIN;
+  }
 
-  static CONSTEXPR_QUALIFIER half max() noexcept { return SYCL_HLF_MAX; }
+  static __SYCL_CONSTEXPR_ON_DEVICE const half max() noexcept {
+    return SYCL_HLF_MAX;
+  }
 
-  static CONSTEXPR_QUALIFIER half lowest() noexcept { return -SYCL_HLF_MAX; }
+  static __SYCL_CONSTEXPR_ON_DEVICE const half lowest() noexcept {
+    return -SYCL_HLF_MAX;
+  }
 
-  static CONSTEXPR_QUALIFIER half epsilon() noexcept {
+  static __SYCL_CONSTEXPR_ON_DEVICE const half epsilon() noexcept {
     return SYCL_HLF_EPSILON;
   }
 
-  static CONSTEXPR_QUALIFIER half round_error() noexcept { return 0.5F; }
+  static __SYCL_CONSTEXPR_ON_DEVICE const half round_error() noexcept {
+    return 0.5F;
+  }
 
-  static CONSTEXPR_QUALIFIER half infinity() noexcept {
+  static __SYCL_CONSTEXPR_ON_DEVICE const half infinity() noexcept {
     return __builtin_huge_valf();
   }
 
-  static CONSTEXPR_QUALIFIER half quiet_NaN() noexcept {
+  static __SYCL_CONSTEXPR_ON_DEVICE const half quiet_NaN() noexcept {
     return __builtin_nanf("");
   }
 
-  static CONSTEXPR_QUALIFIER half signaling_NaN() noexcept {
+  static __SYCL_CONSTEXPR_ON_DEVICE const half signaling_NaN() noexcept {
     return __builtin_nansf("");
   }
 
-  static CONSTEXPR_QUALIFIER half denorm_min() noexcept { return 5.96046e-08F; }
+  static __SYCL_CONSTEXPR_ON_DEVICE const half denorm_min() noexcept {
+    return 5.96046e-08F;
+  }
 };
 
-#undef CONSTEXPR_QUALIFIER
-
 } // namespace std
 
 inline std::ostream &operator<<(std::ostream &O, half const &rhs) {
@@ -349,3 +356,5 @@ inline std::istream &operator>>(std::istream &I, half &rhs) {
   rhs = ValFloat;
   return I;
 }
+
+#undef __SYCL_CONSTEXPR_ON_DEVICE

sycl/test/regression/constexpr-fp16-numeric-limits.cpp

@@ -0,0 +1,21 @@
+// RUN: %clangxx -fsycl-device-only -fsyntax-only -Xclang -verify %s
+// expected-no-diagnostics
+#include <CL/sycl.hpp>
+
+int main() {
+  constexpr cl::sycl::half L1 = std::numeric_limits<cl::sycl::half>::min();
+  constexpr cl::sycl::half L2 = std::numeric_limits<cl::sycl::half>::max();
+  constexpr cl::sycl::half L3 = std::numeric_limits<cl::sycl::half>::lowest();
+  constexpr cl::sycl::half L4 = std::numeric_limits<cl::sycl::half>::epsilon();
+  constexpr cl::sycl::half L5 =
+      std::numeric_limits<cl::sycl::half>::round_error();
+  constexpr cl::sycl::half L6 = std::numeric_limits<cl::sycl::half>::infinity();
+  constexpr cl::sycl::half L7 =
+      std::numeric_limits<cl::sycl::half>::quiet_NaN();
+  constexpr cl::sycl::half L8 =
+      std::numeric_limits<cl::sycl::half>::signaling_NaN();
+  constexpr cl::sycl::half L9 =
+      std::numeric_limits<cl::sycl::half>::denorm_min();
+
+  return 0;
+}