Eigen  3.3.5
Half.h
1 // This file is part of Eigen, a lightweight C++ template library
2 // for linear algebra.
3 //
4 // This Source Code Form is subject to the terms of the Mozilla
5 // Public License v. 2.0. If a copy of the MPL was not distributed
6 // with this file, You can obtain one at http://mozilla.org/MPL/2.0/.
7 //
8 // The conversion routines are Copyright (c) Fabian Giesen, 2016.
9 // The original license follows:
10 //
11 // Copyright (c) Fabian Giesen, 2016
12 // All rights reserved.
13 // Redistribution and use in source and binary forms, with or without
14 // modification, are permitted.
15 // THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS
16 // "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT
17 // LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR
18 // A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT
19 // HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL,
20 // SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT
21 // LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE,
22 // DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY
23 // THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
24 // (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
25 // OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
26 
27 
28 // Standard 16-bit float type, mostly useful for GPUs. Defines a new
29 // type Eigen::half (inheriting from CUDA's __half struct) with
30 // operator overloads such that it behaves basically as an arithmetic
31 // type. It will be quite slow on CPUs (so it is recommended to stay
32 // in fp32 for CPUs, except for simple parameter conversions, I/O
33 // to disk and the likes), but fast on GPUs.
34 
35 
36 #ifndef EIGEN_HALF_CUDA_H
37 #define EIGEN_HALF_CUDA_H
38 
39 #if __cplusplus > 199711L
40 #define EIGEN_EXPLICIT_CAST(tgt_type) explicit operator tgt_type()
41 #else
42 #define EIGEN_EXPLICIT_CAST(tgt_type) operator tgt_type()
43 #endif
44 
45 
46 namespace Eigen {
47 
48 struct half;
49 
50 namespace half_impl {
51 
52 #if !defined(EIGEN_HAS_CUDA_FP16)
53 
54 // Make our own __half definition that is similar to CUDA's.
55 struct __half {
56  EIGEN_DEVICE_FUNC __half() {}
57  explicit EIGEN_DEVICE_FUNC __half(unsigned short raw) : x(raw) {}
58  unsigned short x;
59 };
60 
61 #endif
62 
63 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC __half raw_uint16_to_half(unsigned short x);
64 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC __half float_to_half_rtne(float ff);
65 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC float half_to_float(__half h);
66 
67 struct half_base : public __half {
68  EIGEN_DEVICE_FUNC half_base() {}
69  EIGEN_DEVICE_FUNC half_base(const half_base& h) : __half(h) {}
70  EIGEN_DEVICE_FUNC half_base(const __half& h) : __half(h) {}
71 };
72 
73 } // namespace half_impl
74 
75 // Class definition.
76 struct half : public half_impl::half_base {
77  #if !defined(EIGEN_HAS_CUDA_FP16)
78  typedef half_impl::__half __half;
79  #endif
80 
81  EIGEN_DEVICE_FUNC half() {}
82 
83  EIGEN_DEVICE_FUNC half(const __half& h) : half_impl::half_base(h) {}
84  EIGEN_DEVICE_FUNC half(const half& h) : half_impl::half_base(h) {}
85 
86  explicit EIGEN_DEVICE_FUNC half(bool b)
87  : half_impl::half_base(half_impl::raw_uint16_to_half(b ? 0x3c00 : 0)) {}
88  template<class T>
89  explicit EIGEN_DEVICE_FUNC half(const T& val)
90  : half_impl::half_base(half_impl::float_to_half_rtne(static_cast<float>(val))) {}
91  explicit EIGEN_DEVICE_FUNC half(float f)
92  : half_impl::half_base(half_impl::float_to_half_rtne(f)) {}
93 
94  EIGEN_DEVICE_FUNC EIGEN_EXPLICIT_CAST(bool) const {
95  // +0.0 and -0.0 become false, everything else becomes true.
96  return (x & 0x7fff) != 0;
97  }
98  EIGEN_DEVICE_FUNC EIGEN_EXPLICIT_CAST(signed char) const {
99  return static_cast<signed char>(half_impl::half_to_float(*this));
100  }
101  EIGEN_DEVICE_FUNC EIGEN_EXPLICIT_CAST(unsigned char) const {
102  return static_cast<unsigned char>(half_impl::half_to_float(*this));
103  }
104  EIGEN_DEVICE_FUNC EIGEN_EXPLICIT_CAST(short) const {
105  return static_cast<short>(half_impl::half_to_float(*this));
106  }
107  EIGEN_DEVICE_FUNC EIGEN_EXPLICIT_CAST(unsigned short) const {
108  return static_cast<unsigned short>(half_impl::half_to_float(*this));
109  }
110  EIGEN_DEVICE_FUNC EIGEN_EXPLICIT_CAST(int) const {
111  return static_cast<int>(half_impl::half_to_float(*this));
112  }
113  EIGEN_DEVICE_FUNC EIGEN_EXPLICIT_CAST(unsigned int) const {
114  return static_cast<unsigned int>(half_impl::half_to_float(*this));
115  }
116  EIGEN_DEVICE_FUNC EIGEN_EXPLICIT_CAST(long) const {
117  return static_cast<long>(half_impl::half_to_float(*this));
118  }
119  EIGEN_DEVICE_FUNC EIGEN_EXPLICIT_CAST(unsigned long) const {
120  return static_cast<unsigned long>(half_impl::half_to_float(*this));
121  }
122  EIGEN_DEVICE_FUNC EIGEN_EXPLICIT_CAST(long long) const {
123  return static_cast<long long>(half_impl::half_to_float(*this));
124  }
125  EIGEN_DEVICE_FUNC EIGEN_EXPLICIT_CAST(unsigned long long) const {
126  return static_cast<unsigned long long>(half_to_float(*this));
127  }
128  EIGEN_DEVICE_FUNC EIGEN_EXPLICIT_CAST(float) const {
129  return half_impl::half_to_float(*this);
130  }
131  EIGEN_DEVICE_FUNC EIGEN_EXPLICIT_CAST(double) const {
132  return static_cast<double>(half_impl::half_to_float(*this));
133  }
134 
135  EIGEN_DEVICE_FUNC half& operator=(const half& other) {
136  x = other.x;
137  return *this;
138  }
139 };
140 
141 namespace half_impl {
142 
143 #if defined(EIGEN_HAS_CUDA_FP16) && defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 530
144 
145 // Intrinsics for native fp16 support. Note that on current hardware,
146 // these are no faster than fp32 arithmetic (you need to use the half2
147 // versions to get the ALU speed increased), but you do save the
148 // conversion steps back and forth.
149 
150 EIGEN_STRONG_INLINE __device__ half operator + (const half& a, const half& b) {
151  return __hadd(a, b);
152 }
153 EIGEN_STRONG_INLINE __device__ half operator * (const half& a, const half& b) {
154  return __hmul(a, b);
155 }
156 EIGEN_STRONG_INLINE __device__ half operator - (const half& a, const half& b) {
157  return __hsub(a, b);
158 }
159 EIGEN_STRONG_INLINE __device__ half operator / (const half& a, const half& b) {
160  float num = __half2float(a);
161  float denom = __half2float(b);
162  return __float2half(num / denom);
163 }
164 EIGEN_STRONG_INLINE __device__ half operator - (const half& a) {
165  return __hneg(a);
166 }
167 EIGEN_STRONG_INLINE __device__ half& operator += (half& a, const half& b) {
168  a = a + b;
169  return a;
170 }
171 EIGEN_STRONG_INLINE __device__ half& operator *= (half& a, const half& b) {
172  a = a * b;
173  return a;
174 }
175 EIGEN_STRONG_INLINE __device__ half& operator -= (half& a, const half& b) {
176  a = a - b;
177  return a;
178 }
179 EIGEN_STRONG_INLINE __device__ half& operator /= (half& a, const half& b) {
180  a = a / b;
181  return a;
182 }
183 EIGEN_STRONG_INLINE __device__ bool operator == (const half& a, const half& b) {
184  return __heq(a, b);
185 }
186 EIGEN_STRONG_INLINE __device__ bool operator != (const half& a, const half& b) {
187  return __hne(a, b);
188 }
189 EIGEN_STRONG_INLINE __device__ bool operator < (const half& a, const half& b) {
190  return __hlt(a, b);
191 }
192 EIGEN_STRONG_INLINE __device__ bool operator <= (const half& a, const half& b) {
193  return __hle(a, b);
194 }
195 EIGEN_STRONG_INLINE __device__ bool operator > (const half& a, const half& b) {
196  return __hgt(a, b);
197 }
198 EIGEN_STRONG_INLINE __device__ bool operator >= (const half& a, const half& b) {
199  return __hge(a, b);
200 }
201 
202 #else // Emulate support for half floats
203 
204 // Definitions for CPUs and older CUDA, mostly working through conversion
205 // to/from fp32.
206 
207 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half operator + (const half& a, const half& b) {
208  return half(float(a) + float(b));
209 }
210 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half operator * (const half& a, const half& b) {
211  return half(float(a) * float(b));
212 }
213 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half operator - (const half& a, const half& b) {
214  return half(float(a) - float(b));
215 }
216 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half operator / (const half& a, const half& b) {
217  return half(float(a) / float(b));
218 }
219 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half operator - (const half& a) {
220  half result;
221  result.x = a.x ^ 0x8000;
222  return result;
223 }
224 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half& operator += (half& a, const half& b) {
225  a = half(float(a) + float(b));
226  return a;
227 }
228 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half& operator *= (half& a, const half& b) {
229  a = half(float(a) * float(b));
230  return a;
231 }
232 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half& operator -= (half& a, const half& b) {
233  a = half(float(a) - float(b));
234  return a;
235 }
236 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half& operator /= (half& a, const half& b) {
237  a = half(float(a) / float(b));
238  return a;
239 }
240 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC bool operator == (const half& a, const half& b) {
241  return numext::equal_strict(float(a),float(b));
242 }
243 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC bool operator != (const half& a, const half& b) {
244  return numext::not_equal_strict(float(a), float(b));
245 }
246 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC bool operator < (const half& a, const half& b) {
247  return float(a) < float(b);
248 }
249 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC bool operator <= (const half& a, const half& b) {
250  return float(a) <= float(b);
251 }
252 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC bool operator > (const half& a, const half& b) {
253  return float(a) > float(b);
254 }
255 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC bool operator >= (const half& a, const half& b) {
256  return float(a) >= float(b);
257 }
258 
259 #endif // Emulate support for half floats
260 
261 // Division by an index. Do it in full float precision to avoid accuracy
262 // issues in converting the denominator to half.
263 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half operator / (const half& a, Index b) {
264  return half(static_cast<float>(a) / static_cast<float>(b));
265 }
266 
267 // Conversion routines, including fallbacks for the host or older CUDA.
268 // Note that newer Intel CPUs (Haswell or newer) have vectorized versions of
269 // these in hardware. If we need more performance on older/other CPUs, they are
270 // also possible to vectorize directly.
271 
272 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC __half raw_uint16_to_half(unsigned short x) {
273  __half h;
274  h.x = x;
275  return h;
276 }
277 
278 union FP32 {
279  unsigned int u;
280  float f;
281 };
282 
283 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC __half float_to_half_rtne(float ff) {
284 #if defined(EIGEN_HAS_CUDA_FP16) && defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 300
285  return __float2half(ff);
286 
287 #elif defined(EIGEN_HAS_FP16_C)
288  __half h;
289  h.x = _cvtss_sh(ff, 0);
290  return h;
291 
292 #else
293  FP32 f; f.f = ff;
294 
295  const FP32 f32infty = { 255 << 23 };
296  const FP32 f16max = { (127 + 16) << 23 };
297  const FP32 denorm_magic = { ((127 - 15) + (23 - 10) + 1) << 23 };
298  unsigned int sign_mask = 0x80000000u;
299  __half o;
300  o.x = static_cast<unsigned short>(0x0u);
301 
302  unsigned int sign = f.u & sign_mask;
303  f.u ^= sign;
304 
305  // NOTE all the integer compares in this function can be safely
306  // compiled into signed compares since all operands are below
307  // 0x80000000. Important if you want fast straight SSE2 code
308  // (since there's no unsigned PCMPGTD).
309 
310  if (f.u >= f16max.u) { // result is Inf or NaN (all exponent bits set)
311  o.x = (f.u > f32infty.u) ? 0x7e00 : 0x7c00; // NaN->qNaN and Inf->Inf
312  } else { // (De)normalized number or zero
313  if (f.u < (113 << 23)) { // resulting FP16 is subnormal or zero
314  // use a magic value to align our 10 mantissa bits at the bottom of
315  // the float. as long as FP addition is round-to-nearest-even this
316  // just works.
317  f.f += denorm_magic.f;
318 
319  // and one integer subtract of the bias later, we have our final float!
320  o.x = static_cast<unsigned short>(f.u - denorm_magic.u);
321  } else {
322  unsigned int mant_odd = (f.u >> 13) & 1; // resulting mantissa is odd
323 
324  // update exponent, rounding bias part 1
325  f.u += ((unsigned int)(15 - 127) << 23) + 0xfff;
326  // rounding bias part 2
327  f.u += mant_odd;
328  // take the bits!
329  o.x = static_cast<unsigned short>(f.u >> 13);
330  }
331  }
332 
333  o.x |= static_cast<unsigned short>(sign >> 16);
334  return o;
335 #endif
336 }
337 
338 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC float half_to_float(__half h) {
339 #if defined(EIGEN_HAS_CUDA_FP16) && defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 300
340  return __half2float(h);
341 
342 #elif defined(EIGEN_HAS_FP16_C)
343  return _cvtsh_ss(h.x);
344 
345 #else
346  const FP32 magic = { 113 << 23 };
347  const unsigned int shifted_exp = 0x7c00 << 13; // exponent mask after shift
348  FP32 o;
349 
350  o.u = (h.x & 0x7fff) << 13; // exponent/mantissa bits
351  unsigned int exp = shifted_exp & o.u; // just the exponent
352  o.u += (127 - 15) << 23; // exponent adjust
353 
354  // handle exponent special cases
355  if (exp == shifted_exp) { // Inf/NaN?
356  o.u += (128 - 16) << 23; // extra exp adjust
357  } else if (exp == 0) { // Zero/Denormal?
358  o.u += 1 << 23; // extra exp adjust
359  o.f -= magic.f; // renormalize
360  }
361 
362  o.u |= (h.x & 0x8000) << 16; // sign bit
363  return o.f;
364 #endif
365 }
366 
367 // --- standard functions ---
368 
369 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC bool (isinf)(const half& a) {
370  return (a.x & 0x7fff) == 0x7c00;
371 }
372 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC bool (isnan)(const half& a) {
373 #if defined(EIGEN_HAS_CUDA_FP16) && defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 530
374  return __hisnan(a);
375 #else
376  return (a.x & 0x7fff) > 0x7c00;
377 #endif
378 }
379 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC bool (isfinite)(const half& a) {
380  return !(isinf EIGEN_NOT_A_MACRO (a)) && !(isnan EIGEN_NOT_A_MACRO (a));
381 }
382 
383 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half abs(const half& a) {
384  half result;
385  result.x = a.x & 0x7FFF;
386  return result;
387 }
388 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half exp(const half& a) {
389 #if EIGEN_CUDACC_VER >= 80000 && defined EIGEN_CUDA_ARCH && EIGEN_CUDA_ARCH >= 530
390  return half(hexp(a));
391 #else
392  return half(::expf(float(a)));
393 #endif
394 }
395 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half log(const half& a) {
396 #if defined(EIGEN_HAS_CUDA_FP16) && EIGEN_CUDACC_VER >= 80000 && defined(EIGEN_CUDA_ARCH) && EIGEN_CUDA_ARCH >= 530
397  return half(::hlog(a));
398 #else
399  return half(::logf(float(a)));
400 #endif
401 }
402 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half log1p(const half& a) {
403  return half(numext::log1p(float(a)));
404 }
405 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half log10(const half& a) {
406  return half(::log10f(float(a)));
407 }
408 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half sqrt(const half& a) {
409 #if EIGEN_CUDACC_VER >= 80000 && defined EIGEN_CUDA_ARCH && EIGEN_CUDA_ARCH >= 530
410  return half(hsqrt(a));
411 #else
412  return half(::sqrtf(float(a)));
413 #endif
414 }
415 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half pow(const half& a, const half& b) {
416  return half(::powf(float(a), float(b)));
417 }
418 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half sin(const half& a) {
419  return half(::sinf(float(a)));
420 }
421 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half cos(const half& a) {
422  return half(::cosf(float(a)));
423 }
424 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half tan(const half& a) {
425  return half(::tanf(float(a)));
426 }
427 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half tanh(const half& a) {
428  return half(::tanhf(float(a)));
429 }
430 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half floor(const half& a) {
431 #if EIGEN_CUDACC_VER >= 80000 && defined EIGEN_CUDA_ARCH && EIGEN_CUDA_ARCH >= 300
432  return half(hfloor(a));
433 #else
434  return half(::floorf(float(a)));
435 #endif
436 }
437 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half ceil(const half& a) {
438 #if EIGEN_CUDACC_VER >= 80000 && defined EIGEN_CUDA_ARCH && EIGEN_CUDA_ARCH >= 300
439  return half(hceil(a));
440 #else
441  return half(::ceilf(float(a)));
442 #endif
443 }
444 
445 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half (min)(const half& a, const half& b) {
446 #if defined(EIGEN_HAS_CUDA_FP16) && defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 530
447  return __hlt(b, a) ? b : a;
448 #else
449  const float f1 = static_cast<float>(a);
450  const float f2 = static_cast<float>(b);
451  return f2 < f1 ? b : a;
452 #endif
453 }
454 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half (max)(const half& a, const half& b) {
455 #if defined(EIGEN_HAS_CUDA_FP16) && defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 530
456  return __hlt(a, b) ? b : a;
457 #else
458  const float f1 = static_cast<float>(a);
459  const float f2 = static_cast<float>(b);
460  return f1 < f2 ? b : a;
461 #endif
462 }
463 
464 EIGEN_ALWAYS_INLINE std::ostream& operator << (std::ostream& os, const half& v) {
465  os << static_cast<float>(v);
466  return os;
467 }
468 
469 } // end namespace half_impl
470 
471 // import Eigen::half_impl::half into Eigen namespace
472 // using half_impl::half;
473 
474 namespace internal {
475 
476 template<>
477 struct random_default_impl<half, false, false>
478 {
479  static inline half run(const half& x, const half& y)
480  {
481  return x + (y-x) * half(float(std::rand()) / float(RAND_MAX));
482  }
483  static inline half run()
484  {
485  return run(half(-1.f), half(1.f));
486  }
487 };
488 
489 template<> struct is_arithmetic<half> { enum { value = true }; };
490 
491 } // end namespace internal
492 
493 } // end namespace Eigen
494 
495 namespace std {
496 template<>
497 struct numeric_limits<Eigen::half> {
498  static const bool is_specialized = true;
499  static const bool is_signed = true;
500  static const bool is_integer = false;
501  static const bool is_exact = false;
502  static const bool has_infinity = true;
503  static const bool has_quiet_NaN = true;
504  static const bool has_signaling_NaN = true;
505  static const float_denorm_style has_denorm = denorm_present;
506  static const bool has_denorm_loss = false;
507  static const std::float_round_style round_style = std::round_to_nearest;
508  static const bool is_iec559 = false;
509  static const bool is_bounded = false;
510  static const bool is_modulo = false;
511  static const int digits = 11;
512  static const int digits10 = 3; // according to http://half.sourceforge.net/structstd_1_1numeric__limits_3_01half__float_1_1half_01_4.html
513  static const int max_digits10 = 5; // according to http://half.sourceforge.net/structstd_1_1numeric__limits_3_01half__float_1_1half_01_4.html
514  static const int radix = 2;
515  static const int min_exponent = -13;
516  static const int min_exponent10 = -4;
517  static const int max_exponent = 16;
518  static const int max_exponent10 = 4;
519  static const bool traps = true;
520  static const bool tinyness_before = false;
521 
522  static Eigen::half (min)() { return Eigen::half_impl::raw_uint16_to_half(0x400); }
523  static Eigen::half lowest() { return Eigen::half_impl::raw_uint16_to_half(0xfbff); }
524  static Eigen::half (max)() { return Eigen::half_impl::raw_uint16_to_half(0x7bff); }
525  static Eigen::half epsilon() { return Eigen::half_impl::raw_uint16_to_half(0x0800); }
526  static Eigen::half round_error() { return Eigen::half(0.5); }
527  static Eigen::half infinity() { return Eigen::half_impl::raw_uint16_to_half(0x7c00); }
528  static Eigen::half quiet_NaN() { return Eigen::half_impl::raw_uint16_to_half(0x7e00); }
529  static Eigen::half signaling_NaN() { return Eigen::half_impl::raw_uint16_to_half(0x7e00); }
530  static Eigen::half denorm_min() { return Eigen::half_impl::raw_uint16_to_half(0x1); }
531 };
532 }
533 
534 namespace Eigen {
535 
536 template<> struct NumTraits<Eigen::half>
537  : GenericNumTraits<Eigen::half>
538 {
539  enum {
540  IsSigned = true,
541  IsInteger = false,
542  IsComplex = false,
543  RequireInitialization = false
544  };
545 
546  EIGEN_DEVICE_FUNC static EIGEN_STRONG_INLINE Eigen::half epsilon() {
547  return half_impl::raw_uint16_to_half(0x0800);
548  }
549  EIGEN_DEVICE_FUNC static EIGEN_STRONG_INLINE Eigen::half dummy_precision() { return Eigen::half(1e-2f); }
550  EIGEN_DEVICE_FUNC static EIGEN_STRONG_INLINE Eigen::half highest() {
551  return half_impl::raw_uint16_to_half(0x7bff);
552  }
553  EIGEN_DEVICE_FUNC static EIGEN_STRONG_INLINE Eigen::half lowest() {
554  return half_impl::raw_uint16_to_half(0xfbff);
555  }
556  EIGEN_DEVICE_FUNC static EIGEN_STRONG_INLINE Eigen::half infinity() {
557  return half_impl::raw_uint16_to_half(0x7c00);
558  }
559  EIGEN_DEVICE_FUNC static EIGEN_STRONG_INLINE Eigen::half quiet_NaN() {
560  return half_impl::raw_uint16_to_half(0x7c01);
561  }
562 };
563 
564 } // end namespace Eigen
565 
566 // C-like standard mathematical functions and trancendentals.
567 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC Eigen::half fabsh(const Eigen::half& a) {
568  Eigen::half result;
569  result.x = a.x & 0x7FFF;
570  return result;
571 }
572 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC Eigen::half exph(const Eigen::half& a) {
573  return Eigen::half(::expf(float(a)));
574 }
575 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC Eigen::half logh(const Eigen::half& a) {
576 #if EIGEN_CUDACC_VER >= 80000 && defined(EIGEN_CUDA_ARCH) && EIGEN_CUDA_ARCH >= 530
577  return Eigen::half(::hlog(a));
578 #else
579  return Eigen::half(::logf(float(a)));
580 #endif
581 }
582 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC Eigen::half sqrth(const Eigen::half& a) {
583  return Eigen::half(::sqrtf(float(a)));
584 }
585 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC Eigen::half powh(const Eigen::half& a, const Eigen::half& b) {
586  return Eigen::half(::powf(float(a), float(b)));
587 }
588 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC Eigen::half floorh(const Eigen::half& a) {
589  return Eigen::half(::floorf(float(a)));
590 }
591 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC Eigen::half ceilh(const Eigen::half& a) {
592  return Eigen::half(::ceilf(float(a)));
593 }
594 
595 namespace std {
596 
597 #if __cplusplus > 199711L
598 template <>
599 struct hash<Eigen::half> {
600  EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE std::size_t operator()(const Eigen::half& a) const {
601  return static_cast<std::size_t>(a.x);
602  }
603 };
604 #endif
605 
606 } // end namespace std
607 
608 
609 // Add the missing shfl_xor intrinsic
610 #if defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 300
611 __device__ EIGEN_STRONG_INLINE Eigen::half __shfl_xor(Eigen::half var, int laneMask, int width=warpSize) {
612  return static_cast<Eigen::half>(__shfl_xor(static_cast<float>(var), laneMask, width));
613 }
614 #endif
615 
616 // ldg() has an overload for __half, but we also need one for Eigen::half.
617 #if defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 350
618 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC Eigen::half __ldg(const Eigen::half* ptr) {
619  return Eigen::half_impl::raw_uint16_to_half(
620  __ldg(reinterpret_cast<const unsigned short*>(ptr)));
621 }
622 #endif
623 
624 
625 #if defined(__CUDA_ARCH__)
626 namespace Eigen {
627 namespace numext {
628 
629 template<>
630 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
631 bool (isnan)(const Eigen::half& h) {
632  return (half_impl::isnan)(h);
633 }
634 
635 template<>
636 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
637 bool (isinf)(const Eigen::half& h) {
638  return (half_impl::isinf)(h);
639 }
640 
641 template<>
642 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
643 bool (isfinite)(const Eigen::half& h) {
644  return (half_impl::isfinite)(h);
645 }
646 
647 } // namespace Eigen
648 } // namespace numext
649 #endif
650 
651 #endif // EIGEN_HALF_CUDA_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_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:306
Definition: Half.h:495
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_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:543
EIGEN_DEFAULT_DENSE_INDEX_TYPE Index
The Index type as used for the API.
Definition: Meta.h:33
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_abs_op< typename Derived::Scalar >, const Derived > abs(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_sign_op< typename Derived::Scalar >, const Derived > sign(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)