DPC++ Runtime
Runtime libraries for oneAPI DPC++
|
|
Go to the documentation of this file.
20 #if !__has_builtin(__builtin_expect)
21 #define __builtin_expect(a, b) (a)
24 #ifdef __SYCL_DEVICE_ONLY__
27 #define __SYCL_CONSTEXPR_HALF constexpr
28 #elif __cpp_lib_bit_cast || __has_builtin(__builtin_bit_cast)
29 #define __SYCL_CONSTEXPR_HALF constexpr
31 #define __SYCL_CONSTEXPR_HALF
37 namespace ext::intel::esimd::detail {
38 class WrapperElementTypeProxy;
44 const uint32_t Bits = sycl::bit_cast<uint32_t>(Val);
47 const uint16_t Sign = (Bits & 0x80000000) >> 16;
49 const uint32_t Frac32 = Bits & 0x7fffff;
51 const uint8_t Exp32 = (Bits & 0x7f800000) >> 23;
52 const int16_t Exp32Diff = Exp32 - 127;
55 uint16_t Exp16 = 0, Frac16 = 0;
62 Exp16 = Exp32Diff + 15;
64 Frac16 = Frac32 >> 13;
68 uint32_t roundBits = Frac32 & 0x1fff;
69 uint32_t halfway = 0x1000;
70 if (roundBits > halfway)
73 else if (roundBits == halfway)
77 Frac16 = (Frac32 | (uint32_t(1) << 23)) >> (-Exp32Diff - 1);
97 const uint32_t Sign =
static_cast<uint32_t
>(Val & 0x8000) << 16;
99 const uint8_t Exp16 = (Val & 0x7c00) >> 10;
101 uint16_t Frac16 = Val & 0x3ff;
109 Exp32 =
static_cast<uint32_t
>(Exp16) + 112;
119 }
while ((Frac16 & 0x400) != 0x400);
122 Exp32 = 113 - OffSet;
125 uint32_t Frac32 = Frac16 << 13;
129 Bits |= (Exp32 << 23);
131 const float Result = sycl::bit_cast<float>(Bits);
135 namespace host_half_impl {
141 constexpr
half(
const half &) =
default;
150 *
this =
operator float() +
static_cast<float>(rhs);
155 *
this =
operator float() -
static_cast<float>(rhs);
160 *
this =
operator float() *
static_cast<float>(rhs);
165 *
this =
operator float() /
static_cast<float>(rhs);
201 template <
typename Key>
friend struct std::hash;
204 constexpr
explicit half(uint16_t x) : Buf(x) {}
206 friend class sycl::ext::intel::esimd::detail::WrapperElementTypeProxy;
214 namespace half_impl {
231 #ifdef __SYCL_DEVICE_ONLY__
256 for (
size_t i = 0; i < NumElements; ++i) {
269 #ifndef __SYCL_DEVICE_ONLY__
272 class [[__sycl_detail__::__uses_aspects__(aspect::fp16)]]
half {
276 constexpr
half(
const half &) =
default;
283 #ifndef __SYCL_DEVICE_ONLY__
288 #endif // __SYCL_DEVICE_ONLY__
334 return half(-other.Data);
338 #define OP(op, op_eq) \
339 __SYCL_CONSTEXPR_HALF friend half operator op(const half lhs, \
345 __SYCL_CONSTEXPR_HALF friend double operator op(const half lhs, \
346 const double rhs) { \
351 __SYCL_CONSTEXPR_HALF friend double operator op(const double lhs, \
357 __SYCL_CONSTEXPR_HALF friend float operator op(const half lhs, \
363 __SYCL_CONSTEXPR_HALF friend float operator op(const float lhs, \
369 __SYCL_CONSTEXPR_HALF friend half operator op(const half lhs, \
375 __SYCL_CONSTEXPR_HALF friend half operator op(const int lhs, \
381 __SYCL_CONSTEXPR_HALF friend half operator op(const half lhs, \
387 __SYCL_CONSTEXPR_HALF friend half operator op(const long lhs, \
393 __SYCL_CONSTEXPR_HALF friend half operator op(const half lhs, \
394 const long long rhs) { \
399 __SYCL_CONSTEXPR_HALF friend half operator op(const long long lhs, \
405 __SYCL_CONSTEXPR_HALF friend half operator op(const half &lhs, \
406 const unsigned int &rhs) { \
411 __SYCL_CONSTEXPR_HALF friend half operator op(const unsigned int &lhs, \
417 __SYCL_CONSTEXPR_HALF friend half operator op(const half &lhs, \
418 const unsigned long &rhs) { \
423 __SYCL_CONSTEXPR_HALF friend half operator op(const unsigned long &lhs, \
429 __SYCL_CONSTEXPR_HALF friend half operator op( \
430 const half &lhs, const unsigned long long &rhs) { \
435 __SYCL_CONSTEXPR_HALF friend half operator op(const unsigned long long &lhs, \
450 __SYCL_CONSTEXPR_HALF friend bool operator op(const half &lhs, \
452 return lhs.Data op rhs.Data; \
454 __SYCL_CONSTEXPR_HALF friend bool operator op(const half &lhs, \
455 const double &rhs) { \
456 return lhs.Data op rhs; \
458 __SYCL_CONSTEXPR_HALF friend bool operator op(const double &lhs, \
460 return lhs op rhs.Data; \
462 __SYCL_CONSTEXPR_HALF friend bool operator op(const half &lhs, \
463 const float &rhs) { \
464 return lhs.Data op rhs; \
466 __SYCL_CONSTEXPR_HALF friend bool operator op(const float &lhs, \
468 return lhs op rhs.Data; \
470 __SYCL_CONSTEXPR_HALF friend bool operator op(const half &lhs, \
472 return lhs.Data op rhs; \
474 __SYCL_CONSTEXPR_HALF friend bool operator op(const int &lhs, \
476 return lhs op rhs.Data; \
478 __SYCL_CONSTEXPR_HALF friend bool operator op(const half &lhs, \
480 return lhs.Data op rhs; \
482 __SYCL_CONSTEXPR_HALF friend bool operator op(const long &lhs, \
484 return lhs op rhs.Data; \
486 __SYCL_CONSTEXPR_HALF friend bool operator op(const half &lhs, \
487 const long long &rhs) { \
488 return lhs.Data op rhs; \
490 __SYCL_CONSTEXPR_HALF friend bool operator op(const long long &lhs, \
492 return lhs op rhs.Data; \
494 __SYCL_CONSTEXPR_HALF friend bool operator op(const half &lhs, \
495 const unsigned int &rhs) { \
496 return lhs.Data op rhs; \
498 __SYCL_CONSTEXPR_HALF friend bool operator op(const unsigned int &lhs, \
500 return lhs op rhs.Data; \
502 __SYCL_CONSTEXPR_HALF friend bool operator op(const half &lhs, \
503 const unsigned long &rhs) { \
504 return lhs.Data op rhs; \
506 __SYCL_CONSTEXPR_HALF friend bool operator op(const unsigned long &lhs, \
508 return lhs op rhs.Data; \
510 __SYCL_CONSTEXPR_HALF friend bool operator op( \
511 const half &lhs, const unsigned long long &rhs) { \
512 return lhs.Data op rhs; \
514 __SYCL_CONSTEXPR_HALF friend bool operator op(const unsigned long long &lhs, \
516 return lhs op rhs.Data; \
529 return static_cast<float>(Data);
535 O << static_cast<float>(rhs);
540 float ValFloat = 0.0f;
546 template <
typename Key>
friend struct std::hash;
548 friend class sycl::ext::intel::esimd::detail::WrapperElementTypeProxy;
564 return static_cast<float>(val);
578 return hash<uint16_t>{}(
reinterpret_cast<const uint16_t &
>(Key));
587 static constexpr
bool is_specialized =
true;
588 static constexpr
bool is_signed =
true;
589 static constexpr
bool is_integer =
false;
590 static constexpr
bool is_exact =
false;
591 static constexpr
bool has_infinity =
true;
592 static constexpr
bool has_quiet_NaN =
true;
593 static constexpr
bool has_signaling_NaN =
true;
594 static constexpr float_denorm_style has_denorm = denorm_present;
595 static constexpr
bool has_denorm_loss =
false;
596 static constexpr
bool tinyness_before =
false;
597 static constexpr
bool traps =
false;
598 static constexpr
int max_exponent10 = 4;
599 static constexpr
int max_exponent = 16;
600 static constexpr
int min_exponent10 = -4;
601 static constexpr
int min_exponent = -13;
602 static constexpr
int radix = 2;
603 static constexpr
int max_digits10 = 5;
604 static constexpr
int digits = 11;
605 static constexpr
bool is_bounded =
true;
606 static constexpr
int digits10 = 3;
607 static constexpr
bool is_modulo =
false;
608 static constexpr
bool is_iec559 =
true;
609 static constexpr float_round_style round_style = round_to_nearest;
612 return 6.103515625e-05f;
624 return 9.765625e-04f;
632 #ifdef __SYCL_DEVICE_ONLY__
633 return __builtin_huge_valf();
640 return __builtin_nanf(
"");
644 return __builtin_nansf(
"");
654 #undef __SYCL_CONSTEXPR_HALF
655 #undef _CPP14_CONSTEXPR
__SYCL_CONSTEXPR_HALF half operator++(int)
static const __SYCL_CONSTEXPR_HALF sycl::half quiet_NaN() noexcept
__SYCL_CONSTEXPR_HALF half operator--(int)
constexpr void initialize_data()
__SYCL_CONSTEXPR_HALF float half2Float(const uint16_t &Val)
__SYCL_CONSTEXPR_HALF half & operator++()
half_vec< 16 > Vec16StorageT
size_t operator()(sycl::half const &Key) const noexcept
__SYCL_CONSTEXPR_HALF half_vec()
#define __SYCL_INLINE_VER_NAMESPACE(X)
__SYCL_CONSTEXPR_HALF half & operator-=(const half &rhs)
---— Error handling, matching OpenCL plugin semantics.
__SYCL_CONSTEXPR_HALF half operator--(int)
#define __SYCL_CONSTEXPR_HALF
__SYCL_CONSTEXPR_HALF friend half operator-(const half other)
__SYCL_CONSTEXPR_HALF half & operator--()
__SYCL_CONSTEXPR_HALF half & operator++()
__SYCL_CONSTEXPR_HALF half & operator-=(const half &rhs)
__SYCL_CONSTEXPR_HALF half & operator+=(const half &rhs)
half_vec< 2 > Vec2StorageT
annotated_arg & operator=(annotated_arg &)=default
__SYCL_CONSTEXPR_HALF half(const float &rhs)
static const __SYCL_CONSTEXPR_HALF sycl::half round_error() noexcept
__SYCL_CONSTEXPR_HALF half & operator--()
static const __SYCL_CONSTEXPR_HALF sycl::half lowest() noexcept
sycl::detail::half_impl::half half
__SYCL_CONSTEXPR_HALF half & operator*=(const half &rhs)
half_vec< 3 > Vec3StorageT
float cast_if_host_half(half_impl::half val)
half_vec< 8 > Vec8StorageT
constexpr half(uint16_t x)
static const __SYCL_CONSTEXPR_HALF sycl::half signaling_NaN() noexcept
static const __SYCL_CONSTEXPR_HALF sycl::half denorm_min() noexcept
friend std::ostream & operator<<(std::ostream &O, sycl::half const &rhs)
__SYCL_CONSTEXPR_HALF half operator++(int)
__SYCL_CONSTEXPR_HALF half & operator+=(const half &rhs)
half_vec< 4 > Vec4StorageT
__SYCL_CONSTEXPR_HALF half & operator/=(const half &rhs)
static constexpr const sycl::half infinity() noexcept
annotated_ptr & operator--() noexcept
static const __SYCL_CONSTEXPR_HALF sycl::half epsilon() noexcept
detail::host_half_impl::half StorageT
annotated_ptr & operator++() noexcept
#define __builtin_expect(a, b)
__SYCL_CONSTEXPR_HALF uint16_t float2Half(const float &Val)
__SYCL_CONSTEXPR_HALF half(const float &rhs)
__SYCL_CONSTEXPR_HALF half & operator*=(const half &rhs)
constexpr half(const detail::host_half_impl::half &rhs)
__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 half & operator/=(const half &rhs)
constexpr half & operator-()
friend std::istream & operator>>(std::istream &I, sycl::half &rhs)