From dea5b7f2d3bdc51750ba08aef6f6ecf906b864aa Mon Sep 17 00:00:00 2001 From: "Larsen, Steffen" Date: Tue, 30 Apr 2024 00:02:46 -0700 Subject: [PATCH 1/6] [SYCL] Change internal host vec representation This commit changes the internal representation of sycl::vec to use sycl::half directly instead of the wrapper implementation used inside them. This avoids strict alias violation in the operator[] for the host implementation. Signed-off-by: Larsen, Steffen --- sycl/include/sycl/half_type.hpp | 36 +++++++++++++++++++-------------- sycl/include/sycl/vector.hpp | 12 ++++++----- 2 files changed, 28 insertions(+), 20 deletions(-) diff --git a/sycl/include/sycl/half_type.hpp b/sycl/include/sycl/half_type.hpp index f3a04684c5f58..951146f2cdfbb 100644 --- a/sycl/include/sycl/half_type.hpp +++ b/sycl/include/sycl/half_type.hpp @@ -228,9 +228,9 @@ class half; // Several aliases are defined below: // - StorageT: actual representation of half data type. It is used by scalar -// half values and by 'sycl::vec' class. 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 +// 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 // // - BIsRepresentationT: data type which is used by built-in functions. It is // distinguished from StorageT, because on host, we can still operate on the @@ -238,32 +238,38 @@ class half; // type (too many changes required for BIs implementation without any // foreseeable profits) // -// - VecNStorageT - representation of N-element vector of halfs. Follows the -// same logic as StorageT +// - VecElemT: representation of each element in the vector. On device it is +// the same as StorageT to carry a native vector representation, while on +// host it stores the sycl::half implementation directly. +// +// - VecNStorageT: representation of N-element vector of halfs. Follows the +// same logic as VecElemT. #ifdef __SYCL_DEVICE_ONLY__ using StorageT = _Float16; using BIsRepresentationT = _Float16; +using VecElemT = _Float16; -using Vec2StorageT = StorageT __attribute__((ext_vector_type(2))); -using Vec3StorageT = StorageT __attribute__((ext_vector_type(3))); -using Vec4StorageT = StorageT __attribute__((ext_vector_type(4))); -using Vec8StorageT = StorageT __attribute__((ext_vector_type(8))); -using Vec16StorageT = StorageT __attribute__((ext_vector_type(16))); +using Vec2StorageT = VecElemT __attribute__((ext_vector_type(2))); +using Vec3StorageT = VecElemT __attribute__((ext_vector_type(3))); +using Vec4StorageT = VecElemT __attribute__((ext_vector_type(4))); +using Vec8StorageT = VecElemT __attribute__((ext_vector_type(8))); +using Vec16StorageT = VecElemT __attribute__((ext_vector_type(16))); #else // SYCL_DEVICE_ONLY using StorageT = detail::host_half_impl::half; // No need to extract underlying data type for built-in functions operating on // host using BIsRepresentationT = half; +using VecElemT = half; // On the host side we cannot use OpenCL cl_half# types as an underlying type // for vec because they are actually defined as an integer type under the // hood. As a result half values will be converted to the integer and passed // as a kernel argument which is expected to be floating point number. -using Vec2StorageT = std::array; -using Vec3StorageT = std::array; -using Vec4StorageT = std::array; -using Vec8StorageT = std::array; -using Vec16StorageT = std::array; +using Vec2StorageT = std::array; +using Vec3StorageT = std::array; +using Vec4StorageT = std::array; +using Vec8StorageT = std::array; +using Vec16StorageT = std::array; #endif // SYCL_DEVICE_ONLY #ifndef __SYCL_DEVICE_ONLY__ diff --git a/sycl/include/sycl/vector.hpp b/sycl/include/sycl/vector.hpp index f1bf7fcfcc24d..7c4c509062b37 100644 --- a/sycl/include/sycl/vector.hpp +++ b/sycl/include/sycl/vector.hpp @@ -300,9 +300,9 @@ struct VecStorage< // Single element half template <> struct VecStorage { - using DataType = sycl::detail::half_impl::StorageT; + using DataType = sycl::detail::half_impl::VecElemT; #ifdef __SYCL_DEVICE_ONLY__ - using VectorDataType = sycl::detail::half_impl::StorageT; + using VectorDataType = sycl::detail::half_impl::VecElemT; #endif // __SYCL_DEVICE_ONLY__ }; @@ -365,10 +365,12 @@ template class vec { // in the class, so vec should be equal to float16 in memory. using DataType = typename detail::VecStorage::DataType; +#ifdef __SYCL_DEVICE_ONLY__ + static constexpr bool IsHostHalf = false; +#else static constexpr bool IsHostHalf = - std::is_same_v && - std::is_same_v; + std::is_same_v; +#endif static constexpr bool IsBfloat16 = std::is_same_v; From 6d283a73668c670115b4f5ed16f5ec07b18d8acc Mon Sep 17 00:00:00 2001 From: "Larsen, Steffen" Date: Mon, 29 Apr 2024 08:55:36 -0700 Subject: [PATCH 2/6] [SYCL][ABI-Break] Fold host_half_impl::half into half_impl::half This commit folds the implementation of host_half_impl::half into half_impl::half and making the vector element representation the same as the half representation. This allows us to avoid strict alias violation for half vectors in their operator[] implementations. Note that this is marked as an ABI break as it removes symbols on Windows, despite these symbols never being in the library. Signed-off-by: Larsen, Steffen --- .../intel/esimd/detail/half_type_traits.hpp | 11 +- sycl/include/sycl/half_type.hpp | 199 ++++++++---------- sycl/include/sycl/known_identity.hpp | 4 +- sycl/test/abi/sycl_symbols_windows.dump | 13 -- sycl/tools/abi_check.py | 17 -- 5 files changed, 94 insertions(+), 150 deletions(-) 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 951146f2cdfbb..10679ab9dd004 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 @@ -254,8 +175,8 @@ using Vec3StorageT = VecElemT __attribute__((ext_vector_type(3))); using Vec4StorageT = VecElemT __attribute__((ext_vector_type(4))); using Vec8StorageT = VecElemT __attribute__((ext_vector_type(8))); using Vec16StorageT = VecElemT __attribute__((ext_vector_type(16))); -#else // SYCL_DEVICE_ONLY -using StorageT = detail::host_half_impl::half; +#else // SYCL_DEVICE_ONLY +using StorageT = uint16_t; // No need to extract underlying data type for built-in functions operating on // host using BIsRepresentationT = half; @@ -272,6 +193,12 @@ using Vec8StorageT = std::array; 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 @@ -282,18 +209,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; @@ -313,6 +238,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++() { @@ -336,9 +282,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) \ @@ -455,71 +409,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(!=) @@ -531,9 +485,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, @@ -554,8 +512,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__ + constexpr StorageT getFPRep() const { return Data; } +#else + constexpr 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 @@ -638,7 +620,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/test/abi/sycl_symbols_windows.dump b/sycl/test/abi/sycl_symbols_windows.dump index 69473362c1985..c42a5cebf2cca 100644 --- a/sycl/test/abi/sycl_symbols_windows.dump +++ b/sycl/test/abi/sycl_symbols_windows.dump @@ -571,8 +571,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 @@ -782,7 +780,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 ??4host_selector@_V1@sycl@@QEAAAEAV012@$$QEAV012@@Z ??4host_selector@_V1@sycl@@QEAAAEAV012@AEBV012@@Z ??4image_mem@experimental@oneapi@ext@_V1@sycl@@QEAAAEAV012345@$$QEAV012345@@Z @@ -834,12 +831,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 @@ -848,10 +839,6 @@ ??Rfilter_selector@oneapi@ext@_V1@sycl@@UEBAHAEBVdevice@34@@Z ??Rgpu_selector@_V1@sycl@@UEBAHAEBVdevice@12@@Z ??Rhost_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@ ??_7SYCLMemObjT@detail@_V1@sycl@@6B@ ??_7accelerator_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 From dae8afdcda0bbcd3af43b6bd3df0c2df5da9f3f3 Mon Sep 17 00:00:00 2001 From: "Larsen, Steffen" Date: Tue, 30 Apr 2024 04:01:31 -0700 Subject: [PATCH 3/6] Amend new preview vector file Signed-off-by: Larsen, Steffen --- sycl/include/sycl/vector_preview.hpp | 8 +++++--- 1 file changed, 5 insertions(+), 3 deletions(-) diff --git a/sycl/include/sycl/vector_preview.hpp b/sycl/include/sycl/vector_preview.hpp index f1bf7fcfcc24d..f0e4888207625 100644 --- a/sycl/include/sycl/vector_preview.hpp +++ b/sycl/include/sycl/vector_preview.hpp @@ -365,10 +365,12 @@ template class vec { // in the class, so vec should be equal to float16 in memory. using DataType = typename detail::VecStorage::DataType; +#ifdef __SYCL_DEVICE_ONLY__ + static constexpr bool IsHostHalf = false; +#else static constexpr bool IsHostHalf = - std::is_same_v && - std::is_same_v; + std::is_same_v; +#endif static constexpr bool IsBfloat16 = std::is_same_v; From 9bf5634ad63dc2735f5c2fb2098c1a2446d9f1e5 Mon Sep 17 00:00:00 2001 From: "Larsen, Steffen" Date: Tue, 30 Apr 2024 04:45:56 -0700 Subject: [PATCH 4/6] Make getFPRep conditionally constexpr Signed-off-by: Larsen, Steffen --- sycl/include/sycl/half_type.hpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/sycl/include/sycl/half_type.hpp b/sycl/include/sycl/half_type.hpp index 10679ab9dd004..5f56e3ecde7aa 100644 --- a/sycl/include/sycl/half_type.hpp +++ b/sycl/include/sycl/half_type.hpp @@ -516,9 +516,9 @@ class [[__sycl_detail__::__uses_aspects__(aspect::fp16)]] half { // 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__ - constexpr StorageT getFPRep() const { return Data; } + __SYCL_CONSTEXPR_HALF StorageT getFPRep() const { return Data; } #else - constexpr float getFPRep() const { return operator float(); } + __SYCL_CONSTEXPR_HALF float getFPRep() const { return operator float(); } #endif #ifndef __SYCL_DEVICE_ONLY__ From 0ef5a633755f39a57d274f0117cb546a524279d4 Mon Sep 17 00:00:00 2001 From: "Larsen, Steffen" Date: Wed, 1 May 2024 00:26:57 -0700 Subject: [PATCH 5/6] Amend 1D half vec storage type Signed-off-by: Larsen, Steffen --- sycl/include/sycl/vector_preview.hpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/sycl/include/sycl/vector_preview.hpp b/sycl/include/sycl/vector_preview.hpp index f0e4888207625..7c4c509062b37 100644 --- a/sycl/include/sycl/vector_preview.hpp +++ b/sycl/include/sycl/vector_preview.hpp @@ -300,9 +300,9 @@ struct VecStorage< // Single element half template <> struct VecStorage { - using DataType = sycl::detail::half_impl::StorageT; + using DataType = sycl::detail::half_impl::VecElemT; #ifdef __SYCL_DEVICE_ONLY__ - using VectorDataType = sycl::detail::half_impl::StorageT; + using VectorDataType = sycl::detail::half_impl::VecElemT; #endif // __SYCL_DEVICE_ONLY__ }; From e3dbc2f24b210e81c04c265349560cbed712d97e Mon Sep 17 00:00:00 2001 From: "Larsen, Steffen" Date: Thu, 27 Jun 2024 09:37:21 -0700 Subject: [PATCH 6/6] Fix formatting Signed-off-by: Larsen, Steffen --- sycl/include/sycl/vector_preview.hpp | 8 ++------ 1 file changed, 2 insertions(+), 6 deletions(-) diff --git a/sycl/include/sycl/vector_preview.hpp b/sycl/include/sycl/vector_preview.hpp index 253f368150f8c..6646f01a00454 100644 --- a/sycl/include/sycl/vector_preview.hpp +++ b/sycl/include/sycl/vector_preview.hpp @@ -93,12 +93,8 @@ class SwizzleOp; template class GetOp { public: using DataT = T; - DataT getValue(size_t) const { - return (DataT)0; - } - DataT operator()(DataT, DataT) { - return (DataT)0; - } + DataT getValue(size_t) const { return (DataT)0; } + DataT operator()(DataT, DataT) { return (DataT)0; } }; } // namespace detail