Skip to content

Commit

Permalink
[SYCL][ABI-Break] Fold host_half_impl::half into half_impl::half (#13597
Browse files Browse the repository at this point in the history 
)

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 <steffen.larsen@intel.com>
  • Loading branch information
@steffenlarsen
steffenlarsen authored Jul 16, 2024
1 parent cb3f5cd commit d2f6fb3
Show file tree
Hide file tree
Showing 6 changed files with 95 additions and 161 deletions.
11 changes: 1 addition & 10 deletions sycl/include/sycl/ext/intel/esimd/detail/half_type_traits.hpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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__

Expand Down Expand Up @@ -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) {
Expand Down
197 changes: 90 additions & 107 deletions sycl/include/sycl/half_type.hpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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
Expand Down Expand Up @@ -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;
Expand All @@ -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
Expand All @@ -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;
Expand All @@ -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++() {
Expand All @@ -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) \
Expand Down Expand Up @@ -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(!=)
Expand All @@ -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,
Expand All @@ -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
Expand Down Expand Up @@ -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
}

Expand Down
4 changes: 2 additions & 2 deletions sycl/include/sycl/known_identity.hpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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
};

Expand Down Expand Up @@ -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
};

Expand Down
Loading

0 comments on commit d2f6fb3

Please sign in to comment.