[SYCL][ABI-Break] Fold host_half_impl::half into half_impl::half (intel#13597)

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 <[email protected]>

Author: steffenlarsen

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 <int N> struct vector_conversion_traits<sycl::half, N> {
 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) {

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<float>(rhs);
-    return *this;
-  }
-
-  __SYCL_CONSTEXPR_HALF half &operator-=(const half &rhs) {
-    *this = operator float() - static_cast<float>(rhs);
-    return *this;
-  }
-
-  __SYCL_CONSTEXPR_HALF half &operator*=(const half &rhs) {
-    *this = operator float() * static_cast<float>(rhs);
-    return *this;
-  }
-
-  __SYCL_CONSTEXPR_HALF half &operator/=(const half &rhs) {
-    *this = operator float() / static_cast<float>(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 <typename Key> 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<VecElemT, 16>;
 
 #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<float>(rhs);
+    return *this;
+  }
+
+  __SYCL_CONSTEXPR_HALF half &operator-=(const half &rhs) {
+    *this = operator float() - static_cast<float>(rhs);
+    return *this;
+  }
+
+  __SYCL_CONSTEXPR_HALF half &operator*=(const half &rhs) {
+    *this = operator float() * static_cast<float>(rhs);
+    return *this;
+  }
+
+  __SYCL_CONSTEXPR_HALF half &operator/=(const half &rhs) {
+    *this = operator float() / static_cast<float>(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<float>(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<sycl::half> {
 #ifdef __SYCL_DEVICE_ONLY__
     return __builtin_huge_valf();
 #else
-    return sycl::detail::host_half_impl::half(static_cast<uint16_t>(0x7C00));
+    return sycl::detail::half_impl::CreateHostHalfRaw(
+        static_cast<uint16_t>(0x7C00));
 #endif
   }
 

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<uint16_t>(0));
+      sycl::detail::half_impl::CreateHostHalfRaw(static_cast<uint16_t>(0));
 #endif
 };
 
@@ -227,7 +227,7 @@ struct known_identity_impl<
 #ifdef __SYCL_DEVICE_ONLY__
       1;
 #else
-      sycl::detail::host_half_impl::half(static_cast<uint16_t>(0x3C00));
+      sycl::detail::half_impl::CreateHostHalfRaw(static_cast<uint16_t>(0x3C00));
 #endif
 };