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