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