From 8f5fe4de8e32146ca1a4fafe1f7f9e1a393d3161 Mon Sep 17 00:00:00 2001 From: y Date: Tue, 21 Apr 2026 08:10:59 -0700 Subject: [PATCH 1/2] [SYCL] Split sycl::half into an internal impl header Move the sycl::half implementation into detail/half_type_impl.hpp and keep sycl/half_type.hpp as the public wrapper for stream operators and std::hash. This lets internal headers, including builtins, depend on the half implementation without pulling in the full public half header that relies on std::optional and iostream. This reduces compile time by another 30ms on builtins. --- .../include/sycl/detail/builtins/builtins.hpp | 2 +- sycl/include/sycl/detail/half_type_impl.hpp | 566 +++++++++++++++++ .../intel/esimd/detail/half_type_traits.hpp | 2 +- .../sycl/ext/intel/math/imf_half_trivial.hpp | 2 +- .../experimental/complex/detail/common.hpp | 2 +- sycl/include/sycl/half_type.hpp | 584 +----------------- sycl/include/sycl/known_identity.hpp | 2 +- sycl/test/include_deps/sycl_accessor.hpp.cpp | 1 + .../include_deps/sycl_detail_core.hpp.cpp | 1 + .../sycl_khr_includes_accessor.hpp.cpp | 1 + .../sycl_khr_includes_context.hpp.cpp | 1 + .../sycl_khr_includes_device.hpp.cpp | 1 + .../sycl_khr_includes_event.hpp.cpp | 1 + ...sycl_khr_includes_group_algorithms.hpp.cpp | 1 + .../sycl_khr_includes_groups.hpp.cpp | 1 + .../sycl_khr_includes_half.hpp.cpp | 2 +- .../sycl_khr_includes_handler.hpp.cpp | 1 + .../sycl_khr_includes_images.hpp.cpp | 1 + .../sycl_khr_includes_interop_handle.hpp.cpp | 1 + .../sycl_khr_includes_kernel_bundle.hpp.cpp | 1 + .../sycl_khr_includes_math.hpp.cpp | 14 +- .../sycl_khr_includes_platform.hpp.cpp | 1 + .../sycl_khr_includes_queue.hpp.cpp | 1 + .../sycl_khr_includes_reduction.hpp.cpp | 1 + .../sycl_khr_includes_stream.hpp.cpp | 1 + .../sycl_khr_includes_usm.hpp.cpp | 15 +- 26 files changed, 618 insertions(+), 589 deletions(-) create mode 100644 sycl/include/sycl/detail/half_type_impl.hpp diff --git a/sycl/include/sycl/detail/builtins/builtins.hpp b/sycl/include/sycl/detail/builtins/builtins.hpp index 04658bf84e7b0..c2033bfe2e697 100644 --- a/sycl/include/sycl/detail/builtins/builtins.hpp +++ b/sycl/include/sycl/detail/builtins/builtins.hpp @@ -64,11 +64,11 @@ #pragma once #include +#include #include #include #include #include -#include #include #include diff --git a/sycl/include/sycl/detail/half_type_impl.hpp b/sycl/include/sycl/detail/half_type_impl.hpp new file mode 100644 index 0000000000000..494d43dec0f1f --- /dev/null +++ b/sycl/include/sycl/detail/half_type_impl.hpp @@ -0,0 +1,566 @@ +//==----------- half_type_impl.hpp - SYCL half type implementation ---------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#pragma once + +#include // for bit_cast +#include + +#ifdef __SYCL_DEVICE_ONLY__ +#include +#endif + +#include +#include // for float_denorm_style, float_round_style + +#if !defined(__has_builtin) || !__has_builtin(__builtin_expect) +#define __builtin_expect(a, b) (a) +#endif + +#ifdef __SYCL_DEVICE_ONLY__ +// `constexpr` could work because the implicit conversion from `float` to +// `_Float16` can be `constexpr`. +#define __SYCL_CONSTEXPR_HALF constexpr +#elif __cpp_lib_bit_cast || \ + (defined(__has_builtin) && __has_builtin(__builtin_bit_cast)) +#define __SYCL_CONSTEXPR_HALF constexpr +#else +#define __SYCL_CONSTEXPR_HALF +#endif + +namespace sycl { +inline namespace _V1 { + +namespace ext::intel::esimd::detail { +class WrapperElementTypeProxy; +} // namespace ext::intel::esimd::detail + +namespace detail { + +inline __SYCL_CONSTEXPR_HALF uint16_t float2Half(const float &Val) { + const uint32_t Bits = sycl::bit_cast(Val); + + // Extract the sign from the float value + const uint16_t Sign = (Bits & 0x80000000) >> 16; + // Extract the fraction from the float value + const uint32_t Frac32 = Bits & 0x7fffff; + // Extract the exponent from the float value + const uint8_t Exp32 = (Bits & 0x7f800000) >> 23; + const int16_t Exp32Diff = Exp32 - 127; + + // initialize to 0, covers the case for 0 and small numbers + uint16_t Exp16 = 0, Frac16 = 0; + + if (__builtin_expect(Exp32Diff > 15, 0)) { + // Infinity and big numbers convert to infinity + Exp16 = 0x1f; + } else if (__builtin_expect(Exp32Diff > -14, 0)) { + // normal range for half type + Exp16 = static_cast(Exp32Diff) + 15; + // convert 23-bit mantissa to 10-bit mantissa. + Frac16 = static_cast(Frac32 >> 13); + // Round the mantissa as given in OpenCL spec section : 6.1.1.1 The half + // data type. + // Round to nearest. + uint32_t roundBits = Frac32 & 0x1fff; + uint32_t halfway = 0x1000; + if (roundBits > halfway) + Frac16 += 1; + // Tie to even. + else if (roundBits == halfway) + Frac16 += Frac16 & 1; + } else if (__builtin_expect(Exp32Diff > -25, 0)) { + // subnormals + Frac16 = static_cast((Frac32 | (uint32_t(1) << 23)) >> + (-Exp32Diff - 1)); + } + + if (__builtin_expect(Exp32 == 0xff && Frac32 != 0, 0)) { + // corner case: FP32 is NaN + Exp16 = 0x1F; + Frac16 = 0x200; + } + + // Compose the final FP16 binary + uint16_t Ret = 0; + Ret |= Sign; + Ret |= Exp16 << 10; + Ret += Frac16; // Add the carry bit from operation Frac16 += 1; + + return Ret; +} + +inline __SYCL_CONSTEXPR_HALF float half2Float(const uint16_t &Val) { + // Extract the sign from the bits. It is 1 if the sign is negative + const uint32_t Sign = static_cast(Val & 0x8000) << 16; + // Extract the exponent from the bits + const uint8_t Exp16 = (Val & 0x7c00) >> 10; + // Extract the fraction from the bits + uint16_t Frac16 = Val & 0x3ff; + + uint32_t Exp32 = 0; + if (__builtin_expect(Exp16 == 0x1f, 0)) { + Exp32 = 0xff; + } else if (__builtin_expect(Exp16 == 0, 0)) { + Exp32 = 0; + } else { + Exp32 = static_cast(Exp16) + 112; + } + // corner case: subnormal -> normal + // The denormal number of FP16 can be represented by FP32, therefore we need + // to recover the exponent and recalculate the fration. + if (__builtin_expect(Exp16 == 0 && Frac16 != 0, 0)) { + uint8_t OffSet = 0; + do { + ++OffSet; + Frac16 <<= 1; + } while ((Frac16 & 0x400) != 0x400); + // mask the 9th bit + Frac16 &= 0x3ff; + Exp32 = 113 - OffSet; + } + + uint32_t Frac32 = static_cast(Frac16 << 13); + + uint32_t Bits = 0; + Bits |= Sign; + Bits |= (Exp32 << 23); + Bits |= Frac32; + const float Result = sycl::bit_cast(Bits); + return Result; +} + +namespace half_impl { +class half; + +// Creation token to disambiguate constructors. +struct RawHostHalfToken { + constexpr explicit RawHostHalfToken(uint16_t Val) : Value{Val} {} + uint16_t Value; +}; + +#ifndef __SYCL_DEVICE_ONLY__ +class half { +#else +class [[__sycl_detail__::__uses_aspects__(aspect::fp16)]] half { +#endif +public: + half() = default; + constexpr half(const half &) = default; + constexpr half(half &&) = default; + +#ifdef __SYCL_DEVICE_ONLY__ + __SYCL_CONSTEXPR_HALF half(const float &rhs) : Data(rhs) {} +#else + __SYCL_CONSTEXPR_HALF half(const float &rhs) : Data(float2Half(rhs)) {} +#endif // __SYCL_DEVICE_ONLY__ + + constexpr half &operator=(const half &rhs) = default; + + // Operator +=, -=, *=, /= +#ifdef __SYCL_DEVICE_ONLY__ + __SYCL_CONSTEXPR_HALF half &operator+=(const half &rhs) { + Data += rhs.Data; + return *this; + } + + __SYCL_CONSTEXPR_HALF half &operator-=(const half &rhs) { + Data -= rhs.Data; + return *this; + } + + __SYCL_CONSTEXPR_HALF half &operator*=(const half &rhs) { + Data *= rhs.Data; + return *this; + } + + __SYCL_CONSTEXPR_HALF half &operator/=(const half &rhs) { + Data /= rhs.Data; + return *this; + } +#else + __SYCL_CONSTEXPR_HALF half &operator+=(const half &rhs) { + *this = operator float() + static_cast(rhs); + return *this; + } + + __SYCL_CONSTEXPR_HALF half &operator-=(const half &rhs) { + *this = operator float() - static_cast(rhs); + return *this; + } + + __SYCL_CONSTEXPR_HALF half &operator*=(const half &rhs) { + *this = operator float() * static_cast(rhs); + return *this; + } + + __SYCL_CONSTEXPR_HALF half &operator/=(const half &rhs) { + *this = operator float() / static_cast(rhs); + return *this; + } +#endif // __SYCL_DEVICE_ONLY__ + + // Operator ++, -- + __SYCL_CONSTEXPR_HALF half &operator++() { + *this += 1; + return *this; + } + + __SYCL_CONSTEXPR_HALF half operator++(int) { + half ret(*this); + operator++(); + return ret; + } + + __SYCL_CONSTEXPR_HALF half &operator--() { + *this -= 1; + return *this; + } + + __SYCL_CONSTEXPR_HALF half operator--(int) { + half ret(*this); + operator--(); + return ret; + } + + // Operator neg +#ifdef __SYCL_DEVICE_ONLY__ + __SYCL_CONSTEXPR_HALF friend half operator-(const half other) { + return half(-other.Data); + } +#else + __SYCL_CONSTEXPR_HALF friend half operator-(const half other) { + return half(RawHostHalfToken(other.Data ^ 0x8000)); + } +#endif // __SYCL_DEVICE_ONLY__ + +// Operator +, -, *, / +#define OP(op, op_eq) \ + __SYCL_CONSTEXPR_HALF friend half operator op(const half lhs, \ + const half rhs) { \ + half rtn = lhs; \ + rtn op_eq rhs; \ + return rtn; \ + } \ + __SYCL_CONSTEXPR_HALF friend double operator op(const half lhs, \ + const double rhs) { \ + double rtn = lhs; \ + rtn op_eq rhs; \ + return rtn; \ + } \ + __SYCL_CONSTEXPR_HALF friend double operator op(const double lhs, \ + const half rhs) { \ + double rtn = lhs; \ + rtn op_eq rhs; \ + return rtn; \ + } \ + __SYCL_CONSTEXPR_HALF friend float operator op(const half lhs, \ + const float rhs) { \ + float rtn = lhs; \ + rtn op_eq rhs; \ + return rtn; \ + } \ + __SYCL_CONSTEXPR_HALF friend float operator op(const float lhs, \ + const half rhs) { \ + float rtn = lhs; \ + rtn op_eq rhs; \ + return rtn; \ + } \ + __SYCL_CONSTEXPR_HALF friend half operator op(const half lhs, \ + const int rhs) { \ + half rtn = lhs; \ + rtn op_eq half(static_cast(rhs)); \ + return rtn; \ + } \ + __SYCL_CONSTEXPR_HALF friend half operator op(const int lhs, \ + const half rhs) { \ + half rtn(static_cast(lhs)); \ + rtn op_eq rhs; \ + return rtn; \ + } \ + __SYCL_CONSTEXPR_HALF friend half operator op(const half lhs, \ + const long rhs) { \ + half rtn = lhs; \ + rtn op_eq half(static_cast(rhs)); \ + return rtn; \ + } \ + __SYCL_CONSTEXPR_HALF friend half operator op(const long lhs, \ + const half rhs) { \ + half rtn(static_cast(lhs)); \ + rtn op_eq rhs; \ + return rtn; \ + } \ + __SYCL_CONSTEXPR_HALF friend half operator op(const half lhs, \ + const long long rhs) { \ + half rtn = lhs; \ + rtn op_eq half(static_cast(rhs)); \ + return rtn; \ + } \ + __SYCL_CONSTEXPR_HALF friend half operator op(const long long lhs, \ + const half rhs) { \ + half rtn(static_cast(lhs)); \ + rtn op_eq rhs; \ + return rtn; \ + } \ + __SYCL_CONSTEXPR_HALF friend half operator op(const half &lhs, \ + const unsigned int &rhs) { \ + half rtn = lhs; \ + rtn op_eq half(static_cast(rhs)); \ + return rtn; \ + } \ + __SYCL_CONSTEXPR_HALF friend half operator op(const unsigned int &lhs, \ + const half &rhs) { \ + half rtn(static_cast(lhs)); \ + rtn op_eq rhs; \ + return rtn; \ + } \ + __SYCL_CONSTEXPR_HALF friend half operator op(const half &lhs, \ + const unsigned long &rhs) { \ + half rtn = lhs; \ + rtn op_eq half(static_cast(rhs)); \ + return rtn; \ + } \ + __SYCL_CONSTEXPR_HALF friend half operator op(const unsigned long &lhs, \ + const half &rhs) { \ + half rtn(static_cast(lhs)); \ + rtn op_eq rhs; \ + return rtn; \ + } \ + __SYCL_CONSTEXPR_HALF friend half operator op( \ + const half &lhs, const unsigned long long &rhs) { \ + half rtn = lhs; \ + rtn op_eq half(static_cast(rhs)); \ + return rtn; \ + } \ + __SYCL_CONSTEXPR_HALF friend half operator op(const unsigned long long &lhs, \ + const half &rhs) { \ + half rtn(static_cast(lhs)); \ + rtn op_eq rhs; \ + return rtn; \ + } + OP(+, +=) + OP(-, -=) + OP(*, *=) + OP(/, /=) + +#undef OP + +// Operator ==, !=, <, >, <=, >= +#define OP(op) \ + __SYCL_CONSTEXPR_HALF friend bool operator op(const half &lhs, \ + const half &rhs) { \ + return lhs.getFPRep() op rhs.getFPRep(); \ + } \ + __SYCL_CONSTEXPR_HALF friend bool operator op(const half &lhs, \ + const double &rhs) { \ + return lhs.getFPRep() op rhs; \ + } \ + __SYCL_CONSTEXPR_HALF friend bool operator op(const double &lhs, \ + const half &rhs) { \ + return lhs op rhs.getFPRep(); \ + } \ + __SYCL_CONSTEXPR_HALF friend bool operator op(const half &lhs, \ + const float &rhs) { \ + return lhs.getFPRep() op rhs; \ + } \ + __SYCL_CONSTEXPR_HALF friend bool operator op(const float &lhs, \ + const half &rhs) { \ + return lhs op rhs.getFPRep(); \ + } \ + __SYCL_CONSTEXPR_HALF friend bool operator op(const half &lhs, \ + const int &rhs) { \ + return lhs.getFPRep() op static_cast(rhs); \ + } \ + __SYCL_CONSTEXPR_HALF friend bool operator op(const int &lhs, \ + const half &rhs) { \ + return static_cast(lhs) op rhs.getFPRep(); \ + } \ + __SYCL_CONSTEXPR_HALF friend bool operator op(const half &lhs, \ + const long &rhs) { \ + return lhs.getFPRep() op static_cast(rhs); \ + } \ + __SYCL_CONSTEXPR_HALF friend bool operator op(const long &lhs, \ + const half &rhs) { \ + return static_cast(lhs) op rhs.getFPRep(); \ + } \ + __SYCL_CONSTEXPR_HALF friend bool operator op(const half &lhs, \ + const long long &rhs) { \ + return lhs.getFPRep() op static_cast(rhs); \ + } \ + __SYCL_CONSTEXPR_HALF friend bool operator op(const long long &lhs, \ + const half &rhs) { \ + return static_cast(lhs) op rhs.getFPRep(); \ + } \ + __SYCL_CONSTEXPR_HALF friend bool operator op(const half &lhs, \ + const unsigned int &rhs) { \ + return lhs.getFPRep() op static_cast(rhs); \ + } \ + __SYCL_CONSTEXPR_HALF friend bool operator op(const unsigned int &lhs, \ + const half &rhs) { \ + return static_cast(lhs) op rhs.getFPRep(); \ + } \ + __SYCL_CONSTEXPR_HALF friend bool operator op(const half &lhs, \ + const unsigned long &rhs) { \ + return lhs.getFPRep() op static_cast(rhs); \ + } \ + __SYCL_CONSTEXPR_HALF friend bool operator op(const unsigned long &lhs, \ + const half &rhs) { \ + return static_cast(lhs) op rhs.getFPRep(); \ + } \ + __SYCL_CONSTEXPR_HALF friend bool operator op( \ + const half &lhs, const unsigned long long &rhs) { \ + return lhs.getFPRep() op static_cast(rhs); \ + } \ + __SYCL_CONSTEXPR_HALF friend bool operator op(const unsigned long long &lhs, \ + const half &rhs) { \ + return static_cast(lhs) op rhs.getFPRep(); \ + } + OP(==) + OP(!=) + OP(<) + OP(>) + OP(<=) + OP(>=) + +#undef OP + + // Operator float +#ifdef __SYCL_DEVICE_ONLY__ + __SYCL_CONSTEXPR_HALF operator float() const { + return static_cast(Data); + } +#else + __SYCL_CONSTEXPR_HALF operator float() const { return half2Float(Data); } +#endif // __SYCL_DEVICE_ONLY__ + + friend class sycl::ext::intel::esimd::detail::WrapperElementTypeProxy; + +private: + // When doing operations, we cannot simply work with Data on host as + // it is an integer. Instead, convert it to float. On device we can work with + // Data as it is already a floating point representation. +#ifdef __SYCL_DEVICE_ONLY__ + __SYCL_CONSTEXPR_HALF StorageT getFPRep() const { return Data; } +#else + __SYCL_CONSTEXPR_HALF float getFPRep() const { return operator float(); } +#endif + +#ifndef __SYCL_DEVICE_ONLY__ + // Because sycl::bit_cast might not be constexpr on certain systems, + // implementation needs shortcut for creating a host sycl::half directly from + // a uint16_t representation. + constexpr explicit half(RawHostHalfToken X) : Data(X.Value) {} + + friend constexpr inline half CreateHostHalfRaw(uint16_t X); +#endif // __SYCL_DEVICE_ONLY__ + + StorageT Data; +}; + +#ifndef __SYCL_DEVICE_ONLY__ +constexpr inline half CreateHostHalfRaw(uint16_t X) { + return half(RawHostHalfToken(X)); +} +#endif // __SYCL_DEVICE_ONLY__ +} // namespace half_impl + +// According to the C++ standard, math functions from cmath/math.h should work +// only on arithmetic types. We can't specify half type as arithmetic/floating +// point(via std::is_floating_point) since only float, double and long double +// types are "floating point" according to the standard. In order to use half +// type with these math functions we cast half to float using template +// function helper. +template inline T cast_if_host_half(T val) { return val; } + +inline float cast_if_host_half(half_impl::half val) { + return static_cast(val); +} + +} // namespace detail + +} // namespace _V1 +} // namespace sycl + +// Partial specialization of some functions in namespace `std` +namespace std { + +// Partial specialization of `std::numeric` +template <> struct numeric_limits { + // All following values are either calculated based on description of each + // function/value on https://en.cppreference.com/w/cpp/types/numeric_limits, + // or cl_platform.h. + static constexpr bool is_specialized = true; + static constexpr bool is_signed = true; + static constexpr bool is_integer = false; + static constexpr bool is_exact = false; + static constexpr bool has_infinity = true; + static constexpr bool has_quiet_NaN = true; + static constexpr bool has_signaling_NaN = true; + static constexpr float_denorm_style has_denorm = denorm_present; + static constexpr bool has_denorm_loss = false; + static constexpr bool tinyness_before = false; + static constexpr bool traps = false; + static constexpr int max_exponent10 = 4; + static constexpr int max_exponent = 16; + static constexpr int min_exponent10 = -4; + static constexpr int min_exponent = -13; + static constexpr int radix = 2; + static constexpr int max_digits10 = 5; + static constexpr int digits = 11; + static constexpr bool is_bounded = true; + static constexpr int digits10 = 3; + static constexpr bool is_modulo = false; + static constexpr bool is_iec559 = true; + static constexpr float_round_style round_style = round_to_nearest; + + static __SYCL_CONSTEXPR_HALF const sycl::half(min)() noexcept { + return 6.103515625e-05f; // half minimum value + } + + static __SYCL_CONSTEXPR_HALF const sycl::half(max)() noexcept { + return 65504.0f; // half maximum value + } + + static __SYCL_CONSTEXPR_HALF const sycl::half lowest() noexcept { + return -65504.0f; // -1*(half maximum value) + } + + static __SYCL_CONSTEXPR_HALF const sycl::half epsilon() noexcept { + return 9.765625e-04f; // half epsilon + } + + static __SYCL_CONSTEXPR_HALF const sycl::half round_error() noexcept { + return 0.5f; + } + + static constexpr const sycl::half infinity() noexcept { +#ifdef __SYCL_DEVICE_ONLY__ + return __builtin_huge_valf(); +#else + return sycl::detail::half_impl::CreateHostHalfRaw( + static_cast(0x7C00)); +#endif + } + + static __SYCL_CONSTEXPR_HALF const sycl::half quiet_NaN() noexcept { + return __builtin_nanf(""); + } + + static __SYCL_CONSTEXPR_HALF const sycl::half signaling_NaN() noexcept { + return __builtin_nansf(""); + } + + static __SYCL_CONSTEXPR_HALF const sycl::half denorm_min() noexcept { + return 5.96046e-08f; + } +}; + +} // namespace std + +#undef __SYCL_CONSTEXPR_HALF +#undef _CPP14_CONSTEXPR \ No newline at end of file diff --git a/sycl/include/sycl/ext/intel/esimd/detail/half_type_traits.hpp b/sycl/include/sycl/ext/intel/esimd/detail/half_type_traits.hpp index 7fa2fc820a5b3..565d6efd750bb 100644 --- a/sycl/include/sycl/ext/intel/esimd/detail/half_type_traits.hpp +++ b/sycl/include/sycl/ext/intel/esimd/detail/half_type_traits.hpp @@ -12,7 +12,7 @@ #include -#include +#include /// @cond ESIMD_DETAIL diff --git a/sycl/include/sycl/ext/intel/math/imf_half_trivial.hpp b/sycl/include/sycl/ext/intel/math/imf_half_trivial.hpp index 5a6d0b9067ede..cda21a9ba08e0 100644 --- a/sycl/include/sycl/ext/intel/math/imf_half_trivial.hpp +++ b/sycl/include/sycl/ext/intel/math/imf_half_trivial.hpp @@ -11,7 +11,7 @@ #pragma once #include -#include +#include #include namespace sycl { diff --git a/sycl/include/sycl/ext/oneapi/experimental/complex/detail/common.hpp b/sycl/include/sycl/ext/oneapi/experimental/complex/detail/common.hpp index e6d5ad047960c..1f60580afd863 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/complex/detail/common.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/complex/detail/common.hpp @@ -10,7 +10,7 @@ #include -#include +#include namespace sycl { inline namespace _V1 { diff --git a/sycl/include/sycl/half_type.hpp b/sycl/include/sycl/half_type.hpp index c33898a589ee7..f210e16816659 100644 --- a/sycl/include/sycl/half_type.hpp +++ b/sycl/include/sycl/half_type.hpp @@ -8,18 +8,7 @@ #pragma once -#include // for bit_cast -#include // for __SYCL_EXPORT -#include - -#ifdef __SYCL_DEVICE_ONLY__ -#include -#endif - -#include -#include -#include // for float_denorm_style, float_r... -#include +#include // For std::hash, seems to be the most lightweight header provide it under // C++17: @@ -31,499 +20,30 @@ #include #endif -#if !defined(__has_builtin) || !__has_builtin(__builtin_expect) -#define __builtin_expect(a, b) (a) -#endif - -#ifdef __SYCL_DEVICE_ONLY__ -// `constexpr` could work because the implicit conversion from `float` to -// `_Float16` can be `constexpr`. -#define __SYCL_CONSTEXPR_HALF constexpr -#elif __cpp_lib_bit_cast || \ - (defined(__has_builtin) && __has_builtin(__builtin_bit_cast)) -#define __SYCL_CONSTEXPR_HALF constexpr -#else -#define __SYCL_CONSTEXPR_HALF -#endif - namespace sycl { inline namespace _V1 { namespace detail::half_impl { -class half; -} -using half = detail::half_impl::half; - -namespace ext::intel::esimd::detail { -class WrapperElementTypeProxy; -} // namespace ext::intel::esimd::detail - -namespace detail { - -inline __SYCL_CONSTEXPR_HALF uint16_t float2Half(const float &Val) { - const uint32_t Bits = sycl::bit_cast(Val); - - // Extract the sign from the float value - const uint16_t Sign = (Bits & 0x80000000) >> 16; - // Extract the fraction from the float value - const uint32_t Frac32 = Bits & 0x7fffff; - // Extract the exponent from the float value - const uint8_t Exp32 = (Bits & 0x7f800000) >> 23; - const int16_t Exp32Diff = Exp32 - 127; - - // initialize to 0, covers the case for 0 and small numbers - uint16_t Exp16 = 0, Frac16 = 0; - - if (__builtin_expect(Exp32Diff > 15, 0)) { - // Infinity and big numbers convert to infinity - Exp16 = 0x1f; - } else if (__builtin_expect(Exp32Diff > -14, 0)) { - // normal range for half type - Exp16 = static_cast(Exp32Diff) + 15; - // convert 23-bit mantissa to 10-bit mantissa. - Frac16 = static_cast(Frac32 >> 13); - // Round the mantissa as given in OpenCL spec section : 6.1.1.1 The half - // data type. - // Round to nearest. - uint32_t roundBits = Frac32 & 0x1fff; - uint32_t halfway = 0x1000; - if (roundBits > halfway) - Frac16 += 1; - // Tie to even. - else if (roundBits == halfway) - Frac16 += Frac16 & 1; - } else if (__builtin_expect(Exp32Diff > -25, 0)) { - // subnormals - Frac16 = static_cast((Frac32 | (uint32_t(1) << 23)) >> - (-Exp32Diff - 1)); - } - - if (__builtin_expect(Exp32 == 0xff && Frac32 != 0, 0)) { - // corner case: FP32 is NaN - Exp16 = 0x1F; - Frac16 = 0x200; - } - - // Compose the final FP16 binary - uint16_t Ret = 0; - Ret |= Sign; - Ret |= Exp16 << 10; - Ret += Frac16; // Add the carry bit from operation Frac16 += 1; - - return Ret; -} - -inline __SYCL_CONSTEXPR_HALF float half2Float(const uint16_t &Val) { - // Extract the sign from the bits. It is 1 if the sign is negative - const uint32_t Sign = static_cast(Val & 0x8000) << 16; - // Extract the exponent from the bits - const uint8_t Exp16 = (Val & 0x7c00) >> 10; - // Extract the fraction from the bits - uint16_t Frac16 = Val & 0x3ff; - - uint32_t Exp32 = 0; - if (__builtin_expect(Exp16 == 0x1f, 0)) { - Exp32 = 0xff; - } else if (__builtin_expect(Exp16 == 0, 0)) { - Exp32 = 0; - } else { - Exp32 = static_cast(Exp16) + 112; - } - // corner case: subnormal -> normal - // The denormal number of FP16 can be represented by FP32, therefore we need - // to recover the exponent and recalculate the fration. - if (__builtin_expect(Exp16 == 0 && Frac16 != 0, 0)) { - uint8_t OffSet = 0; - do { - ++OffSet; - Frac16 <<= 1; - } while ((Frac16 & 0x400) != 0x400); - // mask the 9th bit - Frac16 &= 0x3ff; - Exp32 = 113 - OffSet; - } - - uint32_t Frac32 = static_cast(Frac16 << 13); - - uint32_t Bits = 0; - Bits |= Sign; - Bits |= (Exp32 << 23); - Bits |= Frac32; - const float Result = sycl::bit_cast(Bits); - return Result; -} - -namespace half_impl { -class half; - -// Creation token to disambiguate constructors. -struct RawHostHalfToken { - constexpr explicit RawHostHalfToken(uint16_t Val) : Value{Val} {} - uint16_t Value; -}; - -#ifndef __SYCL_DEVICE_ONLY__ -class half { -#else -class [[__sycl_detail__::__uses_aspects__(aspect::fp16)]] half { -#endif -public: - half() = default; - constexpr half(const half &) = default; - constexpr half(half &&) = default; - -#ifdef __SYCL_DEVICE_ONLY__ - __SYCL_CONSTEXPR_HALF half(const float &rhs) : Data(rhs) {} -#else - __SYCL_CONSTEXPR_HALF half(const float &rhs) : Data(float2Half(rhs)) {} -#endif // __SYCL_DEVICE_ONLY__ - - constexpr half &operator=(const half &rhs) = default; - - // Operator +=, -=, *=, /= -#ifdef __SYCL_DEVICE_ONLY__ - __SYCL_CONSTEXPR_HALF half &operator+=(const half &rhs) { - Data += rhs.Data; - return *this; - } - - __SYCL_CONSTEXPR_HALF half &operator-=(const half &rhs) { - Data -= rhs.Data; - return *this; - } - - __SYCL_CONSTEXPR_HALF half &operator*=(const half &rhs) { - Data *= rhs.Data; - return *this; - } - - __SYCL_CONSTEXPR_HALF half &operator/=(const half &rhs) { - Data /= rhs.Data; - return *this; - } -#else - __SYCL_CONSTEXPR_HALF half &operator+=(const half &rhs) { - *this = operator float() + static_cast(rhs); - return *this; - } - - __SYCL_CONSTEXPR_HALF half &operator-=(const half &rhs) { - *this = operator float() - static_cast(rhs); - return *this; - } - - __SYCL_CONSTEXPR_HALF half &operator*=(const half &rhs) { - *this = operator float() * static_cast(rhs); - return *this; - } - - __SYCL_CONSTEXPR_HALF half &operator/=(const half &rhs) { - *this = operator float() / static_cast(rhs); - return *this; - } -#endif // __SYCL_DEVICE_ONLY__ - - // Operator ++, -- - __SYCL_CONSTEXPR_HALF half &operator++() { - *this += 1; - return *this; - } - - __SYCL_CONSTEXPR_HALF half operator++(int) { - half ret(*this); - operator++(); - return ret; - } - __SYCL_CONSTEXPR_HALF half &operator--() { - *this -= 1; - return *this; - } - - __SYCL_CONSTEXPR_HALF half operator--(int) { - half ret(*this); - operator--(); - return ret; - } - - // Operator neg -#ifdef __SYCL_DEVICE_ONLY__ - __SYCL_CONSTEXPR_HALF friend half operator-(const half other) { - return half(-other.Data); - } -#else - __SYCL_CONSTEXPR_HALF friend half operator-(const half other) { - return half(RawHostHalfToken(other.Data ^ 0x8000)); - } -#endif // __SYCL_DEVICE_ONLY__ - -// Operator +, -, *, / -#define OP(op, op_eq) \ - __SYCL_CONSTEXPR_HALF friend half operator op(const half lhs, \ - const half rhs) { \ - half rtn = lhs; \ - rtn op_eq rhs; \ - return rtn; \ - } \ - __SYCL_CONSTEXPR_HALF friend double operator op(const half lhs, \ - const double rhs) { \ - double rtn = lhs; \ - rtn op_eq rhs; \ - return rtn; \ - } \ - __SYCL_CONSTEXPR_HALF friend double operator op(const double lhs, \ - const half rhs) { \ - double rtn = lhs; \ - rtn op_eq rhs; \ - return rtn; \ - } \ - __SYCL_CONSTEXPR_HALF friend float operator op(const half lhs, \ - const float rhs) { \ - float rtn = lhs; \ - rtn op_eq rhs; \ - return rtn; \ - } \ - __SYCL_CONSTEXPR_HALF friend float operator op(const float lhs, \ - const half rhs) { \ - float rtn = lhs; \ - rtn op_eq rhs; \ - return rtn; \ - } \ - __SYCL_CONSTEXPR_HALF friend half operator op(const half lhs, \ - const int rhs) { \ - half rtn = lhs; \ - rtn op_eq half(static_cast(rhs)); \ - return rtn; \ - } \ - __SYCL_CONSTEXPR_HALF friend half operator op(const int lhs, \ - const half rhs) { \ - half rtn(static_cast(lhs)); \ - rtn op_eq rhs; \ - return rtn; \ - } \ - __SYCL_CONSTEXPR_HALF friend half operator op(const half lhs, \ - const long rhs) { \ - half rtn = lhs; \ - rtn op_eq half(static_cast(rhs)); \ - return rtn; \ - } \ - __SYCL_CONSTEXPR_HALF friend half operator op(const long lhs, \ - const half rhs) { \ - half rtn(static_cast(lhs)); \ - rtn op_eq rhs; \ - return rtn; \ - } \ - __SYCL_CONSTEXPR_HALF friend half operator op(const half lhs, \ - const long long rhs) { \ - half rtn = lhs; \ - rtn op_eq half(static_cast(rhs)); \ - return rtn; \ - } \ - __SYCL_CONSTEXPR_HALF friend half operator op(const long long lhs, \ - const half rhs) { \ - half rtn(static_cast(lhs)); \ - rtn op_eq rhs; \ - return rtn; \ - } \ - __SYCL_CONSTEXPR_HALF friend half operator op(const half &lhs, \ - const unsigned int &rhs) { \ - half rtn = lhs; \ - rtn op_eq half(static_cast(rhs)); \ - return rtn; \ - } \ - __SYCL_CONSTEXPR_HALF friend half operator op(const unsigned int &lhs, \ - const half &rhs) { \ - half rtn(static_cast(lhs)); \ - rtn op_eq rhs; \ - return rtn; \ - } \ - __SYCL_CONSTEXPR_HALF friend half operator op(const half &lhs, \ - const unsigned long &rhs) { \ - half rtn = lhs; \ - rtn op_eq half(static_cast(rhs)); \ - return rtn; \ - } \ - __SYCL_CONSTEXPR_HALF friend half operator op(const unsigned long &lhs, \ - const half &rhs) { \ - half rtn(static_cast(lhs)); \ - rtn op_eq rhs; \ - return rtn; \ - } \ - __SYCL_CONSTEXPR_HALF friend half operator op( \ - const half &lhs, const unsigned long long &rhs) { \ - half rtn = lhs; \ - rtn op_eq half(static_cast(rhs)); \ - return rtn; \ - } \ - __SYCL_CONSTEXPR_HALF friend half operator op(const unsigned long long &lhs, \ - const half &rhs) { \ - half rtn(static_cast(lhs)); \ - rtn op_eq rhs; \ - return rtn; \ - } - OP(+, +=) - OP(-, -=) - OP(*, *=) - OP(/, /=) - -#undef OP - -// Operator ==, !=, <, >, <=, >= -#define OP(op) \ - __SYCL_CONSTEXPR_HALF friend bool operator op(const half &lhs, \ - const half &rhs) { \ - return lhs.getFPRep() op rhs.getFPRep(); \ - } \ - __SYCL_CONSTEXPR_HALF friend bool operator op(const half &lhs, \ - const double &rhs) { \ - return lhs.getFPRep() op rhs; \ - } \ - __SYCL_CONSTEXPR_HALF friend bool operator op(const double &lhs, \ - const half &rhs) { \ - return lhs op rhs.getFPRep(); \ - } \ - __SYCL_CONSTEXPR_HALF friend bool operator op(const half &lhs, \ - const float &rhs) { \ - return lhs.getFPRep() op rhs; \ - } \ - __SYCL_CONSTEXPR_HALF friend bool operator op(const float &lhs, \ - const half &rhs) { \ - return lhs op rhs.getFPRep(); \ - } \ - __SYCL_CONSTEXPR_HALF friend bool operator op(const half &lhs, \ - const int &rhs) { \ - return lhs.getFPRep() op static_cast(rhs); \ - } \ - __SYCL_CONSTEXPR_HALF friend bool operator op(const int &lhs, \ - const half &rhs) { \ - return static_cast(lhs) op rhs.getFPRep(); \ - } \ - __SYCL_CONSTEXPR_HALF friend bool operator op(const half &lhs, \ - const long &rhs) { \ - return lhs.getFPRep() op static_cast(rhs); \ - } \ - __SYCL_CONSTEXPR_HALF friend bool operator op(const long &lhs, \ - const half &rhs) { \ - return static_cast(lhs) op rhs.getFPRep(); \ - } \ - __SYCL_CONSTEXPR_HALF friend bool operator op(const half &lhs, \ - const long long &rhs) { \ - return lhs.getFPRep() op static_cast(rhs); \ - } \ - __SYCL_CONSTEXPR_HALF friend bool operator op(const long long &lhs, \ - const half &rhs) { \ - return static_cast(lhs) op rhs.getFPRep(); \ - } \ - __SYCL_CONSTEXPR_HALF friend bool operator op(const half &lhs, \ - const unsigned int &rhs) { \ - return lhs.getFPRep() op static_cast(rhs); \ - } \ - __SYCL_CONSTEXPR_HALF friend bool operator op(const unsigned int &lhs, \ - const half &rhs) { \ - return static_cast(lhs) op rhs.getFPRep(); \ - } \ - __SYCL_CONSTEXPR_HALF friend bool operator op(const half &lhs, \ - const unsigned long &rhs) { \ - return lhs.getFPRep() op static_cast(rhs); \ - } \ - __SYCL_CONSTEXPR_HALF friend bool operator op(const unsigned long &lhs, \ - const half &rhs) { \ - return static_cast(lhs) op rhs.getFPRep(); \ - } \ - __SYCL_CONSTEXPR_HALF friend bool operator op( \ - const half &lhs, const unsigned long long &rhs) { \ - return lhs.getFPRep() op static_cast(rhs); \ - } \ - __SYCL_CONSTEXPR_HALF friend bool operator op(const unsigned long long &lhs, \ - const half &rhs) { \ - return static_cast(lhs) op rhs.getFPRep(); \ - } - OP(==) - OP(!=) - OP(<) - OP(>) - OP(<=) - OP(>=) - -#undef OP - - // Operator float #ifdef __SYCL_DEVICE_ONLY__ - __SYCL_CONSTEXPR_HALF operator float() const { - return static_cast(Data); - } +// std::istream/std::ostream aren't usable on device, so don't provide a +// definition to save compile time by using lightweight ``. +std::ostream &operator<<(std::ostream &O, sycl::half const &rhs); +std::istream &operator>>(std::istream &I, sycl::half &rhs); #else - __SYCL_CONSTEXPR_HALF operator float() const { return half2Float(Data); } -#endif // __SYCL_DEVICE_ONLY__ - - // Operator << and >> -#ifdef __SYCL_DEVICE_ONLY__ - // std::istream/std::ostream aren't usable on device, so don't provide a - // definition to save compile time by using lightweight ``. - inline friend std::ostream &operator<<(std::ostream &O, - sycl::half const &rhs); - inline friend std::istream &operator>>(std::istream &I, sycl::half &rhs); -#else - inline friend std::ostream &operator<<(std::ostream &O, - sycl::half const &rhs) { - O << static_cast(rhs); - return O; - } - - inline friend std::istream &operator>>(std::istream &I, sycl::half &rhs) { - float ValFloat = 0.0f; - I >> ValFloat; - rhs = ValFloat; - return I; - } -#endif - - template friend struct std::hash; - - friend class sycl::ext::intel::esimd::detail::WrapperElementTypeProxy; - -private: - // When doing operations, we cannot simply work with Data on host as - // it is an integer. Instead, convert it to float. On device we can work with - // Data as it is already a floating point representation. -#ifdef __SYCL_DEVICE_ONLY__ - __SYCL_CONSTEXPR_HALF StorageT getFPRep() const { return Data; } -#else - __SYCL_CONSTEXPR_HALF float getFPRep() const { return operator float(); } -#endif - -#ifndef __SYCL_DEVICE_ONLY__ - // Because sycl::bit_cast might not be constexpr on certain systems, - // implementation needs shortcut for creating a host sycl::half directly from - // a uint16_t representation. - constexpr explicit half(RawHostHalfToken X) : Data(X.Value) {} - - friend constexpr inline half CreateHostHalfRaw(uint16_t X); -#endif // __SYCL_DEVICE_ONLY__ - - StorageT Data; -}; - -#ifndef __SYCL_DEVICE_ONLY__ -constexpr inline half CreateHostHalfRaw(uint16_t X) { - return half(RawHostHalfToken(X)); +inline std::ostream &operator<<(std::ostream &O, sycl::half const &rhs) { + O << static_cast(rhs); + return O; } -#endif // __SYCL_DEVICE_ONLY__ -} // namespace half_impl - -// According to the C++ standard, math functions from cmath/math.h should work -// only on arithmetic types. We can't specify half type as arithmetic/floating -// point(via std::is_floating_point) since only float, double and long double -// types are "floating point" according to the standard. In order to use half -// type with these math functions we cast half to float using template -// function helper. -template inline T cast_if_host_half(T val) { return val; } -inline float cast_if_host_half(half_impl::half val) { - return static_cast(val); +inline std::istream &operator>>(std::istream &I, sycl::half &rhs) { + float ValFloat = 0.0f; + I >> ValFloat; + rhs = ValFloat; + return I; } +#endif -} // namespace detail - +} // namespace detail::half_impl } // namespace _V1 } // namespace sycl @@ -537,78 +57,4 @@ template <> struct hash { } }; -// Partial specialization of `std::numeric` -template <> struct numeric_limits { - // All following values are either calculated based on description of each - // function/value on https://en.cppreference.com/w/cpp/types/numeric_limits, - // or cl_platform.h. - static constexpr bool is_specialized = true; - static constexpr bool is_signed = true; - static constexpr bool is_integer = false; - static constexpr bool is_exact = false; - static constexpr bool has_infinity = true; - static constexpr bool has_quiet_NaN = true; - static constexpr bool has_signaling_NaN = true; - static constexpr float_denorm_style has_denorm = denorm_present; - static constexpr bool has_denorm_loss = false; - static constexpr bool tinyness_before = false; - static constexpr bool traps = false; - static constexpr int max_exponent10 = 4; - static constexpr int max_exponent = 16; - static constexpr int min_exponent10 = -4; - static constexpr int min_exponent = -13; - static constexpr int radix = 2; - static constexpr int max_digits10 = 5; - static constexpr int digits = 11; - static constexpr bool is_bounded = true; - static constexpr int digits10 = 3; - static constexpr bool is_modulo = false; - static constexpr bool is_iec559 = true; - static constexpr float_round_style round_style = round_to_nearest; - - static __SYCL_CONSTEXPR_HALF const sycl::half(min)() noexcept { - return 6.103515625e-05f; // half minimum value - } - - static __SYCL_CONSTEXPR_HALF const sycl::half(max)() noexcept { - return 65504.0f; // half maximum value - } - - static __SYCL_CONSTEXPR_HALF const sycl::half lowest() noexcept { - return -65504.0f; // -1*(half maximum value) - } - - static __SYCL_CONSTEXPR_HALF const sycl::half epsilon() noexcept { - return 9.765625e-04f; // half epsilon - } - - static __SYCL_CONSTEXPR_HALF const sycl::half round_error() noexcept { - return 0.5f; - } - - static constexpr const sycl::half infinity() noexcept { -#ifdef __SYCL_DEVICE_ONLY__ - return __builtin_huge_valf(); -#else - return sycl::detail::half_impl::CreateHostHalfRaw( - static_cast(0x7C00)); -#endif - } - - static __SYCL_CONSTEXPR_HALF const sycl::half quiet_NaN() noexcept { - return __builtin_nanf(""); - } - - static __SYCL_CONSTEXPR_HALF const sycl::half signaling_NaN() noexcept { - return __builtin_nansf(""); - } - - static __SYCL_CONSTEXPR_HALF const sycl::half denorm_min() noexcept { - return 5.96046e-08f; - } -}; - } // namespace std - -#undef __SYCL_CONSTEXPR_HALF -#undef _CPP14_CONSTEXPR diff --git a/sycl/include/sycl/known_identity.hpp b/sycl/include/sycl/known_identity.hpp index 49a4305508c3f..2db2be9cdf702 100644 --- a/sycl/include/sycl/known_identity.hpp +++ b/sycl/include/sycl/known_identity.hpp @@ -10,8 +10,8 @@ #include // for half #include // for is_genbool, is_genint... +#include // for half #include // for bit_and, bit_or, bit_xor -#include // for half #include // for marray #include // for vec diff --git a/sycl/test/include_deps/sycl_accessor.hpp.cpp b/sycl/test/include_deps/sycl_accessor.hpp.cpp index 6676700cc16a3..549f833b79e9d 100644 --- a/sycl/test/include_deps/sycl_accessor.hpp.cpp +++ b/sycl/test/include_deps/sycl_accessor.hpp.cpp @@ -71,6 +71,7 @@ // CHECK-NEXT: ext/oneapi/matrix/query-types.hpp // CHECK-NEXT: ext/oneapi/bfloat16.hpp // CHECK-NEXT: half_type.hpp +// CHECK-NEXT: detail/half_type_impl.hpp // CHECK-NEXT: ext/oneapi/matrix/matrix-unified-utils.hpp // CHECK-NEXT: info/platform_traits.def // CHECK-NEXT: info/context_traits.def diff --git a/sycl/test/include_deps/sycl_detail_core.hpp.cpp b/sycl/test/include_deps/sycl_detail_core.hpp.cpp index 2b8ac26da9ddc..4b2ed7b5f557c 100644 --- a/sycl/test/include_deps/sycl_detail_core.hpp.cpp +++ b/sycl/test/include_deps/sycl_detail_core.hpp.cpp @@ -72,6 +72,7 @@ // CHECK-NEXT: ext/oneapi/matrix/query-types.hpp // CHECK-NEXT: ext/oneapi/bfloat16.hpp // CHECK-NEXT: half_type.hpp +// CHECK-NEXT: detail/half_type_impl.hpp // CHECK-NEXT: ext/oneapi/matrix/matrix-unified-utils.hpp // CHECK-NEXT: info/platform_traits.def // CHECK-NEXT: info/context_traits.def diff --git a/sycl/test/include_deps/sycl_khr_includes_accessor.hpp.cpp b/sycl/test/include_deps/sycl_khr_includes_accessor.hpp.cpp index 19b91fe077b3b..adf21e5518287 100644 --- a/sycl/test/include_deps/sycl_khr_includes_accessor.hpp.cpp +++ b/sycl/test/include_deps/sycl_khr_includes_accessor.hpp.cpp @@ -74,6 +74,7 @@ // CHECK-NEXT: ext/oneapi/matrix/query-types.hpp // CHECK-NEXT: ext/oneapi/bfloat16.hpp // CHECK-NEXT: half_type.hpp +// CHECK-NEXT: detail/half_type_impl.hpp // CHECK-NEXT: ext/oneapi/matrix/matrix-unified-utils.hpp // CHECK-NEXT: info/platform_traits.def // CHECK-NEXT: info/context_traits.def diff --git a/sycl/test/include_deps/sycl_khr_includes_context.hpp.cpp b/sycl/test/include_deps/sycl_khr_includes_context.hpp.cpp index 9f10455466756..9231c65b08d09 100644 --- a/sycl/test/include_deps/sycl_khr_includes_context.hpp.cpp +++ b/sycl/test/include_deps/sycl_khr_includes_context.hpp.cpp @@ -36,6 +36,7 @@ // CHECK-NEXT: aliases.hpp // CHECK-NEXT: bit_cast.hpp // CHECK-NEXT: half_type.hpp +// CHECK-NEXT: detail/half_type_impl.hpp // CHECK-NEXT: detail/fwd/half.hpp // CHECK-NEXT: ext/oneapi/matrix/matrix-unified-utils.hpp // CHECK-NEXT: __spirv/spirv_types.hpp diff --git a/sycl/test/include_deps/sycl_khr_includes_device.hpp.cpp b/sycl/test/include_deps/sycl_khr_includes_device.hpp.cpp index cd1f28df83872..e5a91c198eb36 100644 --- a/sycl/test/include_deps/sycl_khr_includes_device.hpp.cpp +++ b/sycl/test/include_deps/sycl_khr_includes_device.hpp.cpp @@ -35,6 +35,7 @@ // CHECK-NEXT: aliases.hpp // CHECK-NEXT: bit_cast.hpp // CHECK-NEXT: half_type.hpp +// CHECK-NEXT: detail/half_type_impl.hpp // CHECK-NEXT: detail/fwd/half.hpp // CHECK-NEXT: ext/oneapi/matrix/matrix-unified-utils.hpp // CHECK-NEXT: __spirv/spirv_types.hpp diff --git a/sycl/test/include_deps/sycl_khr_includes_event.hpp.cpp b/sycl/test/include_deps/sycl_khr_includes_event.hpp.cpp index 7cb44b891d9fd..16453f24326bc 100644 --- a/sycl/test/include_deps/sycl_khr_includes_event.hpp.cpp +++ b/sycl/test/include_deps/sycl_khr_includes_event.hpp.cpp @@ -35,6 +35,7 @@ // CHECK-NEXT: aliases.hpp // CHECK-NEXT: bit_cast.hpp // CHECK-NEXT: half_type.hpp +// CHECK-NEXT: detail/half_type_impl.hpp // CHECK-NEXT: detail/fwd/half.hpp // CHECK-NEXT: ext/oneapi/matrix/matrix-unified-utils.hpp // CHECK-NEXT: __spirv/spirv_types.hpp diff --git a/sycl/test/include_deps/sycl_khr_includes_group_algorithms.hpp.cpp b/sycl/test/include_deps/sycl_khr_includes_group_algorithms.hpp.cpp index 7936929596a17..241a1df7abd2a 100644 --- a/sycl/test/include_deps/sycl_khr_includes_group_algorithms.hpp.cpp +++ b/sycl/test/include_deps/sycl_khr_includes_group_algorithms.hpp.cpp @@ -37,6 +37,7 @@ // CHECK-NEXT: device_event.hpp // CHECK-NEXT: pointers.hpp // CHECK-NEXT: half_type.hpp +// CHECK-NEXT: detail/half_type_impl.hpp // CHECK-NEXT: bit_cast.hpp // CHECK-NEXT: aspects.hpp // CHECK-NEXT: info/aspects.def diff --git a/sycl/test/include_deps/sycl_khr_includes_groups.hpp.cpp b/sycl/test/include_deps/sycl_khr_includes_groups.hpp.cpp index 931fed70bc7f1..012a7b26dad7f 100644 --- a/sycl/test/include_deps/sycl_khr_includes_groups.hpp.cpp +++ b/sycl/test/include_deps/sycl_khr_includes_groups.hpp.cpp @@ -37,6 +37,7 @@ // CHECK-NEXT: detail/string.hpp // CHECK-NEXT: functional.hpp // CHECK-NEXT: half_type.hpp +// CHECK-NEXT: detail/half_type_impl.hpp // CHECK-NEXT: bit_cast.hpp // CHECK-NEXT: aspects.hpp // CHECK-NEXT: info/aspects.def diff --git a/sycl/test/include_deps/sycl_khr_includes_half.hpp.cpp b/sycl/test/include_deps/sycl_khr_includes_half.hpp.cpp index f6026b94c8697..5e96570a1abf8 100644 --- a/sycl/test/include_deps/sycl_khr_includes_half.hpp.cpp +++ b/sycl/test/include_deps/sycl_khr_includes_half.hpp.cpp @@ -9,8 +9,8 @@ // CHECK-NEXT: detail/defines_elementary.hpp // CHECK-NEXT: feature_test.hpp // CHECK-NEXT: half_type.hpp +// CHECK-NEXT: detail/half_type_impl.hpp // CHECK-NEXT: bit_cast.hpp -// CHECK-NEXT: detail/export.hpp // CHECK-NEXT: detail/fwd/half.hpp // CHECK-NEXT: aspects.hpp // CHECK-NEXT: detail/defines.hpp diff --git a/sycl/test/include_deps/sycl_khr_includes_handler.hpp.cpp b/sycl/test/include_deps/sycl_khr_includes_handler.hpp.cpp index a4c53ca4f9439..41a55bb6bc72e 100644 --- a/sycl/test/include_deps/sycl_khr_includes_handler.hpp.cpp +++ b/sycl/test/include_deps/sycl_khr_includes_handler.hpp.cpp @@ -75,6 +75,7 @@ // CHECK-NEXT: ext/oneapi/matrix/query-types.hpp // CHECK-NEXT: ext/oneapi/bfloat16.hpp // CHECK-NEXT: half_type.hpp +// CHECK-NEXT: detail/half_type_impl.hpp // CHECK-NEXT: ext/oneapi/matrix/matrix-unified-utils.hpp // CHECK-NEXT: info/platform_traits.def // CHECK-NEXT: info/context_traits.def diff --git a/sycl/test/include_deps/sycl_khr_includes_images.hpp.cpp b/sycl/test/include_deps/sycl_khr_includes_images.hpp.cpp index 07dfaf2272174..83de974206575 100644 --- a/sycl/test/include_deps/sycl_khr_includes_images.hpp.cpp +++ b/sycl/test/include_deps/sycl_khr_includes_images.hpp.cpp @@ -75,6 +75,7 @@ // CHECK-NEXT: ext/oneapi/matrix/query-types.hpp // CHECK-NEXT: ext/oneapi/bfloat16.hpp // CHECK-NEXT: half_type.hpp +// CHECK-NEXT: detail/half_type_impl.hpp // CHECK-NEXT: ext/oneapi/matrix/matrix-unified-utils.hpp // CHECK-NEXT: info/platform_traits.def // CHECK-NEXT: info/context_traits.def diff --git a/sycl/test/include_deps/sycl_khr_includes_interop_handle.hpp.cpp b/sycl/test/include_deps/sycl_khr_includes_interop_handle.hpp.cpp index eacb3c8070d1c..b32f52fe080c0 100644 --- a/sycl/test/include_deps/sycl_khr_includes_interop_handle.hpp.cpp +++ b/sycl/test/include_deps/sycl_khr_includes_interop_handle.hpp.cpp @@ -75,6 +75,7 @@ // CHECK-NEXT: ext/oneapi/matrix/query-types.hpp // CHECK-NEXT: ext/oneapi/bfloat16.hpp // CHECK-NEXT: half_type.hpp +// CHECK-NEXT: detail/half_type_impl.hpp // CHECK-NEXT: ext/oneapi/matrix/matrix-unified-utils.hpp // CHECK-NEXT: info/platform_traits.def // CHECK-NEXT: info/context_traits.def diff --git a/sycl/test/include_deps/sycl_khr_includes_kernel_bundle.hpp.cpp b/sycl/test/include_deps/sycl_khr_includes_kernel_bundle.hpp.cpp index b82f38f1cfdd6..6ca81fb4900ad 100644 --- a/sycl/test/include_deps/sycl_khr_includes_kernel_bundle.hpp.cpp +++ b/sycl/test/include_deps/sycl_khr_includes_kernel_bundle.hpp.cpp @@ -35,6 +35,7 @@ // CHECK-NEXT: aliases.hpp // CHECK-NEXT: bit_cast.hpp // CHECK-NEXT: half_type.hpp +// CHECK-NEXT: detail/half_type_impl.hpp // CHECK-NEXT: detail/fwd/half.hpp // CHECK-NEXT: ext/oneapi/matrix/matrix-unified-utils.hpp // CHECK-NEXT: __spirv/spirv_types.hpp diff --git a/sycl/test/include_deps/sycl_khr_includes_math.hpp.cpp b/sycl/test/include_deps/sycl_khr_includes_math.hpp.cpp index cbf657caec2c8..b47d43b3501c8 100644 --- a/sycl/test/include_deps/sycl_khr_includes_math.hpp.cpp +++ b/sycl/test/include_deps/sycl_khr_includes_math.hpp.cpp @@ -12,6 +12,13 @@ // CHECK-NEXT: detail/builtins/builtins.hpp // CHECK-NEXT: detail/fwd/multi_ptr.hpp // CHECK-NEXT: access/access.hpp +// CHECK-NEXT: detail/half_type_impl.hpp +// CHECK-NEXT: bit_cast.hpp +// CHECK-NEXT: detail/fwd/half.hpp +// CHECK-NEXT: aspects.hpp +// CHECK-NEXT: detail/defines.hpp +// CHECK-NEXT: info/aspects.def +// CHECK-NEXT: info/aspects_deprecated.def // CHECK-NEXT: detail/loop.hpp // CHECK-NEXT: detail/type_traits.hpp // CHECK-NEXT: detail/type_traits/bool_traits.hpp @@ -20,22 +27,15 @@ // CHECK-NEXT: detail/named_swizzles_mixin.hpp // CHECK-NEXT: detail/vector_traits.hpp // CHECK-NEXT: aliases.hpp -// CHECK-NEXT: bit_cast.hpp // CHECK-NEXT: detail/common.hpp // CHECK-NEXT: detail/assert.hpp // CHECK-NEXT: __spirv/spirv_vars.hpp // CHECK-NEXT: detail/export.hpp // CHECK-NEXT: detail/nd_loop.hpp // CHECK-NEXT: detail/fwd/accessor.hpp -// CHECK-NEXT: detail/defines.hpp -// CHECK-NEXT: detail/fwd/half.hpp // CHECK-NEXT: detail/generic_type_traits.hpp // CHECK-NEXT: detail/type_traits/integer_traits.hpp // CHECK-NEXT: detail/memcpy.hpp -// CHECK-NEXT: half_type.hpp -// CHECK-NEXT: aspects.hpp -// CHECK-NEXT: info/aspects.def -// CHECK-NEXT: info/aspects_deprecated.def // CHECK-NEXT: marray.hpp // CHECK-NEXT: detail/builtins/common_functions.inc // CHECK-NEXT: detail/builtins/helper_macros.hpp diff --git a/sycl/test/include_deps/sycl_khr_includes_platform.hpp.cpp b/sycl/test/include_deps/sycl_khr_includes_platform.hpp.cpp index 09e00c5a240bb..40c2e129fb113 100644 --- a/sycl/test/include_deps/sycl_khr_includes_platform.hpp.cpp +++ b/sycl/test/include_deps/sycl_khr_includes_platform.hpp.cpp @@ -35,6 +35,7 @@ // CHECK-NEXT: aliases.hpp // CHECK-NEXT: bit_cast.hpp // CHECK-NEXT: half_type.hpp +// CHECK-NEXT: detail/half_type_impl.hpp // CHECK-NEXT: detail/fwd/half.hpp // CHECK-NEXT: ext/oneapi/matrix/matrix-unified-utils.hpp // CHECK-NEXT: __spirv/spirv_types.hpp diff --git a/sycl/test/include_deps/sycl_khr_includes_queue.hpp.cpp b/sycl/test/include_deps/sycl_khr_includes_queue.hpp.cpp index 8ddf2a9576ed0..d55ffd18b5702 100644 --- a/sycl/test/include_deps/sycl_khr_includes_queue.hpp.cpp +++ b/sycl/test/include_deps/sycl_khr_includes_queue.hpp.cpp @@ -77,6 +77,7 @@ // CHECK-NEXT: ext/oneapi/matrix/query-types.hpp // CHECK-NEXT: ext/oneapi/bfloat16.hpp // CHECK-NEXT: half_type.hpp +// CHECK-NEXT: detail/half_type_impl.hpp // CHECK-NEXT: ext/oneapi/matrix/matrix-unified-utils.hpp // CHECK-NEXT: info/platform_traits.def // CHECK-NEXT: info/context_traits.def diff --git a/sycl/test/include_deps/sycl_khr_includes_reduction.hpp.cpp b/sycl/test/include_deps/sycl_khr_includes_reduction.hpp.cpp index 6775a70dbd859..71f1416277a22 100644 --- a/sycl/test/include_deps/sycl_khr_includes_reduction.hpp.cpp +++ b/sycl/test/include_deps/sycl_khr_includes_reduction.hpp.cpp @@ -77,6 +77,7 @@ // CHECK-NEXT: ext/oneapi/matrix/query-types.hpp // CHECK-NEXT: ext/oneapi/bfloat16.hpp // CHECK-NEXT: half_type.hpp +// CHECK-NEXT: detail/half_type_impl.hpp // CHECK-NEXT: ext/oneapi/matrix/matrix-unified-utils.hpp // CHECK-NEXT: info/platform_traits.def // CHECK-NEXT: info/context_traits.def diff --git a/sycl/test/include_deps/sycl_khr_includes_stream.hpp.cpp b/sycl/test/include_deps/sycl_khr_includes_stream.hpp.cpp index 8b69fc29fe202..abc80a216df5c 100644 --- a/sycl/test/include_deps/sycl_khr_includes_stream.hpp.cpp +++ b/sycl/test/include_deps/sycl_khr_includes_stream.hpp.cpp @@ -75,6 +75,7 @@ // CHECK-NEXT: ext/oneapi/matrix/query-types.hpp // CHECK-NEXT: ext/oneapi/bfloat16.hpp // CHECK-NEXT: half_type.hpp +// CHECK-NEXT: detail/half_type_impl.hpp // CHECK-NEXT: ext/oneapi/matrix/matrix-unified-utils.hpp // CHECK-NEXT: info/platform_traits.def // CHECK-NEXT: info/context_traits.def diff --git a/sycl/test/include_deps/sycl_khr_includes_usm.hpp.cpp b/sycl/test/include_deps/sycl_khr_includes_usm.hpp.cpp index 10d253de26117..019772ccdbba9 100644 --- a/sycl/test/include_deps/sycl_khr_includes_usm.hpp.cpp +++ b/sycl/test/include_deps/sycl_khr_includes_usm.hpp.cpp @@ -13,6 +13,13 @@ // CHECK-NEXT: detail/builtins/builtins.hpp // CHECK-NEXT: detail/fwd/multi_ptr.hpp // CHECK-NEXT: access/access.hpp +// CHECK-NEXT: detail/half_type_impl.hpp +// CHECK-NEXT: bit_cast.hpp +// CHECK-NEXT: detail/fwd/half.hpp +// CHECK-NEXT: aspects.hpp +// CHECK-NEXT: detail/defines.hpp +// CHECK-NEXT: info/aspects.def +// CHECK-NEXT: info/aspects_deprecated.def // CHECK-NEXT: detail/loop.hpp // CHECK-NEXT: detail/type_traits.hpp // CHECK-NEXT: detail/type_traits/bool_traits.hpp @@ -21,22 +28,15 @@ // CHECK-NEXT: detail/named_swizzles_mixin.hpp // CHECK-NEXT: detail/vector_traits.hpp // CHECK-NEXT: aliases.hpp -// CHECK-NEXT: bit_cast.hpp // CHECK-NEXT: detail/common.hpp // CHECK-NEXT: detail/assert.hpp // CHECK-NEXT: __spirv/spirv_vars.hpp // CHECK-NEXT: detail/export.hpp // CHECK-NEXT: detail/nd_loop.hpp // CHECK-NEXT: detail/fwd/accessor.hpp -// CHECK-NEXT: detail/defines.hpp -// CHECK-NEXT: detail/fwd/half.hpp // CHECK-NEXT: detail/generic_type_traits.hpp // CHECK-NEXT: detail/type_traits/integer_traits.hpp // CHECK-NEXT: detail/memcpy.hpp -// CHECK-NEXT: half_type.hpp -// CHECK-NEXT: aspects.hpp -// CHECK-NEXT: info/aspects.def -// CHECK-NEXT: info/aspects_deprecated.def // CHECK-NEXT: marray.hpp // CHECK-NEXT: detail/builtins/common_functions.inc // CHECK-NEXT: detail/builtins/helper_macros.hpp @@ -60,6 +60,7 @@ // CHECK-NEXT: ext/oneapi/experimental/forward_progress.hpp // CHECK-NEXT: ext/oneapi/matrix/query-types.hpp // CHECK-NEXT: ext/oneapi/bfloat16.hpp +// CHECK-NEXT: half_type.hpp // CHECK-NEXT: ext/oneapi/matrix/matrix-unified-utils.hpp // CHECK-NEXT: __spirv/spirv_types.hpp // CHECK-NEXT: info/platform_traits.def From 98fd37cfd38aad7a99b186dca74e1e1e310ac87f Mon Sep 17 00:00:00 2001 From: y Date: Tue, 21 Apr 2026 10:46:09 -0700 Subject: [PATCH 2/2] Fix CI --- sycl/source/builtins/integer_functions.cpp | 2 ++ 1 file changed, 2 insertions(+) diff --git a/sycl/source/builtins/integer_functions.cpp b/sycl/source/builtins/integer_functions.cpp index 862aa97df3f0f..37fe0a9f8029c 100644 --- a/sycl/source/builtins/integer_functions.cpp +++ b/sycl/source/builtins/integer_functions.cpp @@ -9,6 +9,8 @@ #include #include +#include + #include "host_helper_macros.hpp" namespace {