DPC++ Runtime
Runtime libraries for oneAPI DPC++
|
|
Go to the documentation of this file.
19 #if !__has_builtin(__builtin_expect)
20 #define __builtin_expect(a, b) (a)
23 #ifdef __SYCL_DEVICE_ONLY__
26 #define __SYCL_CONSTEXPR_HALF constexpr
27 #elif __cpp_lib_bit_cast || __has_builtin(__builtin_bit_cast)
28 #define __SYCL_CONSTEXPR_HALF constexpr
30 #define __SYCL_CONSTEXPR_HALF
40 class WrapperElementTypeProxy;
49 const uint32_t Bits = sycl::bit_cast<uint32_t>(Val);
52 const uint16_t Sign = (Bits & 0x80000000) >> 16;
54 const uint32_t Frac32 = Bits & 0x7fffff;
56 const uint8_t Exp32 = (Bits & 0x7f800000) >> 23;
57 const int16_t Exp32Diff = Exp32 - 127;
60 uint16_t Exp16 = 0, Frac16 = 0;
67 Exp16 = Exp32Diff + 15;
69 Frac16 = Frac32 >> 13;
72 if (Frac32 >> 12 & 0x01)
76 Frac16 = (Frac32 | (uint32_t(1) << 23)) >> (-Exp32Diff - 1);
96 const uint32_t Sign =
static_cast<uint32_t
>(Val & 0x8000) << 16;
98 const uint8_t Exp16 = (Val & 0x7c00) >> 10;
100 uint16_t Frac16 = Val & 0x3ff;
108 Exp32 =
static_cast<uint32_t
>(Exp16) + 112;
118 }
while ((Frac16 & 0x400) != 0x400);
121 Exp32 = 113 - OffSet;
124 uint32_t Frac32 = Frac16 << 13;
128 Bits |= (Exp32 << 23);
130 const float Result = sycl::bit_cast<float>(Bits);
134 namespace host_half_impl {
140 constexpr
half(
const half &) =
default;
143 half(
const float &rhs);
145 half &operator=(
const half &rhs) =
default;
186 operator float()
const;
188 template <
typename Key>
friend struct std::hash;
191 constexpr
explicit half(uint16_t x) : Buf(x) {}
210 *
this =
operator float() +
static_cast<float>(rhs);
215 *
this =
operator float() -
static_cast<float>(rhs);
220 *
this =
operator float() *
static_cast<float>(rhs);
225 *
this =
operator float() /
static_cast<float>(rhs);
261 template <
typename Key>
friend struct std::hash;
264 constexpr
explicit half_v2(uint16_t x) : Buf(x) {}
266 friend class sycl::ext::intel::esimd::detail::WrapperElementTypeProxy;
274 namespace half_impl {
291 #ifdef __SYCL_DEVICE_ONLY__
316 for (
size_t i = 0; i < NumElements; ++i) {
332 constexpr
half(
const half &) =
default;
337 constexpr
half &operator=(
const half &rhs) =
default;
339 #ifndef __SYCL_DEVICE_ONLY__
344 #endif // __SYCL_DEVICE_ONLY__
390 return half(-other.Data);
394 #define OP(op, op_eq) \
395 __SYCL_CONSTEXPR_HALF friend half operator op(const half lhs, \
401 __SYCL_CONSTEXPR_HALF friend double operator op(const half lhs, \
402 const double rhs) { \
407 __SYCL_CONSTEXPR_HALF friend double operator op(const double lhs, \
413 __SYCL_CONSTEXPR_HALF friend float operator op(const half lhs, \
419 __SYCL_CONSTEXPR_HALF friend float operator op(const float lhs, \
425 __SYCL_CONSTEXPR_HALF friend half operator op(const half lhs, \
431 __SYCL_CONSTEXPR_HALF friend half operator op(const int lhs, \
437 __SYCL_CONSTEXPR_HALF friend half operator op(const half lhs, \
443 __SYCL_CONSTEXPR_HALF friend half operator op(const long lhs, \
449 __SYCL_CONSTEXPR_HALF friend half operator op(const half lhs, \
450 const long long rhs) { \
455 __SYCL_CONSTEXPR_HALF friend half operator op(const long long lhs, \
461 __SYCL_CONSTEXPR_HALF friend half operator op(const half &lhs, \
462 const unsigned int &rhs) { \
467 __SYCL_CONSTEXPR_HALF friend half operator op(const unsigned int &lhs, \
473 __SYCL_CONSTEXPR_HALF friend half operator op(const half &lhs, \
474 const unsigned long &rhs) { \
479 __SYCL_CONSTEXPR_HALF friend half operator op(const unsigned long &lhs, \
485 __SYCL_CONSTEXPR_HALF friend half operator op( \
486 const half &lhs, const unsigned long long &rhs) { \
491 __SYCL_CONSTEXPR_HALF friend half operator op(const unsigned long long &lhs, \
506 __SYCL_CONSTEXPR_HALF friend bool operator op(const half &lhs, \
508 return lhs.Data op rhs.Data; \
510 __SYCL_CONSTEXPR_HALF friend bool operator op(const half &lhs, \
511 const double &rhs) { \
512 return lhs.Data op rhs; \
514 __SYCL_CONSTEXPR_HALF friend bool operator op(const double &lhs, \
516 return lhs op rhs.Data; \
518 __SYCL_CONSTEXPR_HALF friend bool operator op(const half &lhs, \
519 const float &rhs) { \
520 return lhs.Data op rhs; \
522 __SYCL_CONSTEXPR_HALF friend bool operator op(const float &lhs, \
524 return lhs op rhs.Data; \
526 __SYCL_CONSTEXPR_HALF friend bool operator op(const half &lhs, \
528 return lhs.Data op rhs; \
530 __SYCL_CONSTEXPR_HALF friend bool operator op(const int &lhs, \
532 return lhs op rhs.Data; \
534 __SYCL_CONSTEXPR_HALF friend bool operator op(const half &lhs, \
536 return lhs.Data op rhs; \
538 __SYCL_CONSTEXPR_HALF friend bool operator op(const long &lhs, \
540 return lhs op rhs.Data; \
542 __SYCL_CONSTEXPR_HALF friend bool operator op(const half &lhs, \
543 const long long &rhs) { \
544 return lhs.Data op rhs; \
546 __SYCL_CONSTEXPR_HALF friend bool operator op(const long long &lhs, \
548 return lhs op rhs.Data; \
550 __SYCL_CONSTEXPR_HALF friend bool operator op(const half &lhs, \
551 const unsigned int &rhs) { \
552 return lhs.Data op rhs; \
554 __SYCL_CONSTEXPR_HALF friend bool operator op(const unsigned int &lhs, \
556 return lhs op rhs.Data; \
558 __SYCL_CONSTEXPR_HALF friend bool operator op(const half &lhs, \
559 const unsigned long &rhs) { \
560 return lhs.Data op rhs; \
562 __SYCL_CONSTEXPR_HALF friend bool operator op(const unsigned long &lhs, \
564 return lhs op rhs.Data; \
566 __SYCL_CONSTEXPR_HALF friend bool operator op( \
567 const half &lhs, const unsigned long long &rhs) { \
568 return lhs.Data op rhs; \
570 __SYCL_CONSTEXPR_HALF friend bool operator op(const unsigned long long &lhs, \
572 return lhs op rhs.Data; \
585 return static_cast<float>(Data);
588 template <
typename Key>
friend struct std::hash;
590 friend class sycl::ext::intel::esimd::detail::WrapperElementTypeProxy;
606 return static_cast<float>(val);
620 return hash<uint16_t>{}(
reinterpret_cast<const uint16_t &
>(Key));
629 static constexpr
bool is_specialized =
true;
630 static constexpr
bool is_signed =
true;
631 static constexpr
bool is_integer =
false;
632 static constexpr
bool is_exact =
false;
633 static constexpr
bool has_infinity =
true;
634 static constexpr
bool has_quiet_NaN =
true;
635 static constexpr
bool has_signaling_NaN =
true;
636 static constexpr float_denorm_style has_denorm = denorm_present;
637 static constexpr
bool has_denorm_loss =
false;
638 static constexpr
bool tinyness_before =
false;
639 static constexpr
bool traps =
false;
640 static constexpr
int max_exponent10 = 4;
641 static constexpr
int max_exponent = 16;
642 static constexpr
int min_exponent10 = -4;
643 static constexpr
int min_exponent = -13;
644 static constexpr
int radix = 2;
645 static constexpr
int max_digits10 = 5;
646 static constexpr
int digits = 11;
647 static constexpr
bool is_bounded =
true;
648 static constexpr
int digits10 = 3;
649 static constexpr
bool is_modulo =
false;
650 static constexpr
bool is_iec559 =
true;
651 static constexpr float_round_style round_style = round_to_nearest;
654 return 6.103515625e-05f;
666 return 9.765625e-04f;
674 #ifdef __SYCL_DEVICE_ONLY__
675 return __builtin_huge_valf();
678 static_cast<uint16_t
>(0x7C00));
683 return __builtin_nanf(
"");
687 return __builtin_nansf(
"");
698 O << static_cast<float>(rhs);
703 float ValFloat = 0.0f;
709 #undef __SYCL_CONSTEXPR_HALF
710 #undef _CPP14_CONSTEXPR
half_vec< 16 > Vec16StorageT
__SYCL_CONSTEXPR_HALF half & operator-=(const half &rhs)
__SYCL_CONSTEXPR_HALF half_v2 & operator*=(const half_v2 &rhs)
constexpr half(uint16_t x)
static const __SYCL_CONSTEXPR_HALF cl::sycl::half lowest() noexcept
half_vec< 2 > Vec2StorageT
float cast_if_host_half(half_impl::half val)
constexpr half(const detail::host_half_impl::half_v2 &rhs)
__SYCL_CONSTEXPR_HALF half & operator--()
std::istream & operator>>(std::istream &I, cl::sycl::half &rhs)
__SYCL_CONSTEXPR_HALF half_v2 operator++(int)
__SYCL_CONSTEXPR_HALF half_v2 & operator++()
constexpr void initialize_data()
#define __SYCL_CONSTEXPR_HALF
__SYCL_CONSTEXPR_HALF half operator--(int)
cl::sycl::detail::half_impl::half half
static const __SYCL_CONSTEXPR_HALF cl::sycl::half quiet_NaN() noexcept
static const __SYCL_CONSTEXPR_HALF cl::sycl::half denorm_min() noexcept
half_vec< 3 > Vec3StorageT
__SYCL_CONSTEXPR_HALF half_vec()
__SYCL_CONSTEXPR_HALF half & operator+=(const half &rhs)
constexpr half_v2(uint16_t x)
half_vec< 8 > Vec8StorageT
__SYCL_CONSTEXPR_HALF half(const float &rhs)
We provide new interfaces for matrix muliply in this patch:
__SYCL_CONSTEXPR_HALF half_v2 & operator/=(const half_v2 &rhs)
__SYCL_CONSTEXPR_HALF half & operator++()
__SYCL_CONSTEXPR_HALF half_v2 & operator+=(const half_v2 &rhs)
static const __SYCL_CONSTEXPR_HALF cl::sycl::half round_error() noexcept
__SYCL_CONSTEXPR_HALF uint16_t float2Half(const float &Val)
__attribute__((always_inline)) auto invoke_simd(sycl
The invoke_simd free function invokes a SIMD function using all work-items in a sub_group.
__SYCL_CONSTEXPR_HALF float half2Float(const uint16_t &Val)
__SYCL_CONSTEXPR_HALF half_v2 operator--(int)
half_vec< 4 > Vec4StorageT
__SYCL_CONSTEXPR_HALF half & operator/=(const half &rhs)
constexpr half & operator-()
__SYCL_CONSTEXPR_HALF friend half operator-(const half other)
detail::host_half_impl::half_v2 StorageT
__SYCL_CONSTEXPR_HALF half_v2 & operator-=(const half_v2 &rhs)
std::ostream & operator<<(std::ostream &O, cl::sycl::half const &rhs)
#define __builtin_expect(a, b)
static const __SYCL_CONSTEXPR_HALF cl::sycl::half signaling_NaN() noexcept
static const __SYCL_CONSTEXPR_HALF cl::sycl::half epsilon() noexcept
__SYCL_CONSTEXPR_HALF half_v2(const float &rhs)
constexpr half_v2 & operator-()
size_t operator()(cl::sycl::half const &Key) const noexcept
__SYCL_CONSTEXPR_HALF half & operator*=(const half &rhs)
__SYCL_CONSTEXPR_HALF half operator++(int)
__SYCL_CONSTEXPR_HALF half_v2 & operator--()
static constexpr const cl::sycl::half infinity() noexcept
#define __SYCL_INLINE_NAMESPACE(X)