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 e2d3f6c63a56a..e7866303c11ed 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 @@ -21,16 +21,11 @@ inline namespace _V1 { namespace ext::intel::esimd::detail { // Standalone definitions to use w/o instantiating element_type_traits. -#ifdef __SYCL_DEVICE_ONLY__ -// Can't use sycl::detail::half_impl::StorageT as RawT for both host and -// device as it still maps to struct on/ host (even though the struct is a -// trivial wrapper around uint16_t), and for ESIMD we need a type which can be -// an element of clang vector. using half_raw_type = sycl::detail::half_impl::StorageT; +#ifdef __SYCL_DEVICE_ONLY__ // On device, _Float16 is native Cpp type, so it is the enclosing C++ type using half_enclosing_cpp_type = half_raw_type; #else -using half_raw_type = uint16_t; using half_enclosing_cpp_type = float; #endif // __SYCL_DEVICE_ONLY__ @@ -86,11 +81,7 @@ template struct vector_conversion_traits { class WrapperElementTypeProxy { public: static ESIMD_INLINE half_raw_type bitcast_to_raw_scalar(sycl::half Val) { -#ifdef __SYCL_DEVICE_ONLY__ return Val.Data; -#else - return Val.Data.Buf; -#endif // __SYCL_DEVICE_ONLY__ } static ESIMD_INLINE sycl::half bitcast_to_wrapper_scalar(half_raw_type Val) { diff --git a/sycl/include/sycl/half_type.hpp b/sycl/include/sycl/half_type.hpp index 799ff9fb186e9..f97fca761f9ee 100644 --- a/sycl/include/sycl/half_type.hpp +++ b/sycl/include/sycl/half_type.hpp @@ -144,93 +144,14 @@ inline __SYCL_CONSTEXPR_HALF float half2Float(const uint16_t &Val) { return Result; } -namespace host_half_impl { - -// The main host half class -class __SYCL_EXPORT half { -public: - half() = default; - constexpr half(const half &) = default; - constexpr half(half &&) = default; - - __SYCL_CONSTEXPR_HALF half(const float &rhs) : Buf(float2Half(rhs)) {} - - constexpr half &operator=(const half &rhs) = default; - - // Operator +=, -=, *=, /= - __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; - } - - // 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 - constexpr half &operator-() { - Buf ^= 0x8000; - return *this; - } - - // Operator float - __SYCL_CONSTEXPR_HALF operator float() const { return half2Float(Buf); } - - template friend struct std::hash; - - // Initialize underlying data - constexpr explicit half(uint16_t x) : Buf(x) {} - - friend class sycl::ext::intel::esimd::detail::WrapperElementTypeProxy; - -private: - uint16_t Buf; -}; - -} // namespace host_half_impl - namespace half_impl { class half; // Several aliases are defined below: // - StorageT: actual representation of half data type. It is used by scalar // half values. On device side, it points to some native half data type, while -// on host some custom data type is used to emulate operations of 16-bit -// floating-point values +// on host it is represented by a 16-bit integer that the implementation +// manipulates to emulate half-precision floating-point behavior. // // - BIsRepresentationT: data type which is used by built-in functions. It is // distinguished from StorageT, because on host, we can still operate on the @@ -258,7 +179,7 @@ using Vec16StorageT = VecElemT __attribute__((ext_vector_type(16))); #endif // __INTEL_PREVIEW_BREAKING_CHANGES #else // SYCL_DEVICE_ONLY -using StorageT = detail::host_half_impl::half; +using StorageT = uint16_t; // No need to extract underlying data type for built-in functions operating on // host using BIsRepresentationT = half; @@ -278,6 +199,12 @@ using Vec16StorageT = std::array; #endif // SYCL_DEVICE_ONLY +// 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 @@ -288,18 +215,16 @@ class [[__sycl_detail__::__uses_aspects__(aspect::fp16)]] half { 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; -#ifndef __SYCL_DEVICE_ONLY__ - // Since StorageT and BIsRepresentationT are different on host, these two - // helpers are required for 'vec' class - constexpr half(const detail::host_half_impl::half &rhs) : Data(rhs) {} - constexpr operator detail::host_half_impl::half() const { return Data; } -#endif // __SYCL_DEVICE_ONLY__ - // Operator +=, -=, *=, /= +#ifdef __SYCL_DEVICE_ONLY__ __SYCL_CONSTEXPR_HALF half &operator+=(const half &rhs) { Data += rhs.Data; return *this; @@ -319,6 +244,27 @@ class [[__sycl_detail__::__uses_aspects__(aspect::fp16)]] half { 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++() { @@ -342,9 +288,17 @@ class [[__sycl_detail__::__uses_aspects__(aspect::fp16)]] half { 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) \ @@ -461,71 +415,71 @@ class [[__sycl_detail__::__uses_aspects__(aspect::fp16)]] half { #define OP(op) \ __SYCL_CONSTEXPR_HALF friend bool operator op(const half &lhs, \ const half &rhs) { \ - return lhs.Data op rhs.Data; \ + return lhs.getFPRep() op rhs.getFPRep(); \ } \ __SYCL_CONSTEXPR_HALF friend bool operator op(const half &lhs, \ const double &rhs) { \ - return lhs.Data op rhs; \ + return lhs.getFPRep() op rhs; \ } \ __SYCL_CONSTEXPR_HALF friend bool operator op(const double &lhs, \ const half &rhs) { \ - return lhs op rhs.Data; \ + return lhs op rhs.getFPRep(); \ } \ __SYCL_CONSTEXPR_HALF friend bool operator op(const half &lhs, \ const float &rhs) { \ - return lhs.Data op rhs; \ + return lhs.getFPRep() op rhs; \ } \ __SYCL_CONSTEXPR_HALF friend bool operator op(const float &lhs, \ const half &rhs) { \ - return lhs op rhs.Data; \ + return lhs op rhs.getFPRep(); \ } \ __SYCL_CONSTEXPR_HALF friend bool operator op(const half &lhs, \ const int &rhs) { \ - return lhs.Data op rhs; \ + return lhs.getFPRep() op rhs; \ } \ __SYCL_CONSTEXPR_HALF friend bool operator op(const int &lhs, \ const half &rhs) { \ - return lhs op rhs.Data; \ + return lhs op rhs.getFPRep(); \ } \ __SYCL_CONSTEXPR_HALF friend bool operator op(const half &lhs, \ const long &rhs) { \ - return lhs.Data op rhs; \ + return lhs.getFPRep() op rhs; \ } \ __SYCL_CONSTEXPR_HALF friend bool operator op(const long &lhs, \ const half &rhs) { \ - return lhs op rhs.Data; \ + return lhs op rhs.getFPRep(); \ } \ __SYCL_CONSTEXPR_HALF friend bool operator op(const half &lhs, \ const long long &rhs) { \ - return lhs.Data op rhs; \ + return lhs.getFPRep() op rhs; \ } \ __SYCL_CONSTEXPR_HALF friend bool operator op(const long long &lhs, \ const half &rhs) { \ - return lhs op rhs.Data; \ + return lhs op rhs.getFPRep(); \ } \ __SYCL_CONSTEXPR_HALF friend bool operator op(const half &lhs, \ const unsigned int &rhs) { \ - return lhs.Data op rhs; \ + return lhs.getFPRep() op rhs; \ } \ __SYCL_CONSTEXPR_HALF friend bool operator op(const unsigned int &lhs, \ const half &rhs) { \ - return lhs op rhs.Data; \ + return lhs op rhs.getFPRep(); \ } \ __SYCL_CONSTEXPR_HALF friend bool operator op(const half &lhs, \ const unsigned long &rhs) { \ - return lhs.Data op rhs; \ + return lhs.getFPRep() op rhs; \ } \ __SYCL_CONSTEXPR_HALF friend bool operator op(const unsigned long &lhs, \ const half &rhs) { \ - return lhs op rhs.Data; \ + return lhs op rhs.getFPRep(); \ } \ __SYCL_CONSTEXPR_HALF friend bool operator op( \ const half &lhs, const unsigned long long &rhs) { \ - return lhs.Data op rhs; \ + return lhs.getFPRep() op rhs; \ } \ __SYCL_CONSTEXPR_HALF friend bool operator op(const unsigned long long &lhs, \ const half &rhs) { \ - return lhs op rhs.Data; \ + return lhs op rhs.getFPRep(); \ } OP(==) OP(!=) @@ -537,9 +491,13 @@ class [[__sycl_detail__::__uses_aspects__(aspect::fp16)]] half { #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__ // Operator << and >> inline friend std::ostream &operator<<(std::ostream &O, @@ -560,8 +518,32 @@ class [[__sycl_detail__::__uses_aspects__(aspect::fp16)]] half { 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 @@ -644,7 +626,8 @@ template <> struct numeric_limits { #ifdef __SYCL_DEVICE_ONLY__ return __builtin_huge_valf(); #else - return sycl::detail::host_half_impl::half(static_cast(0x7C00)); + return sycl::detail::half_impl::CreateHostHalfRaw( + static_cast(0x7C00)); #endif } diff --git a/sycl/include/sycl/known_identity.hpp b/sycl/include/sycl/known_identity.hpp index edb0537990ccf..32575b94faccf 100644 --- a/sycl/include/sycl/known_identity.hpp +++ b/sycl/include/sycl/known_identity.hpp @@ -187,7 +187,7 @@ struct known_identity_impl< #ifdef __SYCL_DEVICE_ONLY__ 0; #else - sycl::detail::host_half_impl::half(static_cast(0)); + sycl::detail::half_impl::CreateHostHalfRaw(static_cast(0)); #endif }; @@ -227,7 +227,7 @@ struct known_identity_impl< #ifdef __SYCL_DEVICE_ONLY__ 1; #else - sycl::detail::host_half_impl::half(static_cast(0x3C00)); + sycl::detail::half_impl::CreateHostHalfRaw(static_cast(0x3C00)); #endif }; diff --git a/sycl/include/sycl/vector_preview.hpp b/sycl/include/sycl/vector_preview.hpp index d0549dbefd817..527b40deb7461 100644 --- a/sycl/include/sycl/vector_preview.hpp +++ b/sycl/include/sycl/vector_preview.hpp @@ -93,18 +93,8 @@ class SwizzleOp; template class GetOp { public: using DataT = T; - DataT getValue(size_t) const { - if constexpr (std::is_same_v) - return DataT{0.0f}; - else - return (DataT)0; - } - DataT operator()(DataT, DataT) { - if constexpr (std::is_same_v) - return DataT{0.0f}; - else - return (DataT)0; - } + DataT getValue(size_t) const { return (DataT)0; } + DataT operator()(DataT, DataT) { return (DataT)0; } }; } // namespace detail diff --git a/sycl/test/abi/sycl_symbols_windows.dump b/sycl/test/abi/sycl_symbols_windows.dump index 0bc26c55b8a73..2fc9b38132282 100644 --- a/sycl/test/abi/sycl_symbols_windows.dump +++ b/sycl/test/abi/sycl_symbols_windows.dump @@ -564,8 +564,6 @@ ??0gpu_selector@_V1@sycl@@QEAA@$$QEAV012@@Z ??0gpu_selector@_V1@sycl@@QEAA@AEBV012@@Z ??0gpu_selector@_V1@sycl@@QEAA@XZ -??0half@host_half_impl@detail@_V1@sycl@@QEAA@AEBM@Z -??0half@host_half_impl@detail@_V1@sycl@@QEAA@G@Z ??0handler@_V1@sycl@@AEAA@V?$shared_ptr@Vgraph_impl@detail@experimental@oneapi@ext@_V1@sycl@@@std@@@Z ??0handler@_V1@sycl@@AEAA@V?$shared_ptr@Vqueue_impl@detail@_V1@sycl@@@std@@00_N@Z ??0handler@_V1@sycl@@AEAA@V?$shared_ptr@Vqueue_impl@detail@_V1@sycl@@@std@@_N@Z @@ -747,7 +745,6 @@ ??4fusion_wrapper@experimental@codeplay@ext@_V1@sycl@@QEAAAEAV012345@AEBV012345@@Z ??4gpu_selector@_V1@sycl@@QEAAAEAV012@$$QEAV012@@Z ??4gpu_selector@_V1@sycl@@QEAAAEAV012@AEBV012@@Z -??4half@host_half_impl@detail@_V1@sycl@@QEAAAEAV01234@AEBV01234@@Z ??4image_mem@experimental@oneapi@ext@_V1@sycl@@QEAAAEAV012345@$$QEAV012345@@Z ??4image_mem@experimental@oneapi@ext@_V1@sycl@@QEAAAEAV012345@AEBV012345@@Z ??4image_plain@detail@_V1@sycl@@QEAAAEAV0123@$$QEAV0123@@Z @@ -799,12 +796,6 @@ ??9queue@_V1@sycl@@QEBA_NAEBV012@@Z ??9sampler@_V1@sycl@@QEBA_NAEBV012@@Z ??9stream@_V1@sycl@@QEBA_NAEBV012@@Z -??Bhalf@host_half_impl@detail@_V1@sycl@@QEBAMXZ -??Ehalf@host_half_impl@detail@_V1@sycl@@QEAA?AV01234@H@Z -??Ehalf@host_half_impl@detail@_V1@sycl@@QEAAAEAV01234@XZ -??Fhalf@host_half_impl@detail@_V1@sycl@@QEAA?AV01234@H@Z -??Fhalf@host_half_impl@detail@_V1@sycl@@QEAAAEAV01234@XZ -??Ghalf@host_half_impl@detail@_V1@sycl@@QEAAAEAV01234@XZ ??R?$hash@Vqueue@_V1@sycl@@@std@@QEBA_KAEBVqueue@_V1@sycl@@@Z ??Raccelerator_selector@_V1@sycl@@UEBAHAEBVdevice@12@@Z ??Rcpu_selector@_V1@sycl@@UEBAHAEBVdevice@12@@Z @@ -812,10 +803,6 @@ ??Rfilter_selector@ONEAPI@_V1@sycl@@UEBAHAEBVdevice@23@@Z ??Rfilter_selector@oneapi@ext@_V1@sycl@@UEBAHAEBVdevice@34@@Z ??Rgpu_selector@_V1@sycl@@UEBAHAEBVdevice@12@@Z -??Xhalf@host_half_impl@detail@_V1@sycl@@QEAAAEAV01234@AEBV01234@@Z -??Yhalf@host_half_impl@detail@_V1@sycl@@QEAAAEAV01234@AEBV01234@@Z -??Zhalf@host_half_impl@detail@_V1@sycl@@QEAAAEAV01234@AEBV01234@@Z -??_0half@host_half_impl@detail@_V1@sycl@@QEAAAEAV01234@AEBV01234@@Z ??_7SYCLCategory@detail@_V1@sycl@@6B@ ??_7accelerator_selector@_V1@sycl@@6B@ ??_7cpu_selector@_V1@sycl@@6B@ diff --git a/sycl/tools/abi_check.py b/sycl/tools/abi_check.py index 217910b9f9524..35613b407c557 100644 --- a/sycl/tools/abi_check.py +++ b/sycl/tools/abi_check.py @@ -90,23 +90,6 @@ def parse_readobj_output(output): "?Plugin@?1???$getPlugin@$02@pi@detail@_V1@sycl@@YAAEBVplugin@234@XZ@4PEBV5234@EB", "?Plugin@?1???$getPlugin@$05@pi@detail@_V1@sycl@@YAAEBVplugin@234@XZ@4PEBV5234@EB", ] - # Case 2: - # half_type.hpp: - # class __SYCL_EXPORT half { - # ... - # constexpr half(const half &) = default; - # constexpr half(half &&) = default; - # ... - # }; - # - # For some reason MSVC creates exported symbols for the constexpr versions of those defaulted ctors - # although it never calls them at use point. Instead, those trivially copyable/moveable objects are - # memcpy/memmove'ed. We don't expect these symbols are ever referenced directly so having or not - # having them won't cause ABI issues. - ignore_symbols += [ - "??0half@host_half_impl@detail@_V1@sycl@@QEAA@AEBV01234@@Z", - "??0half@host_half_impl@detail@_V1@sycl@@QEAA@$$QEAV01234@@Z", - ] parsed_symbols = [s for s in parsed_symbols if s not in ignore_symbols] return parsed_symbols