diff --git a/sycl/include/CL/sycl/half_type.hpp b/sycl/include/CL/sycl/half_type.hpp index ee25b5a6fabd1..647a0abff7fff 100644 --- a/sycl/include/CL/sycl/half_type.hpp +++ b/sycl/include/CL/sycl/half_type.hpp @@ -18,6 +18,14 @@ #include #include +#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` template <> struct hash { @@ -307,35 +306,43 @@ template <> struct numeric_limits { 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 diff --git a/sycl/test/regression/constexpr-fp16-numeric-limits.cpp b/sycl/test/regression/constexpr-fp16-numeric-limits.cpp new file mode 100644 index 0000000000000..c3da16460ee52 --- /dev/null +++ b/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 + +int main() { + constexpr cl::sycl::half L1 = std::numeric_limits::min(); + constexpr cl::sycl::half L2 = std::numeric_limits::max(); + constexpr cl::sycl::half L3 = std::numeric_limits::lowest(); + constexpr cl::sycl::half L4 = std::numeric_limits::epsilon(); + constexpr cl::sycl::half L5 = + std::numeric_limits::round_error(); + constexpr cl::sycl::half L6 = std::numeric_limits::infinity(); + constexpr cl::sycl::half L7 = + std::numeric_limits::quiet_NaN(); + constexpr cl::sycl::half L8 = + std::numeric_limits::signaling_NaN(); + constexpr cl::sycl::half L9 = + std::numeric_limits::denorm_min(); + + return 0; +}