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