DPC++ Runtime
Runtime libraries for oneAPI Data Parallel C++
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 <array>
16 #include <cmath>
17 #include <cstdint>
18 #include <functional>
19 #include <iostream>
20 #include <limits>
21 
22 #if !__has_builtin(__builtin_expect)
23 #define __builtin_expect(a, b) (a)
24 #endif
25 
26 #ifdef __SYCL_DEVICE_ONLY__
27 // `constexpr` could work because the implicit conversion from `float` to
28 // `_Float16` can be `constexpr`.
29 #define __SYCL_CONSTEXPR_HALF constexpr
30 #elif __cpp_lib_bit_cast || __has_builtin(__builtin_bit_cast)
31 #define __SYCL_CONSTEXPR_HALF constexpr
32 #else
33 #define __SYCL_CONSTEXPR_HALF
34 #endif
35 
37 namespace sycl {
38 namespace detail {
39 
40 inline __SYCL_CONSTEXPR_HALF uint16_t float2Half(const float &Val) {
41  const uint32_t Bits = sycl::bit_cast<uint32_t>(Val);
42 
43  // Extract the sign from the float value
44  const uint16_t Sign = (Bits & 0x80000000) >> 16;
45  // Extract the fraction from the float value
46  const uint32_t Frac32 = Bits & 0x7fffff;
47  // Extract the exponent from the float value
48  const uint8_t Exp32 = (Bits & 0x7f800000) >> 23;
49  const int16_t Exp32Diff = Exp32 - 127;
50 
51  // intialize to 0, covers the case for 0 and small numbers
52  uint16_t Exp16 = 0, Frac16 = 0;
53 
54  if (__builtin_expect(Exp32Diff > 15, 0)) {
55  // Infinity and big numbers convert to infinity
56  Exp16 = 0x1f;
57  } else if (__builtin_expect(Exp32Diff > -14, 0)) {
58  // normal range for half type
59  Exp16 = Exp32Diff + 15;
60  // convert 23-bit mantissa to 10-bit mantissa.
61  Frac16 = Frac32 >> 13;
62  // Round the mantissa as given in OpenCL spec section : 6.1.1.1 The half
63  // data type.
64  if (Frac32 >> 12 & 0x01)
65  Frac16 += 1;
66  } else if (__builtin_expect(Exp32Diff > -24, 0)) {
67  // subnormals
68  Frac16 = (Frac32 | (uint32_t(1) << 23)) >> (-Exp32Diff - 1);
69  }
70 
71  if (__builtin_expect(Exp32 == 0xff && Frac32 != 0, 0)) {
72  // corner case: FP32 is NaN
73  Exp16 = 0x1F;
74  Frac16 = 0x200;
75  }
76 
77  // Compose the final FP16 binary
78  uint16_t Ret = 0;
79  Ret |= Sign;
80  Ret |= Exp16 << 10;
81  Ret += Frac16; // Add the carry bit from operation Frac16 += 1;
82 
83  return Ret;
84 }
85 
86 inline __SYCL_CONSTEXPR_HALF float half2Float(const uint16_t &Val) {
87  // Extract the sign from the bits. It is 1 if the sign is negative
88  const uint32_t Sign = static_cast<uint32_t>(Val & 0x8000) << 16;
89  // Extract the exponent from the bits
90  const uint8_t Exp16 = (Val & 0x7c00) >> 10;
91  // Extract the fraction from the bits
92  uint16_t Frac16 = Val & 0x3ff;
93 
94  uint32_t Exp32 = 0;
95  if (__builtin_expect(Exp16 == 0x1f, 0)) {
96  Exp32 = 0xff;
97  } else if (__builtin_expect(Exp16 == 0, 0)) {
98  Exp32 = 0;
99  } else {
100  Exp32 = static_cast<uint32_t>(Exp16) + 112;
101  }
102  // corner case: subnormal -> normal
103  // The denormal number of FP16 can be represented by FP32, therefore we need
104  // to recover the exponent and recalculate the fration.
105  if (__builtin_expect(Exp16 == 0 && Frac16 != 0, 0)) {
106  uint8_t OffSet = 0;
107  do {
108  ++OffSet;
109  Frac16 <<= 1;
110  } while ((Frac16 & 0x400) != 0x400);
111  // mask the 9th bit
112  Frac16 &= 0x3ff;
113  Exp32 = 113 - OffSet;
114  }
115 
116  uint32_t Frac32 = Frac16 << 13;
117 
118  uint32_t Bits = 0;
119  Bits |= Sign;
120  Bits |= (Exp32 << 23);
121  Bits |= Frac32;
122  const float Result = sycl::bit_cast<float>(Bits);
123  return Result;
124 }
125 
126 namespace host_half_impl {
127 
128 // This class is legacy and it is needed only to avoid breaking ABI
129 class __SYCL_EXPORT half {
130 public:
131  half() = default;
132  constexpr half(const half &) = default;
133  constexpr half(half &&) = default;
134 
135  half(const float &rhs);
136 
137  half &operator=(const half &rhs) = default;
138 
139  // Operator +=, -=, *=, /=
140  half &operator+=(const half &rhs);
141 
142  half &operator-=(const half &rhs);
143 
144  half &operator*=(const half &rhs);
145 
146  half &operator/=(const half &rhs);
147 
148  // Operator ++, --
150  *this += 1;
151  return *this;
152  }
153 
155  half ret(*this);
156  operator++();
157  return ret;
158  }
159 
161  *this -= 1;
162  return *this;
163  }
164 
166  half ret(*this);
167  operator--();
168  return ret;
169  }
170 
171  // Operator neg
172  constexpr half &operator-() {
173  Buf ^= 0x8000;
174  return *this;
175  }
176 
177  // Operator float
178  operator float() const;
179 
180  template <typename Key> friend struct std::hash;
181 
182  // Initialize underlying data
183  constexpr explicit half(uint16_t x) : Buf(x) {}
184 
185 private:
186  uint16_t Buf;
187 };
188 
189 // The main host half class
190 class __SYCL_EXPORT half_v2 {
191 public:
192  half_v2() = default;
193  constexpr half_v2(const half_v2 &) = default;
194  constexpr half_v2(half_v2 &&) = default;
195 
196  __SYCL_CONSTEXPR_HALF half_v2(const float &rhs) : Buf(float2Half(rhs)) {}
197 
198  constexpr half_v2 &operator=(const half_v2 &rhs) = default;
199 
200  // Operator +=, -=, *=, /=
202  *this = operator float() + static_cast<float>(rhs);
203  return *this;
204  }
205 
207  *this = operator float() - static_cast<float>(rhs);
208  return *this;
209  }
210 
212  *this = operator float() * static_cast<float>(rhs);
213  return *this;
214  }
215 
217  *this = operator float() / static_cast<float>(rhs);
218  return *this;
219  }
220 
221  // Operator ++, --
223  *this += 1;
224  return *this;
225  }
226 
228  half_v2 ret(*this);
229  operator++();
230  return ret;
231  }
232 
234  *this -= 1;
235  return *this;
236  }
237 
239  half_v2 ret(*this);
240  operator--();
241  return ret;
242  }
243 
244  // Operator neg
245  constexpr half_v2 &operator-() {
246  Buf ^= 0x8000;
247  return *this;
248  }
249 
250  // Operator float
251  __SYCL_CONSTEXPR_HALF operator float() const { return half2Float(Buf); }
252 
253  template <typename Key> friend struct std::hash;
254 
255  // Initialize underlying data
256  constexpr explicit half_v2(uint16_t x) : Buf(x) {}
257 
258 private:
259  uint16_t Buf;
260 };
261 
262 } // namespace host_half_impl
263 
264 namespace half_impl {
265 class half;
266 
267 // Several aliases are defined below:
268 // - StorageT: actual representation of half data type. It is used by scalar
269 // half values and by 'cl::sycl::vec' class. On device side, it points to some
270 // native half data type, while on host some custom data type is used to
271 // emulate operations of 16-bit floating-point values
272 //
273 // - BIsRepresentationT: data type which is used by built-in functions. It is
274 // distinguished from StorageT, because on host, we can still operate on the
275 // wrapper itself and there is no sense in direct usage of underlying data
276 // type (too many changes required for BIs implementation without any
277 // foreseeable profits)
278 //
279 // - VecNStorageT - representation of N-element vector of halfs. Follows the
280 // same logic as StorageT
281 #ifdef __SYCL_DEVICE_ONLY__
282  using StorageT = _Float16;
283  using BIsRepresentationT = _Float16;
284 
285  using Vec2StorageT = StorageT __attribute__((ext_vector_type(2)));
286  using Vec3StorageT = StorageT __attribute__((ext_vector_type(3)));
287  using Vec4StorageT = StorageT __attribute__((ext_vector_type(4)));
288  using Vec8StorageT = StorageT __attribute__((ext_vector_type(8)));
289  using Vec16StorageT = StorageT __attribute__((ext_vector_type(16)));
290 #else
292 // No need to extract underlying data type for built-in functions operating on
293 // host
295 
296 // On the host side we cannot use OpenCL cl_half# types as an underlying type
297 // for vec because they are actually defined as an integer type under the
298 // hood. As a result half values will be converted to the integer and passed
299 // as a kernel argument which is expected to be floating point number.
300 template <int NumElements> struct half_vec {
302  StorageT s[NumElements];
303 
304  __SYCL_CONSTEXPR_HALF half_vec() : s{0.0f} { initialize_data(); }
305  constexpr void initialize_data() {
306  for (size_t i = 0; i < NumElements; ++i) {
307  s[i] = StorageT(0.0f);
308  }
309  }
310 };
311 
317 #endif
318 
319 class half {
320 public:
321  half() = default;
322  constexpr half(const half &) = default;
323  constexpr half(half &&) = default;
324 
325  __SYCL_CONSTEXPR_HALF half(const float &rhs) : Data(rhs) {}
326 
327  constexpr half &operator=(const half &rhs) = default;
328 
329 #ifndef __SYCL_DEVICE_ONLY__
330  // Since StorageT and BIsRepresentationT are different on host, these two
331  // helpers are required for 'vec' class
332  constexpr half(const detail::host_half_impl::half_v2 &rhs) : Data(rhs){};
333  constexpr operator detail::host_half_impl::half_v2() const { return Data; }
334 #endif // __SYCL_DEVICE_ONLY__
335 
336  // Operator +=, -=, *=, /=
338  Data += rhs.Data;
339  return *this;
340  }
341 
343  Data -= rhs.Data;
344  return *this;
345  }
346 
348  Data *= rhs.Data;
349  return *this;
350  }
351 
353  Data /= rhs.Data;
354  return *this;
355  }
356 
357  // Operator ++, --
359  *this += 1;
360  return *this;
361  }
362 
364  half ret(*this);
365  operator++();
366  return ret;
367  }
368 
370  *this -= 1;
371  return *this;
372  }
373 
375  half ret(*this);
376  operator--();
377  return ret;
378  }
379  constexpr half &operator-() {
380  Data = -Data;
381  return *this;
382  }
383  constexpr half operator-() const {
384  half r = *this;
385  return -r;
386  }
387  // Operator float
388  __SYCL_CONSTEXPR_HALF operator float() const {
389  return static_cast<float>(Data);
390  }
391 
392  template <typename Key> friend struct std::hash;
393 
394 private:
395  StorageT Data;
396 };
397 } // namespace half_impl
398 
399 // Accroding to C++ standard math functions from cmath/math.h should work only
400 // on arithmetic types. We can't specify half type as arithmetic/floating
401 // point(via std::is_floating_point) since only float, double and long double
402 // types are "floating point" according to the standard. In order to use half
403 // type with these math functions we cast half to float using template
404 // function helper.
405 template <typename T> inline T cast_if_host_half(T val) { return val; }
406 
408  return static_cast<float>(val);
409 }
410 
411 } // namespace detail
412 
413 } // namespace sycl
414 } // __SYCL_INLINE_NAMESPACE(cl)
415 
416 // Partial specialization of some functions in namespace `std`
417 namespace std {
418 
419 // Partial specialization of `std::hash<cl::sycl::half>`
420 template <> struct hash<cl::sycl::half> {
421  size_t operator()(cl::sycl::half const &Key) const noexcept {
422  return hash<uint16_t>{}(reinterpret_cast<const uint16_t &>(Key));
423  }
424 };
425 
426 // Partial specialization of `std::numeric<cl::sycl::half>`
427 template <> struct numeric_limits<cl::sycl::half> {
428  // All following values are either calculated based on description of each
429  // function/value on https://en.cppreference.com/w/cpp/types/numeric_limits,
430  // or cl_platform.h.
431  static constexpr bool is_specialized = true;
432  static constexpr bool is_signed = true;
433  static constexpr bool is_integer = false;
434  static constexpr bool is_exact = false;
435  static constexpr bool has_infinity = true;
436  static constexpr bool has_quiet_NaN = true;
437  static constexpr bool has_signaling_NaN = true;
438  static constexpr float_denorm_style has_denorm = denorm_present;
439  static constexpr bool has_denorm_loss = false;
440  static constexpr bool tinyness_before = false;
441  static constexpr bool traps = false;
442  static constexpr int max_exponent10 = 4;
443  static constexpr int max_exponent = 16;
444  static constexpr int min_exponent10 = -4;
445  static constexpr int min_exponent = -13;
446  static constexpr int radix = 2;
447  static constexpr int max_digits10 = 5;
448  static constexpr int digits = 11;
449  static constexpr bool is_bounded = true;
450  static constexpr int digits10 = 3;
451  static constexpr bool is_modulo = false;
452  static constexpr bool is_iec559 = true;
453  static constexpr float_round_style round_style = round_to_nearest;
454 
455  static __SYCL_CONSTEXPR_HALF const cl::sycl::half(min)() noexcept {
456  return 6.103515625e-05f; // half minimum value
457  }
458 
459  static __SYCL_CONSTEXPR_HALF const cl::sycl::half(max)() noexcept {
460  return 65504.0f; // half maximum value
461  }
462 
463  static __SYCL_CONSTEXPR_HALF const cl::sycl::half lowest() noexcept {
464  return -65504.0f; // -1*(half maximum value)
465  }
466 
467  static __SYCL_CONSTEXPR_HALF const cl::sycl::half epsilon() noexcept {
468  return 9.765625e-04f; // half epsilon
469  }
470 
472  return 0.5f;
473  }
474 
475  static constexpr const cl::sycl::half infinity() noexcept {
476 #ifdef __SYCL_DEVICE_ONLY__
477  return __builtin_huge_valf();
478 #else
480  static_cast<uint16_t>(0x7C00));
481 #endif
482  }
483 
485  return __builtin_nanf("");
486  }
487 
489  return __builtin_nansf("");
490  }
491 
493  return 5.96046e-08f;
494  }
495 };
496 
497 } // namespace std
498 
499 inline std::ostream &operator<<(std::ostream &O, cl::sycl::half const &rhs) {
500  O << static_cast<float>(rhs);
501  return O;
502 }
503 
504 inline std::istream &operator>>(std::istream &I, cl::sycl::half &rhs) {
505  float ValFloat = 0.0f;
506  I >> ValFloat;
507  rhs = ValFloat;
508  return I;
509 }
510 
511 #undef __SYCL_CONSTEXPR_HALF
512 #undef _CPP14_CONSTEXPR
cl::sycl::detail::half_impl::Vec16StorageT
half_vec< 16 > Vec16StorageT
Definition: half_type.hpp:316
cl::sycl::detail::host_half_impl::half
Definition: half_type.hpp:129
cl::sycl::detail::half_impl::half::operator-=
__SYCL_CONSTEXPR_HALF half & operator-=(const half &rhs)
Definition: half_type.hpp:342
T
cl::sycl::detail::__attribute__
__attribute__((destructor(110))) static void syclUnload()
Definition: global_handler.cpp:181
cl::sycl::detail::host_half_impl::half::operator++
half operator++(int)
Definition: half_type.hpp:154
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:211
cl::sycl::detail::host_half_impl::half::operator--
half & operator--()
Definition: half_type.hpp:160
cl::sycl::detail::host_half_impl::half::half
constexpr half(uint16_t x)
Definition: half_type.hpp:183
std::numeric_limits< cl::sycl::half >::lowest
static const __SYCL_CONSTEXPR_HALF cl::sycl::half lowest() noexcept
Definition: half_type.hpp:463
cl::sycl::detail::half_impl::Vec2StorageT
half_vec< 2 > Vec2StorageT
Definition: half_type.hpp:312
cl::sycl::detail::cast_if_host_half
float cast_if_host_half(half_impl::half val)
Definition: half_type.hpp:407
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:332
cl::sycl::detail::half_impl::BIsRepresentationT
half BIsRepresentationT
Definition: half_type.hpp:294
cl::sycl::detail::half_impl::half::operator--
__SYCL_CONSTEXPR_HALF half & operator--()
Definition: half_type.hpp:369
operator>>
std::istream & operator>>(std::istream &I, cl::sycl::half &rhs)
Definition: half_type.hpp:504
cl::sycl::detail::host_half_impl::half_v2::operator++
__SYCL_CONSTEXPR_HALF half_v2 operator++(int)
Definition: half_type.hpp:227
cl::sycl::detail::host_half_impl::half_v2::operator++
__SYCL_CONSTEXPR_HALF half_v2 & operator++()
Definition: half_type.hpp:222
cl::sycl::detail::half_impl::half_vec::initialize_data
constexpr void initialize_data()
Definition: half_type.hpp:305
__SYCL_CONSTEXPR_HALF
#define __SYCL_CONSTEXPR_HALF
Definition: half_type.hpp:33
cl::sycl::detail::half_impl::half::operator--
__SYCL_CONSTEXPR_HALF half operator--(int)
Definition: half_type.hpp:374
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:484
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:492
cl::sycl::detail::half_impl::Vec3StorageT
half_vec< 3 > Vec3StorageT
Definition: half_type.hpp:313
cl::sycl::detail::half_impl::half_vec::half_vec
__SYCL_CONSTEXPR_HALF half_vec()
Definition: half_type.hpp:304
cl::sycl::detail::half_impl::half_vec
Definition: half_type.hpp:300
cl::sycl::detail::half_impl::half
Definition: half_type.hpp:319
cl::sycl::detail::half_impl::half::operator+=
__SYCL_CONSTEXPR_HALF half & operator+=(const half &rhs)
Definition: half_type.hpp:337
cl::sycl::detail::host_half_impl::half_v2::half_v2
constexpr half_v2(uint16_t x)
Definition: half_type.hpp:256
cl::sycl::detail::half_impl::Vec8StorageT
half_vec< 8 > Vec8StorageT
Definition: half_type.hpp:315
cl::sycl::detail::half_impl::half::half
__SYCL_CONSTEXPR_HALF half(const float &rhs)
Definition: half_type.hpp:325
defines.hpp
cl::sycl::detail::half_impl::half::operator-
constexpr half & operator-()
Definition: half_type.hpp:379
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:216
cl::sycl::image_channel_order::r
@ r
cl::sycl::detail::half_impl::half::operator++
__SYCL_CONSTEXPR_HALF half & operator++()
Definition: half_type.hpp:358
cl::sycl::detail::half_impl::half::operator-
constexpr half operator-() const
Definition: half_type.hpp:383
cl::sycl::detail::host_half_impl::half_v2::operator+=
__SYCL_CONSTEXPR_HALF half_v2 & operator+=(const half_v2 &rhs)
Definition: half_type.hpp:201
std::numeric_limits< cl::sycl::half >::round_error
static const __SYCL_CONSTEXPR_HALF cl::sycl::half round_error() noexcept
Definition: half_type.hpp:471
cl::sycl::detail::host_half_impl::half_v2
Definition: half_type.hpp:190
cl::sycl::detail::float2Half
__SYCL_CONSTEXPR_HALF uint16_t float2Half(const float &Val)
Definition: half_type.hpp:40
cl::sycl::detail::half2Float
__SYCL_CONSTEXPR_HALF float half2Float(const uint16_t &Val)
Definition: half_type.hpp:86
std
Definition: accessor.hpp:2397
cl::sycl::detail::host_half_impl::half_v2::operator--
__SYCL_CONSTEXPR_HALF half_v2 operator--(int)
Definition: half_type.hpp:238
cl::sycl::detail::half_impl::Vec4StorageT
half_vec< 4 > Vec4StorageT
Definition: half_type.hpp:314
uint16_t
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:352
cl::sycl::detail::host_half_impl::half::operator-
constexpr half & operator-()
Definition: half_type.hpp:172
cl::sycl::detail::half_impl::StorageT
detail::host_half_impl::half_v2 StorageT
Definition: half_type.hpp:291
cl::sycl::detail::host_half_impl::half_v2::operator-=
__SYCL_CONSTEXPR_HALF half_v2 & operator-=(const half_v2 &rhs)
Definition: half_type.hpp:206
operator<<
std::ostream & operator<<(std::ostream &O, cl::sycl::half const &rhs)
Definition: half_type.hpp:499
__builtin_expect
#define __builtin_expect(a, b)
Definition: half_type.hpp:23
cl::sycl::detail::host_half_impl::half::operator++
half & operator++()
Definition: half_type.hpp:149
std::numeric_limits< cl::sycl::half >::signaling_NaN
static const __SYCL_CONSTEXPR_HALF cl::sycl::half signaling_NaN() noexcept
Definition: half_type.hpp:488
cl::sycl::detail::host_half_impl::half::operator--
half operator--(int)
Definition: half_type.hpp:165
std::numeric_limits< cl::sycl::half >::epsilon
static const __SYCL_CONSTEXPR_HALF cl::sycl::half epsilon() noexcept
Definition: half_type.hpp:467
cl::sycl::detail::host_half_impl::half_v2::half_v2
__SYCL_CONSTEXPR_HALF half_v2(const float &rhs)
Definition: half_type.hpp:196
cl::sycl::detail::host_half_impl::half_v2::operator-
constexpr half_v2 & operator-()
Definition: half_type.hpp:245
std::hash< cl::sycl::half >::operator()
size_t operator()(cl::sycl::half const &Key) const noexcept
Definition: half_type.hpp:421
cl::sycl::detail::half_impl::half::operator*=
__SYCL_CONSTEXPR_HALF half & operator*=(const half &rhs)
Definition: half_type.hpp:347
cl::sycl::detail::half_impl::half::operator++
__SYCL_CONSTEXPR_HALF half operator++(int)
Definition: half_type.hpp:363
cl::sycl::detail::host_half_impl::half_v2::operator--
__SYCL_CONSTEXPR_HALF half_v2 & operator--()
Definition: half_type.hpp:233
std::numeric_limits< cl::sycl::half >::infinity
static constexpr const cl::sycl::half infinity() noexcept
Definition: half_type.hpp:475
__SYCL_INLINE_NAMESPACE
#define __SYCL_INLINE_NAMESPACE(X)
Definition: defines_elementary.hpp:12