DPC++ Runtime
Runtime libraries for oneAPI DPC++
half_type.hpp
Go to the documentation of this file.
1 //==-------------- half_type.hpp --- SYCL half type ------------------------==//
2 //
3 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4 // See https://llvm.org/LICENSE.txt for license information.
5 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6 //
7 //===----------------------------------------------------------------------===//
8 
9 #pragma once
10 
11 #include <sycl/bit_cast.hpp> // for bit_cast
12 #include <sycl/detail/export.hpp> // for __SYCL_EXPORT
13 #include <sycl/detail/iostream_proxy.hpp> // for istream, ostream
14 #include <sycl/detail/vector_traits.hpp> // for vector_alignment
15 
16 #ifdef __SYCL_DEVICE_ONLY__
17 #include <sycl/aspects.hpp>
18 #endif
19 
20 #include <cstddef> // for size_t
21 #include <cstdint> // for uint16_t, uint32_t, uint8_t
22 #include <functional> // for hash
23 #include <limits> // for float_denorm_style, float_r...
24 #include <string_view> // for hash
25 #include <type_traits> // for enable_if_t
26 
27 #if !defined(__has_builtin) || !__has_builtin(__builtin_expect)
28 #define __builtin_expect(a, b) (a)
29 #endif
30 
31 #ifdef __SYCL_DEVICE_ONLY__
32 // `constexpr` could work because the implicit conversion from `float` to
33 // `_Float16` can be `constexpr`.
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
38 #else
39 #define __SYCL_CONSTEXPR_HALF
40 #endif
41 
42 namespace sycl {
43 inline namespace _V1 {
44 namespace detail::half_impl {
45 class half;
46 }
48 
49 namespace ext::intel::esimd::detail {
50 class WrapperElementTypeProxy;
51 } // namespace ext::intel::esimd::detail
52 
53 namespace detail {
54 
55 inline __SYCL_CONSTEXPR_HALF uint16_t float2Half(const float &Val) {
56  const uint32_t Bits = sycl::bit_cast<uint32_t>(Val);
57 
58  // Extract the sign from the float value
59  const uint16_t Sign = (Bits & 0x80000000) >> 16;
60  // Extract the fraction from the float value
61  const uint32_t Frac32 = Bits & 0x7fffff;
62  // Extract the exponent from the float value
63  const uint8_t Exp32 = (Bits & 0x7f800000) >> 23;
64  const int16_t Exp32Diff = Exp32 - 127;
65 
66  // initialize to 0, covers the case for 0 and small numbers
67  uint16_t Exp16 = 0, Frac16 = 0;
68 
69  if (__builtin_expect(Exp32Diff > 15, 0)) {
70  // Infinity and big numbers convert to infinity
71  Exp16 = 0x1f;
72  } else if (__builtin_expect(Exp32Diff > -14, 0)) {
73  // normal range for half type
74  Exp16 = Exp32Diff + 15;
75  // convert 23-bit mantissa to 10-bit mantissa.
76  Frac16 = Frac32 >> 13;
77  // Round the mantissa as given in OpenCL spec section : 6.1.1.1 The half
78  // data type.
79  // Round to nearest.
80  uint32_t roundBits = Frac32 & 0x1fff;
81  uint32_t halfway = 0x1000;
82  if (roundBits > halfway)
83  Frac16 += 1;
84  // Tie to even.
85  else if (roundBits == halfway)
86  Frac16 += Frac16 & 1;
87  } else if (__builtin_expect(Exp32Diff > -25, 0)) {
88  // subnormals
89  Frac16 = (Frac32 | (uint32_t(1) << 23)) >> (-Exp32Diff - 1);
90  }
91 
92  if (__builtin_expect(Exp32 == 0xff && Frac32 != 0, 0)) {
93  // corner case: FP32 is NaN
94  Exp16 = 0x1F;
95  Frac16 = 0x200;
96  }
97 
98  // Compose the final FP16 binary
99  uint16_t Ret = 0;
100  Ret |= Sign;
101  Ret |= Exp16 << 10;
102  Ret += Frac16; // Add the carry bit from operation Frac16 += 1;
103 
104  return Ret;
105 }
106 
107 inline __SYCL_CONSTEXPR_HALF float half2Float(const uint16_t &Val) {
108  // Extract the sign from the bits. It is 1 if the sign is negative
109  const uint32_t Sign = static_cast<uint32_t>(Val & 0x8000) << 16;
110  // Extract the exponent from the bits
111  const uint8_t Exp16 = (Val & 0x7c00) >> 10;
112  // Extract the fraction from the bits
113  uint16_t Frac16 = Val & 0x3ff;
114 
115  uint32_t Exp32 = 0;
116  if (__builtin_expect(Exp16 == 0x1f, 0)) {
117  Exp32 = 0xff;
118  } else if (__builtin_expect(Exp16 == 0, 0)) {
119  Exp32 = 0;
120  } else {
121  Exp32 = static_cast<uint32_t>(Exp16) + 112;
122  }
123  // corner case: subnormal -> normal
124  // The denormal number of FP16 can be represented by FP32, therefore we need
125  // to recover the exponent and recalculate the fration.
126  if (__builtin_expect(Exp16 == 0 && Frac16 != 0, 0)) {
127  uint8_t OffSet = 0;
128  do {
129  ++OffSet;
130  Frac16 <<= 1;
131  } while ((Frac16 & 0x400) != 0x400);
132  // mask the 9th bit
133  Frac16 &= 0x3ff;
134  Exp32 = 113 - OffSet;
135  }
136 
137  uint32_t Frac32 = Frac16 << 13;
138 
139  uint32_t Bits = 0;
140  Bits |= Sign;
141  Bits |= (Exp32 << 23);
142  Bits |= Frac32;
143  const float Result = sycl::bit_cast<float>(Bits);
144  return Result;
145 }
146 
147 namespace host_half_impl {
148 
149 // The main host half class
150 class __SYCL_EXPORT half {
151 public:
152  half() = default;
153  constexpr half(const half &) = default;
154  constexpr half(half &&) = default;
155 
156  __SYCL_CONSTEXPR_HALF half(const float &rhs) : Buf(float2Half(rhs)) {}
157 
158  constexpr half &operator=(const half &rhs) = default;
159 
160  // Operator +=, -=, *=, /=
162  *this = operator float() + static_cast<float>(rhs);
163  return *this;
164  }
165 
167  *this = operator float() - static_cast<float>(rhs);
168  return *this;
169  }
170 
172  *this = operator float() * static_cast<float>(rhs);
173  return *this;
174  }
175 
177  *this = operator float() / static_cast<float>(rhs);
178  return *this;
179  }
180 
181  // Operator ++, --
183  *this += 1;
184  return *this;
185  }
186 
188  half ret(*this);
189  operator++();
190  return ret;
191  }
192 
194  *this -= 1;
195  return *this;
196  }
197 
199  half ret(*this);
200  operator--();
201  return ret;
202  }
203 
204  // Operator neg
205  constexpr half &operator-() {
206  Buf ^= 0x8000;
207  return *this;
208  }
209 
210  // Operator float
211  __SYCL_CONSTEXPR_HALF operator float() const { return half2Float(Buf); }
212 
213  template <typename Key> friend struct std::hash;
214 
215  // Initialize underlying data
216  constexpr explicit half(uint16_t x) : Buf(x) {}
217 
218  friend class sycl::ext::intel::esimd::detail::WrapperElementTypeProxy;
219 
220 private:
221  uint16_t Buf;
222 };
223 
224 } // namespace host_half_impl
225 
226 namespace half_impl {
227 class half;
228 
229 // Several aliases are defined below:
230 // - StorageT: actual representation of half data type. It is used by scalar
231 // half values and by 'sycl::vec' class. On device side, it points to some
232 // native half data type, while on host some custom data type is used to
233 // emulate operations of 16-bit floating-point values
234 //
235 // - BIsRepresentationT: data type which is used by built-in functions. It is
236 // distinguished from StorageT, because on host, we can still operate on the
237 // wrapper itself and there is no sense in direct usage of underlying data
238 // type (too many changes required for BIs implementation without any
239 // foreseeable profits)
240 //
241 // - VecNStorageT - representation of N-element vector of halfs. Follows the
242 // same logic as StorageT
243 #ifdef __SYCL_DEVICE_ONLY__
244 using StorageT = _Float16;
245 using BIsRepresentationT = _Float16;
246 
247 using Vec2StorageT = StorageT __attribute__((ext_vector_type(2)));
248 using Vec3StorageT = StorageT __attribute__((ext_vector_type(3)));
249 using Vec4StorageT = StorageT __attribute__((ext_vector_type(4)));
250 using Vec8StorageT = StorageT __attribute__((ext_vector_type(8)));
251 using Vec16StorageT = StorageT __attribute__((ext_vector_type(16)));
252 #else // SYCL_DEVICE_ONLY
254 // No need to extract underlying data type for built-in functions operating on
255 // host
257 
258 // On the host side we cannot use OpenCL cl_half# types as an underlying type
259 // for vec because they are actually defined as an integer type under the
260 // hood. As a result half values will be converted to the integer and passed
261 // as a kernel argument which is expected to be floating point number.
262 #ifdef __INTEL_PREVIEW_BREAKING_CHANGES
263 using Vec2StorageT = std::array<StorageT, 2>;
264 using Vec3StorageT = std::array<StorageT, 3>;
265 using Vec4StorageT = std::array<StorageT, 4>;
266 using Vec8StorageT = std::array<StorageT, 8>;
267 using Vec16StorageT = std::array<StorageT, 16>;
268 #else // __INTEL_PREVIEW_BREAKING_CHANGES
269 template <int NumElements> struct half_vec {
270  alignas(
272 
274  template <typename... Ts,
275  typename = std::enable_if_t<(sizeof...(Ts) == NumElements) &&
276  (std::is_same_v<half, Ts> && ...)>>
277  __SYCL_CONSTEXPR_HALF half_vec(const Ts &...hs) : s{hs...} {}
278 
279  constexpr void initialize_data() {
280  for (size_t i = 0; i < NumElements; ++i) {
281  s[i] = StorageT(0.0f);
282  }
283  }
284 };
285 
291 #endif // __INTEL_PREVIEW_BREAKING_CHANGES
292 #endif // SYCL_DEVICE_ONLY
293 
294 #ifndef __SYCL_DEVICE_ONLY__
295 class half {
296 #else
297 class [[__sycl_detail__::__uses_aspects__(aspect::fp16)]] half {
298 #endif
299 public:
300  half() = default;
301  constexpr half(const half &) = default;
302  constexpr half(half &&) = default;
303 
304  __SYCL_CONSTEXPR_HALF half(const float &rhs) : Data(rhs) {}
305 
306  constexpr half &operator=(const half &rhs) = default;
307 
308 #ifndef __SYCL_DEVICE_ONLY__
309  // Since StorageT and BIsRepresentationT are different on host, these two
310  // helpers are required for 'vec' class
311  constexpr half(const detail::host_half_impl::half &rhs) : Data(rhs) {}
312  constexpr operator detail::host_half_impl::half() const { return Data; }
313 #endif // __SYCL_DEVICE_ONLY__
314 
315  // Operator +=, -=, *=, /=
317  Data += rhs.Data;
318  return *this;
319  }
320 
322  Data -= rhs.Data;
323  return *this;
324  }
325 
327  Data *= rhs.Data;
328  return *this;
329  }
330 
332  Data /= rhs.Data;
333  return *this;
334  }
335 
336  // Operator ++, --
338  *this += 1;
339  return *this;
340  }
341 
343  half ret(*this);
344  operator++();
345  return ret;
346  }
347 
349  *this -= 1;
350  return *this;
351  }
352 
354  half ret(*this);
355  operator--();
356  return ret;
357  }
358  __SYCL_CONSTEXPR_HALF friend half operator-(const half other) {
359  return half(-other.Data);
360  }
361 
362 // Operator +, -, *, /
363 #define OP(op, op_eq) \
364  __SYCL_CONSTEXPR_HALF friend half operator op(const half lhs, \
365  const half rhs) { \
366  half rtn = lhs; \
367  rtn op_eq rhs; \
368  return rtn; \
369  } \
370  __SYCL_CONSTEXPR_HALF friend double operator op(const half lhs, \
371  const double rhs) { \
372  double rtn = lhs; \
373  rtn op_eq rhs; \
374  return rtn; \
375  } \
376  __SYCL_CONSTEXPR_HALF friend double operator op(const double lhs, \
377  const half rhs) { \
378  double rtn = lhs; \
379  rtn op_eq rhs; \
380  return rtn; \
381  } \
382  __SYCL_CONSTEXPR_HALF friend float operator op(const half lhs, \
383  const float rhs) { \
384  float rtn = lhs; \
385  rtn op_eq rhs; \
386  return rtn; \
387  } \
388  __SYCL_CONSTEXPR_HALF friend float operator op(const float lhs, \
389  const half rhs) { \
390  float rtn = lhs; \
391  rtn op_eq rhs; \
392  return rtn; \
393  } \
394  __SYCL_CONSTEXPR_HALF friend half operator op(const half lhs, \
395  const int rhs) { \
396  half rtn = lhs; \
397  rtn op_eq rhs; \
398  return rtn; \
399  } \
400  __SYCL_CONSTEXPR_HALF friend half operator op(const int lhs, \
401  const half rhs) { \
402  half rtn = lhs; \
403  rtn op_eq rhs; \
404  return rtn; \
405  } \
406  __SYCL_CONSTEXPR_HALF friend half operator op(const half lhs, \
407  const long rhs) { \
408  half rtn = lhs; \
409  rtn op_eq rhs; \
410  return rtn; \
411  } \
412  __SYCL_CONSTEXPR_HALF friend half operator op(const long lhs, \
413  const half rhs) { \
414  half rtn = lhs; \
415  rtn op_eq rhs; \
416  return rtn; \
417  } \
418  __SYCL_CONSTEXPR_HALF friend half operator op(const half lhs, \
419  const long long rhs) { \
420  half rtn = lhs; \
421  rtn op_eq rhs; \
422  return rtn; \
423  } \
424  __SYCL_CONSTEXPR_HALF friend half operator op(const long long lhs, \
425  const half rhs) { \
426  half rtn = lhs; \
427  rtn op_eq rhs; \
428  return rtn; \
429  } \
430  __SYCL_CONSTEXPR_HALF friend half operator op(const half &lhs, \
431  const unsigned int &rhs) { \
432  half rtn = lhs; \
433  rtn op_eq rhs; \
434  return rtn; \
435  } \
436  __SYCL_CONSTEXPR_HALF friend half operator op(const unsigned int &lhs, \
437  const half &rhs) { \
438  half rtn = lhs; \
439  rtn op_eq rhs; \
440  return rtn; \
441  } \
442  __SYCL_CONSTEXPR_HALF friend half operator op(const half &lhs, \
443  const unsigned long &rhs) { \
444  half rtn = lhs; \
445  rtn op_eq rhs; \
446  return rtn; \
447  } \
448  __SYCL_CONSTEXPR_HALF friend half operator op(const unsigned long &lhs, \
449  const half &rhs) { \
450  half rtn = lhs; \
451  rtn op_eq rhs; \
452  return rtn; \
453  } \
454  __SYCL_CONSTEXPR_HALF friend half operator op( \
455  const half &lhs, const unsigned long long &rhs) { \
456  half rtn = lhs; \
457  rtn op_eq rhs; \
458  return rtn; \
459  } \
460  __SYCL_CONSTEXPR_HALF friend half operator op(const unsigned long long &lhs, \
461  const half &rhs) { \
462  half rtn = lhs; \
463  rtn op_eq rhs; \
464  return rtn; \
465  }
466  OP(+, +=)
467  OP(-, -=)
468  OP(*, *=)
469  OP(/, /=)
470 
471 #undef OP
472 
473 // Operator ==, !=, <, >, <=, >=
474 #define OP(op) \
475  __SYCL_CONSTEXPR_HALF friend bool operator op(const half &lhs, \
476  const half &rhs) { \
477  return lhs.Data op rhs.Data; \
478  } \
479  __SYCL_CONSTEXPR_HALF friend bool operator op(const half &lhs, \
480  const double &rhs) { \
481  return lhs.Data op rhs; \
482  } \
483  __SYCL_CONSTEXPR_HALF friend bool operator op(const double &lhs, \
484  const half &rhs) { \
485  return lhs op rhs.Data; \
486  } \
487  __SYCL_CONSTEXPR_HALF friend bool operator op(const half &lhs, \
488  const float &rhs) { \
489  return lhs.Data op rhs; \
490  } \
491  __SYCL_CONSTEXPR_HALF friend bool operator op(const float &lhs, \
492  const half &rhs) { \
493  return lhs op rhs.Data; \
494  } \
495  __SYCL_CONSTEXPR_HALF friend bool operator op(const half &lhs, \
496  const int &rhs) { \
497  return lhs.Data op rhs; \
498  } \
499  __SYCL_CONSTEXPR_HALF friend bool operator op(const int &lhs, \
500  const half &rhs) { \
501  return lhs op rhs.Data; \
502  } \
503  __SYCL_CONSTEXPR_HALF friend bool operator op(const half &lhs, \
504  const long &rhs) { \
505  return lhs.Data op rhs; \
506  } \
507  __SYCL_CONSTEXPR_HALF friend bool operator op(const long &lhs, \
508  const half &rhs) { \
509  return lhs op rhs.Data; \
510  } \
511  __SYCL_CONSTEXPR_HALF friend bool operator op(const half &lhs, \
512  const long long &rhs) { \
513  return lhs.Data op rhs; \
514  } \
515  __SYCL_CONSTEXPR_HALF friend bool operator op(const long long &lhs, \
516  const half &rhs) { \
517  return lhs op rhs.Data; \
518  } \
519  __SYCL_CONSTEXPR_HALF friend bool operator op(const half &lhs, \
520  const unsigned int &rhs) { \
521  return lhs.Data op rhs; \
522  } \
523  __SYCL_CONSTEXPR_HALF friend bool operator op(const unsigned int &lhs, \
524  const half &rhs) { \
525  return lhs op rhs.Data; \
526  } \
527  __SYCL_CONSTEXPR_HALF friend bool operator op(const half &lhs, \
528  const unsigned long &rhs) { \
529  return lhs.Data op rhs; \
530  } \
531  __SYCL_CONSTEXPR_HALF friend bool operator op(const unsigned long &lhs, \
532  const half &rhs) { \
533  return lhs op rhs.Data; \
534  } \
535  __SYCL_CONSTEXPR_HALF friend bool operator op( \
536  const half &lhs, const unsigned long long &rhs) { \
537  return lhs.Data op rhs; \
538  } \
539  __SYCL_CONSTEXPR_HALF friend bool operator op(const unsigned long long &lhs, \
540  const half &rhs) { \
541  return lhs op rhs.Data; \
542  }
543  OP(==)
544  OP(!=)
545  OP(<)
546  OP(>)
547  OP(<=)
548  OP(>=)
549 
550 #undef OP
551 
552  // Operator float
553  __SYCL_CONSTEXPR_HALF operator float() const {
554  return static_cast<float>(Data);
555  }
556 
557  // Operator << and >>
558  inline friend std::ostream &operator<<(std::ostream &O,
559  sycl::half const &rhs) {
560  O << static_cast<float>(rhs);
561  return O;
562  }
563 
564  inline friend std::istream &operator>>(std::istream &I, sycl::half &rhs) {
565  float ValFloat = 0.0f;
566  I >> ValFloat;
567  rhs = ValFloat;
568  return I;
569  }
570 
571  template <typename Key> friend struct std::hash;
572 
573  friend class sycl::ext::intel::esimd::detail::WrapperElementTypeProxy;
574 
575 private:
576  StorageT Data;
577 };
578 } // namespace half_impl
579 
580 // According to the C++ standard, math functions from cmath/math.h should work
581 // only on arithmetic types. We can't specify half type as arithmetic/floating
582 // point(via std::is_floating_point) since only float, double and long double
583 // types are "floating point" according to the standard. In order to use half
584 // type with these math functions we cast half to float using template
585 // function helper.
586 template <typename T> inline T cast_if_host_half(T val) { return val; }
587 
589  return static_cast<float>(val);
590 }
591 
592 } // namespace detail
593 
594 } // namespace _V1
595 } // namespace sycl
596 
597 // Partial specialization of some functions in namespace `std`
598 namespace std {
599 
600 // Partial specialization of `std::hash<sycl::half>`
601 template <> struct hash<sycl::half> {
602  size_t operator()(sycl::half const &Key) const noexcept {
603  return hash<uint16_t>{}(reinterpret_cast<const uint16_t &>(Key));
604  }
605 };
606 
607 // Partial specialization of `std::numeric<sycl::half>`
608 template <> struct numeric_limits<sycl::half> {
609  // All following values are either calculated based on description of each
610  // function/value on https://en.cppreference.com/w/cpp/types/numeric_limits,
611  // or cl_platform.h.
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;
635 
637  return 6.103515625e-05f; // half minimum value
638  }
639 
641  return 65504.0f; // half maximum value
642  }
643 
645  return -65504.0f; // -1*(half maximum value)
646  }
647 
649  return 9.765625e-04f; // half epsilon
650  }
651 
653  return 0.5f;
654  }
655 
656  static constexpr const sycl::half infinity() noexcept {
657 #ifdef __SYCL_DEVICE_ONLY__
658  return __builtin_huge_valf();
659 #else
660  return sycl::detail::host_half_impl::half(static_cast<uint16_t>(0x7C00));
661 #endif
662  }
663 
665  return __builtin_nanf("");
666  }
667 
669  return __builtin_nansf("");
670  }
671 
673  return 5.96046e-08f;
674  }
675 };
676 
677 } // namespace std
678 
679 #undef __SYCL_CONSTEXPR_HALF
680 #undef _CPP14_CONSTEXPR
friend std::istream & operator>>(std::istream &I, sycl::half &rhs)
Definition: half_type.hpp:564
__SYCL_CONSTEXPR_HALF half & operator/=(const half &rhs)
Definition: half_type.hpp:331
__SYCL_CONSTEXPR_HALF half & operator+=(const half &rhs)
Definition: half_type.hpp:316
__SYCL_CONSTEXPR_HALF half & operator-=(const half &rhs)
Definition: half_type.hpp:321
__SYCL_CONSTEXPR_HALF half operator--(int)
Definition: half_type.hpp:353
__SYCL_CONSTEXPR_HALF half & operator*=(const half &rhs)
Definition: half_type.hpp:326
__SYCL_CONSTEXPR_HALF friend half operator-(const half other)
Definition: half_type.hpp:358
constexpr half(half &&)=default
constexpr half(const detail::host_half_impl::half &rhs)
Definition: half_type.hpp:311
constexpr half(const half &)=default
constexpr half & operator=(const half &rhs)=default
__SYCL_CONSTEXPR_HALF half & operator--()
Definition: half_type.hpp:348
__SYCL_CONSTEXPR_HALF half(const float &rhs)
Definition: half_type.hpp:304
friend std::ostream & operator<<(std::ostream &O, sycl::half const &rhs)
Definition: half_type.hpp:558
__SYCL_CONSTEXPR_HALF half & operator++()
Definition: half_type.hpp:337
__SYCL_CONSTEXPR_HALF half operator++(int)
Definition: half_type.hpp:342
constexpr half(half &&)=default
constexpr half & operator=(const half &rhs)=default
__SYCL_CONSTEXPR_HALF half & operator-=(const half &rhs)
Definition: half_type.hpp:166
__SYCL_CONSTEXPR_HALF half & operator+=(const half &rhs)
Definition: half_type.hpp:161
__SYCL_CONSTEXPR_HALF half(const float &rhs)
Definition: half_type.hpp:156
__SYCL_CONSTEXPR_HALF half & operator++()
Definition: half_type.hpp:182
__SYCL_CONSTEXPR_HALF half & operator/=(const half &rhs)
Definition: half_type.hpp:176
__SYCL_CONSTEXPR_HALF half operator--(int)
Definition: half_type.hpp:198
__SYCL_CONSTEXPR_HALF half & operator--()
Definition: half_type.hpp:193
__SYCL_CONSTEXPR_HALF half operator++(int)
Definition: half_type.hpp:187
constexpr half(const half &)=default
__SYCL_CONSTEXPR_HALF half & operator*=(const half &rhs)
Definition: half_type.hpp:171
#define __builtin_expect(a, b)
Definition: half_type.hpp:28
#define OP(op, op_eq)
Definition: half_type.hpp:474
#define __SYCL_CONSTEXPR_HALF
Definition: half_type.hpp:39
half_vec< 16 > Vec16StorageT
Definition: half_type.hpp:290
detail::host_half_impl::half StorageT
Definition: half_type.hpp:253
T cast_if_host_half(T val)
Definition: half_type.hpp:586
__SYCL_CONSTEXPR_HALF uint16_t float2Half(const float &Val)
Definition: half_type.hpp:55
__SYCL_CONSTEXPR_HALF float half2Float(const uint16_t &Val)
Definition: half_type.hpp:107
__attribute__((destructor(110))) static void syclUnload()
multi_ptr & operator--()
Definition: multi_ptr.hpp:1027
multi_ptr & operator++()
Definition: multi_ptr.hpp:1018
sycl::detail::half_impl::half half
Definition: aliases.hpp:101
Definition: access.hpp:18
_Abi const simd< _Tp, _Abi > & noexcept
Definition: simd.hpp:1324
size_t operator()(sycl::half const &Key) const noexcept
Definition: half_type.hpp:602
static constexpr const sycl::half infinity() noexcept
Definition: half_type.hpp:656
static __SYCL_CONSTEXPR_HALF const sycl::half round_error() noexcept
Definition: half_type.hpp:652
static __SYCL_CONSTEXPR_HALF const sycl::half quiet_NaN() noexcept
Definition: half_type.hpp:664
static __SYCL_CONSTEXPR_HALF const sycl::half lowest() noexcept
Definition: half_type.hpp:644
static __SYCL_CONSTEXPR_HALF const sycl::half denorm_min() noexcept
Definition: half_type.hpp:672
static __SYCL_CONSTEXPR_HALF const sycl::half epsilon() noexcept
Definition: half_type.hpp:648
static __SYCL_CONSTEXPR_HALF const sycl::half signaling_NaN() noexcept
Definition: half_type.hpp:668
__SYCL_CONSTEXPR_HALF half_vec(const Ts &...hs)
Definition: half_type.hpp:277
__SYCL_CONSTEXPR_HALF half_vec()
Definition: half_type.hpp:273