16 #ifdef __SYCL_DEVICE_ONLY__
24 #include <string_view>
25 #include <type_traits>
27 #if !defined(__has_builtin) || !__has_builtin(__builtin_expect)
28 #define __builtin_expect(a, b) (a)
31 #ifdef __SYCL_DEVICE_ONLY__
34 #define __SYCL_CONSTEXPR_HALF constexpr
35 #elif __cpp_lib_bit_cast || \
36 (defined(__has_builtin) && __has_builtin(__builtin_bit_cast))
37 #define __SYCL_CONSTEXPR_HALF constexpr
39 #define __SYCL_CONSTEXPR_HALF
43 inline namespace _V1 {
44 namespace detail::half_impl {
49 namespace ext::intel::esimd::detail {
50 class WrapperElementTypeProxy;
56 const uint32_t Bits = sycl::bit_cast<uint32_t>(Val);
59 const uint16_t Sign = (Bits & 0x80000000) >> 16;
61 const uint32_t Frac32 = Bits & 0x7fffff;
63 const uint8_t Exp32 = (Bits & 0x7f800000) >> 23;
64 const int16_t Exp32Diff = Exp32 - 127;
67 uint16_t Exp16 = 0, Frac16 = 0;
74 Exp16 = Exp32Diff + 15;
76 Frac16 = Frac32 >> 13;
80 uint32_t roundBits = Frac32 & 0x1fff;
81 uint32_t halfway = 0x1000;
82 if (roundBits > halfway)
85 else if (roundBits == halfway)
89 Frac16 = (Frac32 | (uint32_t(1) << 23)) >> (-Exp32Diff - 1);
109 const uint32_t Sign =
static_cast<uint32_t
>(Val & 0x8000) << 16;
111 const uint8_t Exp16 = (Val & 0x7c00) >> 10;
113 uint16_t Frac16 = Val & 0x3ff;
121 Exp32 =
static_cast<uint32_t
>(Exp16) + 112;
131 }
while ((Frac16 & 0x400) != 0x400);
134 Exp32 = 113 - OffSet;
137 uint32_t Frac32 = Frac16 << 13;
141 Bits |= (Exp32 << 23);
143 const float Result = sycl::bit_cast<float>(Bits);
147 namespace host_half_impl {
162 *
this =
operator float() +
static_cast<float>(rhs);
167 *
this =
operator float() -
static_cast<float>(rhs);
172 *
this =
operator float() *
static_cast<float>(rhs);
177 *
this =
operator float() /
static_cast<float>(rhs);
213 template <
typename Key>
friend struct std::hash;
216 constexpr
explicit half(uint16_t x) : Buf(
x) {}
218 friend class sycl::ext::intel::esimd::detail::WrapperElementTypeProxy;
226 namespace half_impl {
243 #ifdef __SYCL_DEVICE_ONLY__
262 #ifdef __INTEL_PREVIEW_BREAKING_CHANGES
274 template <
typename... Ts,
275 typename = std::enable_if_t<(
sizeof...(Ts) == NumElements) &&
276 (std::is_same_v<half, Ts> && ...)>>
280 for (
size_t i = 0; i < NumElements; ++i) {
294 #ifndef __SYCL_DEVICE_ONLY__
297 class [[__sycl_detail__::__uses_aspects__(aspect::fp16)]]
half {
308 #ifndef __SYCL_DEVICE_ONLY__
359 return half(-other.Data);
363 #define OP(op, op_eq) \
364 __SYCL_CONSTEXPR_HALF friend half operator op(const half lhs, \
370 __SYCL_CONSTEXPR_HALF friend double operator op(const half lhs, \
371 const double rhs) { \
376 __SYCL_CONSTEXPR_HALF friend double operator op(const double lhs, \
382 __SYCL_CONSTEXPR_HALF friend float operator op(const half lhs, \
388 __SYCL_CONSTEXPR_HALF friend float operator op(const float lhs, \
394 __SYCL_CONSTEXPR_HALF friend half operator op(const half lhs, \
400 __SYCL_CONSTEXPR_HALF friend half operator op(const int lhs, \
406 __SYCL_CONSTEXPR_HALF friend half operator op(const half lhs, \
412 __SYCL_CONSTEXPR_HALF friend half operator op(const long lhs, \
418 __SYCL_CONSTEXPR_HALF friend half operator op(const half lhs, \
419 const long long rhs) { \
424 __SYCL_CONSTEXPR_HALF friend half operator op(const long long lhs, \
430 __SYCL_CONSTEXPR_HALF friend half operator op(const half &lhs, \
431 const unsigned int &rhs) { \
436 __SYCL_CONSTEXPR_HALF friend half operator op(const unsigned int &lhs, \
442 __SYCL_CONSTEXPR_HALF friend half operator op(const half &lhs, \
443 const unsigned long &rhs) { \
448 __SYCL_CONSTEXPR_HALF friend half operator op(const unsigned long &lhs, \
454 __SYCL_CONSTEXPR_HALF friend half operator op( \
455 const half &lhs, const unsigned long long &rhs) { \
460 __SYCL_CONSTEXPR_HALF friend half operator op(const unsigned long long &lhs, \
475 __SYCL_CONSTEXPR_HALF friend bool operator op(const half &lhs, \
477 return lhs.Data op rhs.Data; \
479 __SYCL_CONSTEXPR_HALF friend bool operator op(const half &lhs, \
480 const double &rhs) { \
481 return lhs.Data op rhs; \
483 __SYCL_CONSTEXPR_HALF friend bool operator op(const double &lhs, \
485 return lhs op rhs.Data; \
487 __SYCL_CONSTEXPR_HALF friend bool operator op(const half &lhs, \
488 const float &rhs) { \
489 return lhs.Data op rhs; \
491 __SYCL_CONSTEXPR_HALF friend bool operator op(const float &lhs, \
493 return lhs op rhs.Data; \
495 __SYCL_CONSTEXPR_HALF friend bool operator op(const half &lhs, \
497 return lhs.Data op rhs; \
499 __SYCL_CONSTEXPR_HALF friend bool operator op(const int &lhs, \
501 return lhs op rhs.Data; \
503 __SYCL_CONSTEXPR_HALF friend bool operator op(const half &lhs, \
505 return lhs.Data op rhs; \
507 __SYCL_CONSTEXPR_HALF friend bool operator op(const long &lhs, \
509 return lhs op rhs.Data; \
511 __SYCL_CONSTEXPR_HALF friend bool operator op(const half &lhs, \
512 const long long &rhs) { \
513 return lhs.Data op rhs; \
515 __SYCL_CONSTEXPR_HALF friend bool operator op(const long long &lhs, \
517 return lhs op rhs.Data; \
519 __SYCL_CONSTEXPR_HALF friend bool operator op(const half &lhs, \
520 const unsigned int &rhs) { \
521 return lhs.Data op rhs; \
523 __SYCL_CONSTEXPR_HALF friend bool operator op(const unsigned int &lhs, \
525 return lhs op rhs.Data; \
527 __SYCL_CONSTEXPR_HALF friend bool operator op(const half &lhs, \
528 const unsigned long &rhs) { \
529 return lhs.Data op rhs; \
531 __SYCL_CONSTEXPR_HALF friend bool operator op(const unsigned long &lhs, \
533 return lhs op rhs.Data; \
535 __SYCL_CONSTEXPR_HALF friend bool operator op( \
536 const half &lhs, const unsigned long long &rhs) { \
537 return lhs.Data op rhs; \
539 __SYCL_CONSTEXPR_HALF friend bool operator op(const unsigned long long &lhs, \
541 return lhs op rhs.Data; \
554 return static_cast<float>(Data);
560 O << static_cast<float>(rhs);
565 float ValFloat = 0.0f;
571 template <
typename Key>
friend struct std::hash;
573 friend class sycl::ext::intel::esimd::detail::WrapperElementTypeProxy;
589 return static_cast<float>(val);
603 return hash<uint16_t>{}(
reinterpret_cast<const uint16_t &
>(Key));
612 static constexpr
bool is_specialized =
true;
613 static constexpr
bool is_signed =
true;
614 static constexpr
bool is_integer =
false;
615 static constexpr
bool is_exact =
false;
616 static constexpr
bool has_infinity =
true;
617 static constexpr
bool has_quiet_NaN =
true;
618 static constexpr
bool has_signaling_NaN =
true;
619 static constexpr float_denorm_style has_denorm = denorm_present;
620 static constexpr
bool has_denorm_loss =
false;
621 static constexpr
bool tinyness_before =
false;
622 static constexpr
bool traps =
false;
623 static constexpr
int max_exponent10 = 4;
624 static constexpr
int max_exponent = 16;
625 static constexpr
int min_exponent10 = -4;
626 static constexpr
int min_exponent = -13;
627 static constexpr
int radix = 2;
628 static constexpr
int max_digits10 = 5;
629 static constexpr
int digits = 11;
630 static constexpr
bool is_bounded =
true;
631 static constexpr
int digits10 = 3;
632 static constexpr
bool is_modulo =
false;
633 static constexpr
bool is_iec559 =
true;
634 static constexpr float_round_style round_style = round_to_nearest;
637 return 6.103515625e-05f;
649 return 9.765625e-04f;
657 #ifdef __SYCL_DEVICE_ONLY__
658 return __builtin_huge_valf();
665 return __builtin_nanf(
"");
669 return __builtin_nansf(
"");
679 #undef __SYCL_CONSTEXPR_HALF
680 #undef _CPP14_CONSTEXPR
friend std::istream & operator>>(std::istream &I, sycl::half &rhs)
__SYCL_CONSTEXPR_HALF half & operator/=(const half &rhs)
__SYCL_CONSTEXPR_HALF half & operator+=(const half &rhs)
__SYCL_CONSTEXPR_HALF half & operator-=(const half &rhs)
__SYCL_CONSTEXPR_HALF half operator--(int)
__SYCL_CONSTEXPR_HALF half & operator*=(const half &rhs)
__SYCL_CONSTEXPR_HALF friend half operator-(const half other)
constexpr half(half &&)=default
constexpr half(const detail::host_half_impl::half &rhs)
constexpr half(const half &)=default
constexpr half & operator=(const half &rhs)=default
__SYCL_CONSTEXPR_HALF half & operator--()
__SYCL_CONSTEXPR_HALF half(const float &rhs)
friend std::ostream & operator<<(std::ostream &O, sycl::half const &rhs)
__SYCL_CONSTEXPR_HALF half & operator++()
__SYCL_CONSTEXPR_HALF half operator++(int)
constexpr half(half &&)=default
constexpr half & operator=(const half &rhs)=default
__SYCL_CONSTEXPR_HALF half & operator-=(const half &rhs)
constexpr half & operator-()
__SYCL_CONSTEXPR_HALF half & operator+=(const half &rhs)
constexpr half(uint16_t x)
__SYCL_CONSTEXPR_HALF half(const float &rhs)
__SYCL_CONSTEXPR_HALF half & operator++()
__SYCL_CONSTEXPR_HALF half & operator/=(const half &rhs)
__SYCL_CONSTEXPR_HALF half operator--(int)
__SYCL_CONSTEXPR_HALF half & operator--()
__SYCL_CONSTEXPR_HALF half operator++(int)
constexpr half(const half &)=default
__SYCL_CONSTEXPR_HALF half & operator*=(const half &rhs)
#define __builtin_expect(a, b)
#define __SYCL_CONSTEXPR_HALF
half_vec< 16 > Vec16StorageT
half_vec< 4 > Vec4StorageT
half_vec< 3 > Vec3StorageT
detail::host_half_impl::half StorageT
half_vec< 2 > Vec2StorageT
half_vec< 8 > Vec8StorageT
T cast_if_host_half(T val)
__SYCL_CONSTEXPR_HALF uint16_t float2Half(const float &Val)
__SYCL_CONSTEXPR_HALF float half2Float(const uint16_t &Val)
__attribute__((destructor(110))) static void syclUnload()
sycl::detail::half_impl::half half
_Abi const simd< _Tp, _Abi > & noexcept
size_t operator()(sycl::half const &Key) const noexcept
static constexpr const sycl::half infinity() noexcept
static __SYCL_CONSTEXPR_HALF const sycl::half round_error() noexcept
static __SYCL_CONSTEXPR_HALF const sycl::half quiet_NaN() noexcept
static __SYCL_CONSTEXPR_HALF const sycl::half lowest() noexcept
static __SYCL_CONSTEXPR_HALF const sycl::half denorm_min() noexcept
static __SYCL_CONSTEXPR_HALF const sycl::half epsilon() noexcept
static __SYCL_CONSTEXPR_HALF const sycl::half signaling_NaN() noexcept
__SYCL_CONSTEXPR_HALF half_vec(const Ts &...hs)
constexpr void initialize_data()
__SYCL_CONSTEXPR_HALF half_vec()