Please, help us to better know about our user community by answering the following short survey: https://forms.gle/wpyrxWi18ox9Z5ae9
Eigen  3.4.0
BFloat16.h
1 /* Copyright 2017 The TensorFlow Authors. All Rights Reserved.
2 
3 Licensed under the Apache License, Version 2.0 (the "License");
4 you may not use this file except in compliance with the License.
5 You may obtain a copy of the License at
6 
7  http://www.apache.org/licenses/LICENSE-2.0
8 
9 Unless required by applicable law or agreed to in writing, software
10 distributed under the License is distributed on an "AS IS" BASIS,
11 WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
12 See the License for the specific language governing permissions and
13 limitations under the License.
14 ==============================================================================*/
15 
16 #ifndef EIGEN_BFLOAT16_H
17 #define EIGEN_BFLOAT16_H
18 
19 #define BF16_PACKET_FUNCTION(PACKET_F, PACKET_BF16, METHOD) \
20  template <> \
21  EIGEN_DEFINE_FUNCTION_ALLOWING_MULTIPLE_DEFINITIONS EIGEN_UNUSED \
22  PACKET_BF16 METHOD<PACKET_BF16>(const PACKET_BF16& _x) { \
23  return F32ToBf16(METHOD<PACKET_F>(Bf16ToF32(_x))); \
24  }
25 
26 namespace Eigen {
27 
28 struct bfloat16;
29 
30 namespace bfloat16_impl {
31 
32 // Make our own __bfloat16_raw definition.
33 struct __bfloat16_raw {
34  EIGEN_DEVICE_FUNC EIGEN_CONSTEXPR __bfloat16_raw() : value(0) {}
35  explicit EIGEN_DEVICE_FUNC EIGEN_CONSTEXPR __bfloat16_raw(unsigned short raw) : value(raw) {}
36  unsigned short value;
37 };
38 
39 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC EIGEN_CONSTEXPR __bfloat16_raw raw_uint16_to_bfloat16(unsigned short value);
40 template <bool AssumeArgumentIsNormalOrInfinityOrZero>
41 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC __bfloat16_raw float_to_bfloat16_rtne(float ff);
42 // Forward declarations of template specializations, to avoid Visual C++ 2019 errors, saying:
43 // > error C2908: explicit specialization; 'float_to_bfloat16_rtne' has already been instantiated
44 template <>
45 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC __bfloat16_raw float_to_bfloat16_rtne<false>(float ff);
46 template <>
47 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC __bfloat16_raw float_to_bfloat16_rtne<true>(float ff);
48 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC float bfloat16_to_float(__bfloat16_raw h);
49 
50 struct bfloat16_base : public __bfloat16_raw {
51  EIGEN_DEVICE_FUNC EIGEN_CONSTEXPR bfloat16_base() {}
52  EIGEN_DEVICE_FUNC EIGEN_CONSTEXPR bfloat16_base(const __bfloat16_raw& h) : __bfloat16_raw(h) {}
53 };
54 
55 } // namespace bfloat16_impl
56 
57 // Class definition.
58 struct bfloat16 : public bfloat16_impl::bfloat16_base {
59 
60  typedef bfloat16_impl::__bfloat16_raw __bfloat16_raw;
61 
62  EIGEN_DEVICE_FUNC EIGEN_CONSTEXPR bfloat16() {}
63 
64  EIGEN_DEVICE_FUNC EIGEN_CONSTEXPR bfloat16(const __bfloat16_raw& h) : bfloat16_impl::bfloat16_base(h) {}
65 
66  explicit EIGEN_DEVICE_FUNC EIGEN_CONSTEXPR bfloat16(bool b)
67  : bfloat16_impl::bfloat16_base(bfloat16_impl::raw_uint16_to_bfloat16(b ? 0x3f80 : 0)) {}
68 
69  template<class T>
70  explicit EIGEN_DEVICE_FUNC EIGEN_CONSTEXPR bfloat16(T val)
71  : bfloat16_impl::bfloat16_base(bfloat16_impl::float_to_bfloat16_rtne<internal::is_integral<T>::value>(static_cast<float>(val))) {}
72 
73  explicit EIGEN_DEVICE_FUNC bfloat16(float f)
74  : bfloat16_impl::bfloat16_base(bfloat16_impl::float_to_bfloat16_rtne<false>(f)) {}
75 
76  // Following the convention of numpy, converting between complex and
77  // float will lead to loss of imag value.
78  template<typename RealScalar>
79  explicit EIGEN_DEVICE_FUNC EIGEN_CONSTEXPR bfloat16(const std::complex<RealScalar>& val)
80  : bfloat16_impl::bfloat16_base(bfloat16_impl::float_to_bfloat16_rtne<false>(static_cast<float>(val.real()))) {}
81 
82  EIGEN_DEVICE_FUNC operator float() const { // NOLINT: Allow implicit conversion to float, because it is lossless.
83  return bfloat16_impl::bfloat16_to_float(*this);
84  }
85 };
86 } // namespace Eigen
87 
88 namespace std {
89 template<>
90 struct numeric_limits<Eigen::bfloat16> {
91  static const bool is_specialized = true;
92  static const bool is_signed = true;
93  static const bool is_integer = false;
94  static const bool is_exact = false;
95  static const bool has_infinity = true;
96  static const bool has_quiet_NaN = true;
97  static const bool has_signaling_NaN = true;
98  static const float_denorm_style has_denorm = std::denorm_absent;
99  static const bool has_denorm_loss = false;
100  static const std::float_round_style round_style = numeric_limits<float>::round_style;
101  static const bool is_iec559 = false;
102  static const bool is_bounded = true;
103  static const bool is_modulo = false;
104  static const int digits = 8;
105  static const int digits10 = 2;
106  static const int max_digits10 = 4;
107  static const int radix = 2;
108  static const int min_exponent = numeric_limits<float>::min_exponent;
109  static const int min_exponent10 = numeric_limits<float>::min_exponent10;
110  static const int max_exponent = numeric_limits<float>::max_exponent;
111  static const int max_exponent10 = numeric_limits<float>::max_exponent10;
112  static const bool traps = numeric_limits<float>::traps;
113  static const bool tinyness_before = numeric_limits<float>::tinyness_before;
114 
115  static Eigen::bfloat16 (min)() { return Eigen::bfloat16_impl::raw_uint16_to_bfloat16(0x0080); }
116  static Eigen::bfloat16 lowest() { return Eigen::bfloat16_impl::raw_uint16_to_bfloat16(0xff7f); }
117  static Eigen::bfloat16 (max)() { return Eigen::bfloat16_impl::raw_uint16_to_bfloat16(0x7f7f); }
118  static Eigen::bfloat16 epsilon() { return Eigen::bfloat16_impl::raw_uint16_to_bfloat16(0x3c00); }
119  static Eigen::bfloat16 round_error() { return Eigen::bfloat16(0x3f00); }
120  static Eigen::bfloat16 infinity() { return Eigen::bfloat16_impl::raw_uint16_to_bfloat16(0x7f80); }
121  static Eigen::bfloat16 quiet_NaN() { return Eigen::bfloat16_impl::raw_uint16_to_bfloat16(0x7fc0); }
122  static Eigen::bfloat16 signaling_NaN() { return Eigen::bfloat16_impl::raw_uint16_to_bfloat16(0x7f81); }
123  static Eigen::bfloat16 denorm_min() { return Eigen::bfloat16_impl::raw_uint16_to_bfloat16(0x0001); }
124 };
125 
126 // If std::numeric_limits<T> is specialized, should also specialize
127 // std::numeric_limits<const T>, std::numeric_limits<volatile T>, and
128 // std::numeric_limits<const volatile T>
129 // https://stackoverflow.com/a/16519653/
130 template<>
131 struct numeric_limits<const Eigen::bfloat16> : numeric_limits<Eigen::bfloat16> {};
132 template<>
133 struct numeric_limits<volatile Eigen::bfloat16> : numeric_limits<Eigen::bfloat16> {};
134 template<>
135 struct numeric_limits<const volatile Eigen::bfloat16> : numeric_limits<Eigen::bfloat16> {};
136 } // namespace std
137 
138 namespace Eigen {
139 
140 namespace bfloat16_impl {
141 
142 // We need to distinguish ‘clang as the CUDA compiler’ from ‘clang as the host compiler,
143 // invoked by NVCC’ (e.g. on MacOS). The former needs to see both host and device implementation
144 // of the functions, while the latter can only deal with one of them.
145 #if !defined(EIGEN_HAS_NATIVE_BF16) || (EIGEN_COMP_CLANG && !EIGEN_COMP_NVCC) // Emulate support for bfloat16 floats
146 
147 #if EIGEN_COMP_CLANG && defined(EIGEN_CUDACC)
148 // We need to provide emulated *host-side* BF16 operators for clang.
149 #pragma push_macro("EIGEN_DEVICE_FUNC")
150 #undef EIGEN_DEVICE_FUNC
151 #if defined(EIGEN_HAS_CUDA_BF16) && defined(EIGEN_HAS_NATIVE_BF16)
152 #define EIGEN_DEVICE_FUNC __host__
153 #else // both host and device need emulated ops.
154 #define EIGEN_DEVICE_FUNC __host__ __device__
155 #endif
156 #endif
157 
158 // Definitions for CPUs, mostly working through conversion
159 // to/from fp32.
160 
161 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC bfloat16 operator + (const bfloat16& a, const bfloat16& b) {
162  return bfloat16(float(a) + float(b));
163 }
164 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC bfloat16 operator + (const bfloat16& a, const int& b) {
165  return bfloat16(float(a) + static_cast<float>(b));
166 }
167 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC bfloat16 operator + (const int& a, const bfloat16& b) {
168  return bfloat16(static_cast<float>(a) + float(b));
169 }
170 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC bfloat16 operator * (const bfloat16& a, const bfloat16& b) {
171  return bfloat16(float(a) * float(b));
172 }
173 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC bfloat16 operator - (const bfloat16& a, const bfloat16& b) {
174  return bfloat16(float(a) - float(b));
175 }
176 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC bfloat16 operator / (const bfloat16& a, const bfloat16& b) {
177  return bfloat16(float(a) / float(b));
178 }
179 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC bfloat16 operator - (const bfloat16& a) {
180  bfloat16 result;
181  result.value = a.value ^ 0x8000;
182  return result;
183 }
184 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC bfloat16& operator += (bfloat16& a, const bfloat16& b) {
185  a = bfloat16(float(a) + float(b));
186  return a;
187 }
188 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC bfloat16& operator *= (bfloat16& a, const bfloat16& b) {
189  a = bfloat16(float(a) * float(b));
190  return a;
191 }
192 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC bfloat16& operator -= (bfloat16& a, const bfloat16& b) {
193  a = bfloat16(float(a) - float(b));
194  return a;
195 }
196 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC bfloat16& operator /= (bfloat16& a, const bfloat16& b) {
197  a = bfloat16(float(a) / float(b));
198  return a;
199 }
200 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC bfloat16 operator++(bfloat16& a) {
201  a += bfloat16(1);
202  return a;
203 }
204 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC bfloat16 operator--(bfloat16& a) {
205  a -= bfloat16(1);
206  return a;
207 }
208 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC bfloat16 operator++(bfloat16& a, int) {
209  bfloat16 original_value = a;
210  ++a;
211  return original_value;
212 }
213 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC bfloat16 operator--(bfloat16& a, int) {
214  bfloat16 original_value = a;
215  --a;
216  return original_value;
217 }
218 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC bool operator == (const bfloat16& a, const bfloat16& b) {
219  return numext::equal_strict(float(a),float(b));
220 }
221 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC bool operator != (const bfloat16& a, const bfloat16& b) {
222  return numext::not_equal_strict(float(a), float(b));
223 }
224 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC bool operator < (const bfloat16& a, const bfloat16& b) {
225  return float(a) < float(b);
226 }
227 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC bool operator <= (const bfloat16& a, const bfloat16& b) {
228  return float(a) <= float(b);
229 }
230 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC bool operator > (const bfloat16& a, const bfloat16& b) {
231  return float(a) > float(b);
232 }
233 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC bool operator >= (const bfloat16& a, const bfloat16& b) {
234  return float(a) >= float(b);
235 }
236 
237 #if EIGEN_COMP_CLANG && defined(EIGEN_CUDACC)
238 #pragma pop_macro("EIGEN_DEVICE_FUNC")
239 #endif
240 #endif // Emulate support for bfloat16 floats
241 
242 // Division by an index. Do it in full float precision to avoid accuracy
243 // issues in converting the denominator to bfloat16.
244 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC bfloat16 operator / (const bfloat16& a, Index b) {
245  return bfloat16(static_cast<float>(a) / static_cast<float>(b));
246 }
247 
248 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC __bfloat16_raw truncate_to_bfloat16(const float v) {
249  __bfloat16_raw output;
250  if (Eigen::numext::isnan EIGEN_NOT_A_MACRO(v)) {
251  output.value = std::signbit(v) ? 0xFFC0: 0x7FC0;
252  return output;
253  }
254  const uint16_t* p = reinterpret_cast<const uint16_t*>(&v);
255 #if defined(__BYTE_ORDER__) && __BYTE_ORDER__ == __ORDER_BIG_ENDIAN__
256  output.value = p[0];
257 #else
258  output.value = p[1];
259 #endif
260  return output;
261 }
262 
263 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC EIGEN_CONSTEXPR __bfloat16_raw raw_uint16_to_bfloat16(numext::uint16_t value) {
264  return __bfloat16_raw(value);
265 }
266 
267 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC EIGEN_CONSTEXPR numext::uint16_t raw_bfloat16_as_uint16(const __bfloat16_raw& bf) {
268  return bf.value;
269 }
270 
271 // float_to_bfloat16_rtne template specialization that does not make any
272 // assumption about the value of its function argument (ff).
273 template <>
274 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC __bfloat16_raw float_to_bfloat16_rtne<false>(float ff) {
275 #if (defined(EIGEN_HAS_CUDA_BF16) && defined(EIGEN_HAS_HIP_BF16))
276  // Nothing to do here
277 #else
278  __bfloat16_raw output;
279 
280  if (Eigen::numext::isnan EIGEN_NOT_A_MACRO(ff)) {
281  // If the value is a NaN, squash it to a qNaN with msb of fraction set,
282  // this makes sure after truncation we don't end up with an inf.
283  //
284  // qNaN magic: All exponent bits set + most significant bit of fraction
285  // set.
286  output.value = std::signbit(ff) ? 0xFFC0: 0x7FC0;
287  } else {
288  // Fast rounding algorithm that rounds a half value to nearest even. This
289  // reduces expected error when we convert a large number of floats. Here
290  // is how it works:
291  //
292  // Definitions:
293  // To convert a float 32 to bfloat16, a float 32 can be viewed as 32 bits
294  // with the following tags:
295  //
296  // Sign | Exp (8 bits) | Frac (23 bits)
297  // S EEEEEEEE FFFFFFLRTTTTTTTTTTTTTTT
298  //
299  // S: Sign bit.
300  // E: Exponent bits.
301  // F: First 6 bits of fraction.
302  // L: Least significant bit of resulting bfloat16 if we truncate away the
303  // rest of the float32. This is also the 7th bit of fraction
304  // R: Rounding bit, 8th bit of fraction.
305  // T: Sticky bits, rest of fraction, 15 bits.
306  //
307  // To round half to nearest even, there are 3 cases where we want to round
308  // down (simply truncate the result of the bits away, which consists of
309  // rounding bit and sticky bits) and two cases where we want to round up
310  // (truncate then add one to the result).
311  //
312  // The fast converting algorithm simply adds lsb (L) to 0x7fff (15 bits of
313  // 1s) as the rounding bias, adds the rounding bias to the input, then
314  // truncates the last 16 bits away.
315  //
316  // To understand how it works, we can analyze this algorithm case by case:
317  //
318  // 1. L = 0, R = 0:
319  // Expect: round down, this is less than half value.
320  //
321  // Algorithm:
322  // - Rounding bias: 0x7fff + 0 = 0x7fff
323  // - Adding rounding bias to input may create any carry, depending on
324  // whether there is any value set to 1 in T bits.
325  // - R may be set to 1 if there is a carry.
326  // - L remains 0.
327  // - Note that this case also handles Inf and -Inf, where all fraction
328  // bits, including L, R and Ts are all 0. The output remains Inf after
329  // this algorithm.
330  //
331  // 2. L = 1, R = 0:
332  // Expect: round down, this is less than half value.
333  //
334  // Algorithm:
335  // - Rounding bias: 0x7fff + 1 = 0x8000
336  // - Adding rounding bias to input doesn't change sticky bits but
337  // adds 1 to rounding bit.
338  // - L remains 1.
339  //
340  // 3. L = 0, R = 1, all of T are 0:
341  // Expect: round down, this is exactly at half, the result is already
342  // even (L=0).
343  //
344  // Algorithm:
345  // - Rounding bias: 0x7fff + 0 = 0x7fff
346  // - Adding rounding bias to input sets all sticky bits to 1, but
347  // doesn't create a carry.
348  // - R remains 1.
349  // - L remains 0.
350  //
351  // 4. L = 1, R = 1:
352  // Expect: round up, this is exactly at half, the result needs to be
353  // round to the next even number.
354  //
355  // Algorithm:
356  // - Rounding bias: 0x7fff + 1 = 0x8000
357  // - Adding rounding bias to input doesn't change sticky bits, but
358  // creates a carry from rounding bit.
359  // - The carry sets L to 0, creates another carry bit and propagate
360  // forward to F bits.
361  // - If all the F bits are 1, a carry then propagates to the exponent
362  // bits, which then creates the minimum value with the next exponent
363  // value. Note that we won't have the case where exponents are all 1,
364  // since that's either a NaN (handled in the other if condition) or inf
365  // (handled in case 1).
366  //
367  // 5. L = 0, R = 1, any of T is 1:
368  // Expect: round up, this is greater than half.
369  //
370  // Algorithm:
371  // - Rounding bias: 0x7fff + 0 = 0x7fff
372  // - Adding rounding bias to input creates a carry from sticky bits,
373  // sets rounding bit to 0, then create another carry.
374  // - The second carry sets L to 1.
375  //
376  // Examples:
377  //
378  // Exact half value that is already even:
379  // Input:
380  // Sign | Exp (8 bit) | Frac (first 7 bit) | Frac (last 16 bit)
381  // S E E E E E E E E F F F F F F L RTTTTTTTTTTTTTTT
382  // 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 0 1000000000000000
383  //
384  // This falls into case 3. We truncate the rest of 16 bits and no
385  // carry is created into F and L:
386  //
387  // Output:
388  // Sign | Exp (8 bit) | Frac (first 7 bit)
389  // S E E E E E E E E F F F F F F L
390  // 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 0
391  //
392  // Exact half value, round to next even number:
393  // Input:
394  // Sign | Exp (8 bit) | Frac (first 7 bit) | Frac (last 16 bit)
395  // S E E E E E E E E F F F F F F L RTTTTTTTTTTTTTTT
396  // 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1000000000000000
397  //
398  // This falls into case 4. We create a carry from R and T,
399  // which then propagates into L and F:
400  //
401  // Output:
402  // Sign | Exp (8 bit) | Frac (first 7 bit)
403  // S E E E E E E E E F F F F F F L
404  // 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 0
405  //
406  //
407  // Max denormal value round to min normal value:
408  // Input:
409  // Sign | Exp (8 bit) | Frac (first 7 bit) | Frac (last 16 bit)
410  // S E E E E E E E E F F F F F F L RTTTTTTTTTTTTTTT
411  // 0 0 0 0 0 0 0 0 0 1 1 1 1 1 1 1 1111111111111111
412  //
413  // This falls into case 4. We create a carry from R and T,
414  // propagate into L and F, which then propagates into exponent
415  // bits:
416  //
417  // Output:
418  // Sign | Exp (8 bit) | Frac (first 7 bit)
419  // S E E E E E E E E F F F F F F L
420  // 0 0 0 0 0 0 0 0 1 0 0 0 0 0 0 0
421  //
422  // Max normal value round to Inf:
423  // Input:
424  // Sign | Exp (8 bit) | Frac (first 7 bit) | Frac (last 16 bit)
425  // S E E E E E E E E F F F F F F L RTTTTTTTTTTTTTTT
426  // 0 1 1 1 1 1 1 1 0 1 1 1 1 1 1 1 1111111111111111
427  //
428  // This falls into case 4. We create a carry from R and T,
429  // propagate into L and F, which then propagates into exponent
430  // bits:
431  //
432  // Sign | Exp (8 bit) | Frac (first 7 bit)
433  // S E E E E E E E E F F F F F F L
434  // 0 1 1 1 1 1 1 1 1 0 0 0 0 0 0 0
435 
436  // At this point, ff must be either a normal float, or +/-infinity.
437  output = float_to_bfloat16_rtne<true>(ff);
438  }
439  return output;
440 #endif
441 }
442 
443 // float_to_bfloat16_rtne template specialization that assumes that its function
444 // argument (ff) is either a normal floating point number, or +/-infinity, or
445 // zero. Used to improve the runtime performance of conversion from an integer
446 // type to bfloat16.
447 template <>
448 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC __bfloat16_raw float_to_bfloat16_rtne<true>(float ff) {
449 #if (defined(EIGEN_HAS_CUDA_BF16) && defined(EIGEN_HAS_HIP_BF16))
450  // Nothing to do here
451 #else
452  numext::uint32_t input = numext::bit_cast<numext::uint32_t>(ff);
453  __bfloat16_raw output;
454 
455  // Least significant bit of resulting bfloat.
456  numext::uint32_t lsb = (input >> 16) & 1;
457  numext::uint32_t rounding_bias = 0x7fff + lsb;
458  input += rounding_bias;
459  output.value = static_cast<numext::uint16_t>(input >> 16);
460  return output;
461 #endif
462 }
463 
464 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC float bfloat16_to_float(__bfloat16_raw h) {
465  float result = 0;
466  unsigned short* q = reinterpret_cast<unsigned short*>(&result);
467 #if defined(__BYTE_ORDER__) && __BYTE_ORDER__ == __ORDER_BIG_ENDIAN__
468  q[0] = h.value;
469 #else
470  q[1] = h.value;
471 #endif
472  return result;
473 }
474 // --- standard functions ---
475 
476 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC bool (isinf)(const bfloat16& a) {
477  EIGEN_USING_STD(isinf);
478  return (isinf)(float(a));
479 }
480 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC bool (isnan)(const bfloat16& a) {
481  EIGEN_USING_STD(isnan);
482  return (isnan)(float(a));
483 }
484 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC bool (isfinite)(const bfloat16& a) {
485  return !(isinf EIGEN_NOT_A_MACRO (a)) && !(isnan EIGEN_NOT_A_MACRO (a));
486 }
487 
488 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC bfloat16 abs(const bfloat16& a) {
489  bfloat16 result;
490  result.value = a.value & 0x7FFF;
491  return result;
492 }
493 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC bfloat16 exp(const bfloat16& a) {
494  return bfloat16(::expf(float(a)));
495 }
496 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC bfloat16 expm1(const bfloat16& a) {
497  return bfloat16(numext::expm1(float(a)));
498 }
499 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC bfloat16 log(const bfloat16& a) {
500  return bfloat16(::logf(float(a)));
501 }
502 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC bfloat16 log1p(const bfloat16& a) {
503  return bfloat16(numext::log1p(float(a)));
504 }
505 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC bfloat16 log10(const bfloat16& a) {
506  return bfloat16(::log10f(float(a)));
507 }
508 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC bfloat16 log2(const bfloat16& a) {
509  return bfloat16(static_cast<float>(EIGEN_LOG2E) * ::logf(float(a)));
510 }
511 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC bfloat16 sqrt(const bfloat16& a) {
512  return bfloat16(::sqrtf(float(a)));
513 }
514 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC bfloat16 pow(const bfloat16& a, const bfloat16& b) {
515  return bfloat16(::powf(float(a), float(b)));
516 }
517 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC bfloat16 sin(const bfloat16& a) {
518  return bfloat16(::sinf(float(a)));
519 }
520 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC bfloat16 cos(const bfloat16& a) {
521  return bfloat16(::cosf(float(a)));
522 }
523 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC bfloat16 tan(const bfloat16& a) {
524  return bfloat16(::tanf(float(a)));
525 }
526 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC bfloat16 asin(const bfloat16& a) {
527  return bfloat16(::asinf(float(a)));
528 }
529 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC bfloat16 acos(const bfloat16& a) {
530  return bfloat16(::acosf(float(a)));
531 }
532 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC bfloat16 atan(const bfloat16& a) {
533  return bfloat16(::atanf(float(a)));
534 }
535 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC bfloat16 sinh(const bfloat16& a) {
536  return bfloat16(::sinhf(float(a)));
537 }
538 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC bfloat16 cosh(const bfloat16& a) {
539  return bfloat16(::coshf(float(a)));
540 }
541 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC bfloat16 tanh(const bfloat16& a) {
542  return bfloat16(::tanhf(float(a)));
543 }
544 #if EIGEN_HAS_CXX11_MATH
545 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC bfloat16 asinh(const bfloat16& a) {
546  return bfloat16(::asinhf(float(a)));
547 }
548 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC bfloat16 acosh(const bfloat16& a) {
549  return bfloat16(::acoshf(float(a)));
550 }
551 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC bfloat16 atanh(const bfloat16& a) {
552  return bfloat16(::atanhf(float(a)));
553 }
554 #endif
555 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC bfloat16 floor(const bfloat16& a) {
556  return bfloat16(::floorf(float(a)));
557 }
558 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC bfloat16 ceil(const bfloat16& a) {
559  return bfloat16(::ceilf(float(a)));
560 }
561 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC bfloat16 rint(const bfloat16& a) {
562  return bfloat16(::rintf(float(a)));
563 }
564 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC bfloat16 round(const bfloat16& a) {
565  return bfloat16(::roundf(float(a)));
566 }
567 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC bfloat16 fmod(const bfloat16& a, const bfloat16& b) {
568  return bfloat16(::fmodf(float(a), float(b)));
569 }
570 
571 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC bfloat16 (min)(const bfloat16& a, const bfloat16& b) {
572  const float f1 = static_cast<float>(a);
573  const float f2 = static_cast<float>(b);
574  return f2 < f1 ? b : a;
575 }
576 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC bfloat16 (max)(const bfloat16& a, const bfloat16& b) {
577  const float f1 = static_cast<float>(a);
578  const float f2 = static_cast<float>(b);
579  return f1 < f2 ? b : a;
580 }
581 
582 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC bfloat16 fmin(const bfloat16& a, const bfloat16& b) {
583  const float f1 = static_cast<float>(a);
584  const float f2 = static_cast<float>(b);
585  return bfloat16(::fminf(f1, f2));
586 }
587 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC bfloat16 fmax(const bfloat16& a, const bfloat16& b) {
588  const float f1 = static_cast<float>(a);
589  const float f2 = static_cast<float>(b);
590  return bfloat16(::fmaxf(f1, f2));
591 }
592 
593 #ifndef EIGEN_NO_IO
594 EIGEN_ALWAYS_INLINE std::ostream& operator << (std::ostream& os, const bfloat16& v) {
595  os << static_cast<float>(v);
596  return os;
597 }
598 #endif
599 
600 } // namespace bfloat16_impl
601 
602 namespace internal {
603 
604 template<>
605 struct random_default_impl<bfloat16, false, false>
606 {
607  static inline bfloat16 run(const bfloat16& x, const bfloat16& y)
608  {
609  return x + (y-x) * bfloat16(float(std::rand()) / float(RAND_MAX));
610  }
611  static inline bfloat16 run()
612  {
613  return run(bfloat16(-1.f), bfloat16(1.f));
614  }
615 };
616 
617 template<> struct is_arithmetic<bfloat16> { enum { value = true }; };
618 
619 } // namespace internal
620 
621 template<> struct NumTraits<Eigen::bfloat16>
622  : GenericNumTraits<Eigen::bfloat16>
623 {
624  enum {
625  IsSigned = true,
626  IsInteger = false,
627  IsComplex = false,
628  RequireInitialization = false
629  };
630 
631  EIGEN_DEVICE_FUNC EIGEN_CONSTEXPR static EIGEN_STRONG_INLINE Eigen::bfloat16 epsilon() {
632  return bfloat16_impl::raw_uint16_to_bfloat16(0x3c00);
633  }
634  EIGEN_DEVICE_FUNC EIGEN_CONSTEXPR static EIGEN_STRONG_INLINE Eigen::bfloat16 dummy_precision() {
635  return bfloat16_impl::raw_uint16_to_bfloat16(0x3D4D); // bfloat16(5e-2f);
636 
637  }
638  EIGEN_DEVICE_FUNC EIGEN_CONSTEXPR static EIGEN_STRONG_INLINE Eigen::bfloat16 highest() {
639  return bfloat16_impl::raw_uint16_to_bfloat16(0x7F7F);
640  }
641  EIGEN_DEVICE_FUNC EIGEN_CONSTEXPR static EIGEN_STRONG_INLINE Eigen::bfloat16 lowest() {
642  return bfloat16_impl::raw_uint16_to_bfloat16(0xFF7F);
643  }
644  EIGEN_DEVICE_FUNC EIGEN_CONSTEXPR static EIGEN_STRONG_INLINE Eigen::bfloat16 infinity() {
645  return bfloat16_impl::raw_uint16_to_bfloat16(0x7f80);
646  }
647  EIGEN_DEVICE_FUNC EIGEN_CONSTEXPR static EIGEN_STRONG_INLINE Eigen::bfloat16 quiet_NaN() {
648  return bfloat16_impl::raw_uint16_to_bfloat16(0x7fc0);
649  }
650 };
651 
652 } // namespace Eigen
653 
654 namespace Eigen {
655 namespace numext {
656 
657 template<>
658 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
659 bool (isnan)(const Eigen::bfloat16& h) {
660  return (bfloat16_impl::isnan)(h);
661 }
662 
663 template<>
664 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
665 bool (isinf)(const Eigen::bfloat16& h) {
666  return (bfloat16_impl::isinf)(h);
667 }
668 
669 template<>
670 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
671 bool (isfinite)(const Eigen::bfloat16& h) {
672  return (bfloat16_impl::isfinite)(h);
673 }
674 
675 template <>
676 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC Eigen::bfloat16 bit_cast<Eigen::bfloat16, uint16_t>(const uint16_t& src) {
677  return Eigen::bfloat16(Eigen::bfloat16_impl::raw_uint16_to_bfloat16(src));
678 }
679 
680 template <>
681 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC uint16_t bit_cast<uint16_t, Eigen::bfloat16>(const Eigen::bfloat16& src) {
682  return Eigen::bfloat16_impl::raw_bfloat16_as_uint16(src);
683 }
684 
685 } // namespace numext
686 } // namespace Eigen
687 
688 #if EIGEN_HAS_STD_HASH
689 namespace std {
690 template <>
691 struct hash<Eigen::bfloat16> {
692  EIGEN_STRONG_INLINE std::size_t operator()(const Eigen::bfloat16& a) const {
693  return static_cast<std::size_t>(Eigen::numext::bit_cast<Eigen::numext::uint16_t>(a));
694  }
695 };
696 } // namespace std
697 #endif
698 
699 
700 #endif // EIGEN_BFLOAT16_H
const Eigen::CwiseUnaryOp< Eigen::internal::scalar_tanh_op< typename Derived::Scalar >, const Derived > tanh(const Eigen::ArrayBase< Derived > &x)
const Eigen::CwiseUnaryOp< Eigen::internal::scalar_sinh_op< typename Derived::Scalar >, const Derived > sinh(const Eigen::ArrayBase< Derived > &x)
const Eigen::CwiseUnaryOp< Eigen::internal::scalar_isfinite_op< typename Derived::Scalar >, const Derived > isfinite(const Eigen::ArrayBase< Derived > &x)
const Eigen::CwiseUnaryOp< Eigen::internal::scalar_sqrt_op< typename Derived::Scalar >, const Derived > sqrt(const Eigen::ArrayBase< Derived > &x)
Namespace containing all symbols from the Eigen library.
Definition: Core:141
Definition: BFloat16.h:88
const Eigen::CwiseUnaryOp< Eigen::internal::scalar_ceil_op< typename Derived::Scalar >, const Derived > ceil(const Eigen::ArrayBase< Derived > &x)
const Eigen::CwiseUnaryOp< Eigen::internal::scalar_asin_op< typename Derived::Scalar >, const Derived > asin(const Eigen::ArrayBase< Derived > &x)
const Eigen::CwiseUnaryOp< Eigen::internal::scalar_acos_op< typename Derived::Scalar >, const Derived > acos(const Eigen::ArrayBase< Derived > &x)
const Eigen::CwiseUnaryOp< Eigen::internal::scalar_isnan_op< typename Derived::Scalar >, const Derived > isnan(const Eigen::ArrayBase< Derived > &x)
const Eigen::CwiseUnaryOp< Eigen::internal::scalar_cos_op< typename Derived::Scalar >, const Derived > cos(const Eigen::ArrayBase< Derived > &x)
const Product< MatrixDerived, PermutationDerived, AliasFreeProduct > operator*(const MatrixBase< MatrixDerived > &matrix, const PermutationBase< PermutationDerived > &permutation)
Definition: PermutationMatrix.h:515
EIGEN_DEFAULT_DENSE_INDEX_TYPE Index
The Index type as used for the API.
Definition: Meta.h:74
const Eigen::CwiseUnaryOp< Eigen::internal::scalar_round_op< typename Derived::Scalar >, const Derived > round(const Eigen::ArrayBase< Derived > &x)
const Eigen::CwiseUnaryOp< Eigen::internal::scalar_rint_op< typename Derived::Scalar >, const Derived > rint(const Eigen::ArrayBase< Derived > &x)
const Eigen::CwiseUnaryOp< Eigen::internal::scalar_floor_op< typename Derived::Scalar >, const Derived > floor(const Eigen::ArrayBase< Derived > &x)
const Eigen::CwiseUnaryOp< Eigen::internal::scalar_log1p_op< typename Derived::Scalar >, const Derived > log1p(const Eigen::ArrayBase< Derived > &x)
const Eigen::CwiseUnaryOp< Eigen::internal::scalar_isinf_op< typename Derived::Scalar >, const Derived > isinf(const Eigen::ArrayBase< Derived > &x)
const Eigen::CwiseUnaryOp< Eigen::internal::scalar_real_op< typename Derived::Scalar >, const Derived > real(const Eigen::ArrayBase< Derived > &x)
const Eigen::CwiseUnaryOp< Eigen::internal::scalar_abs_op< typename Derived::Scalar >, const Derived > abs(const Eigen::ArrayBase< Derived > &x)
const Eigen::CwiseUnaryOp< Eigen::internal::scalar_cosh_op< typename Derived::Scalar >, const Derived > cosh(const Eigen::ArrayBase< Derived > &x)
Definition: Eigen_Colamd.h:50
const Eigen::CwiseUnaryOp< Eigen::internal::scalar_log_op< typename Derived::Scalar >, const Derived > log(const Eigen::ArrayBase< Derived > &x)
const Eigen::CwiseUnaryOp< Eigen::internal::scalar_tan_op< typename Derived::Scalar >, const Derived > tan(const Eigen::ArrayBase< Derived > &x)
const Eigen::CwiseUnaryOp< Eigen::internal::scalar_expm1_op< typename Derived::Scalar >, const Derived > expm1(const Eigen::ArrayBase< Derived > &x)
const Eigen::CwiseUnaryOp< Eigen::internal::scalar_atanh_op< typename Derived::Scalar >, const Derived > atanh(const Eigen::ArrayBase< Derived > &x)
const Eigen::CwiseUnaryOp< Eigen::internal::scalar_log2_op< typename Derived::Scalar >, const Derived > log2(const Eigen::ArrayBase< Derived > &x)
const Eigen::CwiseUnaryOp< Eigen::internal::scalar_atan_op< typename Derived::Scalar >, const Derived > atan(const Eigen::ArrayBase< Derived > &x)
const Eigen::CwiseUnaryOp< Eigen::internal::scalar_sin_op< typename Derived::Scalar >, const Derived > sin(const Eigen::ArrayBase< Derived > &x)
const Eigen::CwiseUnaryOp< Eigen::internal::scalar_exp_op< typename Derived::Scalar >, const Derived > exp(const Eigen::ArrayBase< Derived > &x)
const Eigen::CwiseUnaryOp< Eigen::internal::scalar_log10_op< typename Derived::Scalar >, const Derived > log10(const Eigen::ArrayBase< Derived > &x)
const Eigen::CwiseUnaryOp< Eigen::internal::scalar_acosh_op< typename Derived::Scalar >, const Derived > acosh(const Eigen::ArrayBase< Derived > &x)
const Eigen::CwiseUnaryOp< Eigen::internal::scalar_asinh_op< typename Derived::Scalar >, const Derived > asinh(const Eigen::ArrayBase< Derived > &x)