Half.h
Go to the documentation of this file.
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 either from CUDA's or HIP'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_H
37 #define EIGEN_HALF_H
38 
39 #include "../../InternalHeaderCheck.h"
40 
41 #if defined(EIGEN_HAS_GPU_FP16) || defined(EIGEN_HAS_ARM64_FP16_SCALAR_ARITHMETIC)
42 // When compiling with GPU support, the "__half_raw" base class as well as
43 // some other routines are defined in the GPU compiler header files
44 // (cuda_fp16.h, hip_fp16.h), and they are not tagged constexpr
45 // As a consequence, we get compile failures when compiling Eigen with
46 // GPU support. Hence the need to disable EIGEN_CONSTEXPR when building
47 // Eigen with GPU support
48  #pragma push_macro("EIGEN_CONSTEXPR")
49  #undef EIGEN_CONSTEXPR
50  #define EIGEN_CONSTEXPR
51 #endif
52 
53 #define F16_PACKET_FUNCTION(PACKET_F, PACKET_F16, METHOD) \
54  template <> \
55  EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC EIGEN_UNUSED \
56  PACKET_F16 METHOD<PACKET_F16>(const PACKET_F16& _x) { \
57  return float2half(METHOD<PACKET_F>(half2float(_x))); \
58  }
59 
60 namespace Eigen {
61 
62 struct half;
63 
64 namespace half_impl {
65 
66 // We want to use the __half_raw struct from the HIP header file only during the device compile phase.
67 // This is required because of a quirk in the way TensorFlow GPU builds are done.
68 // When compiling TensorFlow source code with GPU support, files that
69 // * contain GPU kernels (i.e. *.cu.cc files) are compiled via hipcc
70 // * do not contain GPU kernels ( i.e. *.cc files) are compiled via gcc (typically)
71 //
72 // Tensorflow uses the Eigen::half type as its FP16 type, and there are functions that
73 // * are defined in a file that gets compiled via hipcc AND
74 // * have Eigen::half as a pass-by-value argument AND
75 // * are called in a file that gets compiled via gcc
76 //
77 // In the scenario described above the caller and callee will see different versions
78 // of the Eigen::half base class __half_raw, and they will be compiled by different compilers
79 //
80 // There appears to be an ABI mismatch between gcc and clang (which is called by hipcc) that results in
81 // the callee getting corrupted values for the Eigen::half argument.
82 //
83 // Making the host side compile phase of hipcc use the same Eigen::half impl, as the gcc compile, resolves
84 // this error, and hence the following convoluted #if condition
85 #if !defined(EIGEN_HAS_GPU_FP16) || !defined(EIGEN_GPU_COMPILE_PHASE)
86 // Make our own __half_raw definition that is similar to CUDA's.
87 struct __half_raw {
88 #if (defined(EIGEN_HAS_GPU_FP16) && !defined(EIGEN_GPU_COMPILE_PHASE))
89  // Eigen::half can be used as the datatype for shared memory declarations (in Eigen and TF)
90  // The element type for shared memory cannot have non-trivial constructors
91  // and hence the following special casing (which skips the zero-initilization).
92  // Note that this check gets done even in the host compilation phase, and
93  // hence the need for this
95 #else
97 #endif
98 #if defined(EIGEN_HAS_ARM64_FP16_SCALAR_ARITHMETIC)
99  explicit EIGEN_DEVICE_FUNC EIGEN_CONSTEXPR __half_raw(numext::uint16_t raw) : x(numext::bit_cast<__fp16>(raw)) {
100  }
101  __fp16 x;
102 #else
105 #endif
106 };
107 
108 #elif defined(EIGEN_HAS_HIP_FP16)
109  // Nothing to do here
110  // HIP fp16 header file has a definition for __half_raw
111 #elif defined(EIGEN_HAS_CUDA_FP16)
112  #if EIGEN_CUDA_SDK_VER < 90000
113  // In CUDA < 9.0, __half is the equivalent of CUDA 9's __half_raw
114  typedef __half __half_raw;
115  #endif // defined(EIGEN_HAS_CUDA_FP16)
116 #elif defined(SYCL_DEVICE_ONLY)
117  typedef cl::sycl::half __half_raw;
118 #endif
119 
121 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC __half_raw float_to_half_rtne(float ff);
122 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC float half_to_float(__half_raw h);
123 
124 struct half_base : public __half_raw {
127 
128 #if defined(EIGEN_HAS_GPU_FP16)
129  #if defined(EIGEN_HAS_HIP_FP16)
130  EIGEN_DEVICE_FUNC EIGEN_CONSTEXPR half_base(const __half& h) { x = __half_as_ushort(h); }
131  #elif defined(EIGEN_HAS_CUDA_FP16)
132  #if EIGEN_CUDA_SDK_VER >= 90000
134  #endif
135  #endif
136 #endif
137 };
138 
139 } // namespace half_impl
140 
141 // Class definition.
142 struct half : public half_impl::half_base {
143 
144  // Writing this out as separate #if-else blocks to make the code easier to follow
145  // The same applies to most #if-else blocks in this file
146 #if !defined(EIGEN_HAS_GPU_FP16) || !defined(EIGEN_GPU_COMPILE_PHASE)
147  // Use the same base class for the following two scenarios
148  // * when compiling without GPU support enabled
149  // * during host compile phase when compiling with GPU support enabled
151 #elif defined(EIGEN_HAS_HIP_FP16)
152  // Nothing to do here
153  // HIP fp16 header file has a definition for __half_raw
154 #elif defined(EIGEN_HAS_CUDA_FP16)
155  // Note that EIGEN_CUDA_SDK_VER is set to 0 even when compiling with HIP, so
156  // (EIGEN_CUDA_SDK_VER < 90000) is true even for HIP! So keeping this within
157  // #if defined(EIGEN_HAS_CUDA_FP16) is needed
158  #if defined(EIGEN_CUDA_SDK_VER) && EIGEN_CUDA_SDK_VER < 90000
160  #endif
161 #endif
162 
164 
166 
167 #if defined(EIGEN_HAS_GPU_FP16)
168  #if defined(EIGEN_HAS_HIP_FP16)
169  EIGEN_DEVICE_FUNC EIGEN_CONSTEXPR half(const __half& h) : half_impl::half_base(h) {}
170  #elif defined(EIGEN_HAS_CUDA_FP16)
171  #if defined(EIGEN_CUDA_SDK_VER) && EIGEN_CUDA_SDK_VER >= 90000
172  EIGEN_DEVICE_FUNC EIGEN_CONSTEXPR half(const __half& h) : half_impl::half_base(h) {}
173  #endif
174  #endif
175 #endif
176 
177 
179  : half_impl::half_base(half_impl::raw_uint16_to_half(b ? 0x3c00 : 0)) {}
180  template<class T>
181  explicit EIGEN_DEVICE_FUNC half(T val)
182  : half_impl::half_base(half_impl::float_to_half_rtne(static_cast<float>(val))) {}
183  explicit EIGEN_DEVICE_FUNC half(float f)
184  : half_impl::half_base(half_impl::float_to_half_rtne(f)) {}
185 
186  // Following the convention of numpy, converting between complex and
187  // float will lead to loss of imag value.
188  template<typename RealScalar>
189  explicit EIGEN_DEVICE_FUNC half(std::complex<RealScalar> c)
190  : half_impl::half_base(half_impl::float_to_half_rtne(static_cast<float>(c.real()))) {}
191 
192  EIGEN_DEVICE_FUNC operator float() const { // NOLINT: Allow implicit conversion to float, because it is lossless.
193  return half_impl::half_to_float(*this);
194  }
195 
196 #if defined(EIGEN_HAS_GPU_FP16) && !defined(EIGEN_GPU_COMPILE_PHASE)
197  EIGEN_DEVICE_FUNC operator __half() const {
198  ::__half_raw hr;
199  hr.x = x;
200  return __half(hr);
201  }
202 #endif
203 };
204 
205 // TODO(majnemer): Get rid of this once we can rely on C++17 inline variables do
206 // solve the ODR issue.
207 namespace half_impl {
208 template <typename = void>
210  static EIGEN_CONSTEXPR const bool is_specialized = true;
211  static EIGEN_CONSTEXPR const bool is_signed = true;
212  static EIGEN_CONSTEXPR const bool is_integer = false;
213  static EIGEN_CONSTEXPR const bool is_exact = false;
214  static EIGEN_CONSTEXPR const bool has_infinity = true;
215  static EIGEN_CONSTEXPR const bool has_quiet_NaN = true;
216  static EIGEN_CONSTEXPR const bool has_signaling_NaN = true;
217  static EIGEN_CONSTEXPR const std::float_denorm_style has_denorm = std::denorm_present;
218  static EIGEN_CONSTEXPR const bool has_denorm_loss = false;
219  static EIGEN_CONSTEXPR const std::float_round_style round_style = std::round_to_nearest;
220  static EIGEN_CONSTEXPR const bool is_iec559 = true;
221  // The C++ standard defines this as "true if the set of values representable
222  // by the type is finite." Half has finite precision.
223  static EIGEN_CONSTEXPR const bool is_bounded = true;
224  static EIGEN_CONSTEXPR const bool is_modulo = false;
225  static EIGEN_CONSTEXPR const int digits = 11;
226  static EIGEN_CONSTEXPR const int digits10 = 3; // according to http://half.sourceforge.net/structstd_1_1numeric__limits_3_01half__float_1_1half_01_4.html
227  static EIGEN_CONSTEXPR const int max_digits10 = 5; // according to http://half.sourceforge.net/structstd_1_1numeric__limits_3_01half__float_1_1half_01_4.html
228  static EIGEN_CONSTEXPR const int radix = std::numeric_limits<float>::radix;
229  static EIGEN_CONSTEXPR const int min_exponent = -13;
230  static EIGEN_CONSTEXPR const int min_exponent10 = -4;
231  static EIGEN_CONSTEXPR const int max_exponent = 16;
232  static EIGEN_CONSTEXPR const int max_exponent10 = 4;
233  static EIGEN_CONSTEXPR const bool traps = std::numeric_limits<float>::traps;
234  // IEEE754: "The implementer shall choose how tininess is detected, but shall
235  // detect tininess in the same way for all operations in radix two"
236  static EIGEN_CONSTEXPR const bool tinyness_before = std::numeric_limits<float>::tinyness_before;
237 
247 };
248 
249 template<typename T>
251 template<typename T>
253 template<typename T>
255 template<typename T>
257 template<typename T>
259 template<typename T>
261 template<typename T>
263 template<typename T>
264 EIGEN_CONSTEXPR const std::float_denorm_style numeric_limits_half_impl<T>::has_denorm;
265 template<typename T>
267 template<typename T>
268 EIGEN_CONSTEXPR const std::float_round_style numeric_limits_half_impl<T>::round_style;
269 template<typename T>
271 template<typename T>
273 template<typename T>
275 template<typename T>
277 template<typename T>
279 template<typename T>
281 template<typename T>
283 template<typename T>
285 template<typename T>
287 template<typename T>
289 template<typename T>
291 template<typename T>
293 template<typename T>
295 } // end namespace half_impl
296 } // end namespace Eigen
297 
298 namespace std {
299 // If std::numeric_limits<T> is specialized, should also specialize
300 // std::numeric_limits<const T>, std::numeric_limits<volatile T>, and
301 // std::numeric_limits<const volatile T>
302 // https://stackoverflow.com/a/16519653/
303 template<>
304 class numeric_limits<Eigen::half> : public Eigen::half_impl::numeric_limits_half_impl<> {};
305 template<>
306 class numeric_limits<const Eigen::half> : public numeric_limits<Eigen::half> {};
307 template<>
308 class numeric_limits<volatile Eigen::half> : public numeric_limits<Eigen::half> {};
309 template<>
310 class numeric_limits<const volatile Eigen::half> : public numeric_limits<Eigen::half> {};
311 } // end namespace std
312 
313 namespace Eigen {
314 
315 namespace half_impl {
316 
317 #if (defined(EIGEN_HAS_CUDA_FP16) && defined(EIGEN_CUDA_ARCH) && \
318  EIGEN_CUDA_ARCH >= 530) || \
319  (defined(EIGEN_HAS_HIP_FP16) && defined(HIP_DEVICE_COMPILE))
320 // Note: We deliberately do *not* define this to 1 even if we have Arm's native
321 // fp16 type since GPU halfs are rather different from native CPU halfs.
322 // TODO: Rename to something like EIGEN_HAS_NATIVE_GPU_FP16
323 #define EIGEN_HAS_NATIVE_FP16
324 #endif
325 
326 // Intrinsics for native fp16 support. Note that on current hardware,
327 // these are no faster than fp32 arithmetic (you need to use the half2
328 // versions to get the ALU speed increased), but you do save the
329 // conversion steps back and forth.
330 
331 #if defined(EIGEN_HAS_NATIVE_FP16)
332 EIGEN_STRONG_INLINE __device__ half operator + (const half& a, const half& b) {
333 #if defined(EIGEN_CUDA_SDK_VER) && EIGEN_CUDA_SDK_VER >= 90000
334  return __hadd(::__half(a), ::__half(b));
335 #else
336  return __hadd(a, b);
337 #endif
338 }
339 EIGEN_STRONG_INLINE __device__ half operator * (const half& a, const half& b) {
340  return __hmul(a, b);
341 }
342 EIGEN_STRONG_INLINE __device__ half operator - (const half& a, const half& b) {
343  return __hsub(a, b);
344 }
345 EIGEN_STRONG_INLINE __device__ half operator / (const half& a, const half& b) {
346 #if defined(EIGEN_CUDA_SDK_VER) && EIGEN_CUDA_SDK_VER >= 90000
347  return __hdiv(a, b);
348 #else
349  float num = __half2float(a);
350  float denom = __half2float(b);
351  return __float2half(num / denom);
352 #endif
353 }
354 EIGEN_STRONG_INLINE __device__ half operator - (const half& a) {
355  return __hneg(a);
356 }
357 EIGEN_STRONG_INLINE __device__ half& operator += (half& a, const half& b) {
358  a = a + b;
359  return a;
360 }
361 EIGEN_STRONG_INLINE __device__ half& operator *= (half& a, const half& b) {
362  a = a * b;
363  return a;
364 }
365 EIGEN_STRONG_INLINE __device__ half& operator -= (half& a, const half& b) {
366  a = a - b;
367  return a;
368 }
369 EIGEN_STRONG_INLINE __device__ half& operator /= (half& a, const half& b) {
370  a = a / b;
371  return a;
372 }
373 EIGEN_STRONG_INLINE __device__ bool operator == (const half& a, const half& b) {
374  return __heq(a, b);
375 }
376 EIGEN_STRONG_INLINE __device__ bool operator != (const half& a, const half& b) {
377  return __hne(a, b);
378 }
379 EIGEN_STRONG_INLINE __device__ bool operator < (const half& a, const half& b) {
380  return __hlt(a, b);
381 }
382 EIGEN_STRONG_INLINE __device__ bool operator <= (const half& a, const half& b) {
383  return __hle(a, b);
384 }
385 EIGEN_STRONG_INLINE __device__ bool operator > (const half& a, const half& b) {
386  return __hgt(a, b);
387 }
388 EIGEN_STRONG_INLINE __device__ bool operator >= (const half& a, const half& b) {
389  return __hge(a, b);
390 }
391 #endif
392 
393 #if defined(EIGEN_HAS_ARM64_FP16_SCALAR_ARITHMETIC) && !defined(EIGEN_GPU_COMPILE_PHASE)
394 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half operator + (const half& a, const half& b) {
395  return half(vaddh_f16(a.x, b.x));
396 }
397 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half operator * (const half& a, const half& b) {
398  return half(vmulh_f16(a.x, b.x));
399 }
400 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half operator - (const half& a, const half& b) {
401  return half(vsubh_f16(a.x, b.x));
402 }
403 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half operator / (const half& a, const half& b) {
404  return half(vdivh_f16(a.x, b.x));
405 }
406 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half operator - (const half& a) {
407  return half(vnegh_f16(a.x));
408 }
409 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half& operator += (half& a, const half& b) {
410  a = half(vaddh_f16(a.x, b.x));
411  return a;
412 }
413 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half& operator *= (half& a, const half& b) {
414  a = half(vmulh_f16(a.x, b.x));
415  return a;
416 }
417 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half& operator -= (half& a, const half& b) {
418  a = half(vsubh_f16(a.x, b.x));
419  return a;
420 }
421 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half& operator /= (half& a, const half& b) {
422  a = half(vdivh_f16(a.x, b.x));
423  return a;
424 }
425 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC bool operator == (const half& a, const half& b) {
426  return vceqh_f16(a.x, b.x);
427 }
428 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC bool operator != (const half& a, const half& b) {
429  return !vceqh_f16(a.x, b.x);
430 }
431 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC bool operator < (const half& a, const half& b) {
432  return vclth_f16(a.x, b.x);
433 }
434 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC bool operator <= (const half& a, const half& b) {
435  return vcleh_f16(a.x, b.x);
436 }
437 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC bool operator > (const half& a, const half& b) {
438  return vcgth_f16(a.x, b.x);
439 }
440 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC bool operator >= (const half& a, const half& b) {
441  return vcgeh_f16(a.x, b.x);
442 }
443 // We need to distinguish ‘clang as the CUDA compiler’ from ‘clang as the host compiler,
444 // invoked by NVCC’ (e.g. on MacOS). The former needs to see both host and device implementation
445 // of the functions, while the latter can only deal with one of them.
446 #elif !defined(EIGEN_HAS_NATIVE_FP16) || (EIGEN_COMP_CLANG && !EIGEN_COMP_NVCC) // Emulate support for half floats
447 
448 #if EIGEN_COMP_CLANG && defined(EIGEN_GPUCC)
449 // We need to provide emulated *host-side* FP16 operators for clang.
450 #pragma push_macro("EIGEN_DEVICE_FUNC")
451 #undef EIGEN_DEVICE_FUNC
452 #if defined(EIGEN_HAS_CUDA_FP16) && defined(EIGEN_HAS_NATIVE_FP16)
453 #define EIGEN_DEVICE_FUNC __host__
454 #else // both host and device need emulated ops.
455 #define EIGEN_DEVICE_FUNC __host__ __device__
456 #endif
457 #endif
458 
459 // Definitions for CPUs and older HIP+CUDA, mostly working through conversion
460 // to/from fp32.
461 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half operator + (const half& a, const half& b) {
462  return half(float(a) + float(b));
463 }
464 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half operator * (const half& a, const half& b) {
465  return half(float(a) * float(b));
466 }
467 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half operator - (const half& a, const half& b) {
468  return half(float(a) - float(b));
469 }
470 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half operator / (const half& a, const half& b) {
471  return half(float(a) / float(b));
472 }
473 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half operator - (const half& a) {
474  half result;
475  result.x = a.x ^ 0x8000;
476  return result;
477 }
478 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half& operator += (half& a, const half& b) {
479  a = half(float(a) + float(b));
480  return a;
481 }
482 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half& operator *= (half& a, const half& b) {
483  a = half(float(a) * float(b));
484  return a;
485 }
486 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half& operator -= (half& a, const half& b) {
487  a = half(float(a) - float(b));
488  return a;
489 }
490 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half& operator /= (half& a, const half& b) {
491  a = half(float(a) / float(b));
492  return a;
493 }
494 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC bool operator == (const half& a, const half& b) {
495  return numext::equal_strict(float(a),float(b));
496 }
497 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC bool operator != (const half& a, const half& b) {
498  return numext::not_equal_strict(float(a), float(b));
499 }
500 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC bool operator < (const half& a, const half& b) {
501  return float(a) < float(b);
502 }
503 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC bool operator <= (const half& a, const half& b) {
504  return float(a) <= float(b);
505 }
506 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC bool operator > (const half& a, const half& b) {
507  return float(a) > float(b);
508 }
509 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC bool operator >= (const half& a, const half& b) {
510  return float(a) >= float(b);
511 }
512 
513 #if EIGEN_COMP_CLANG && defined(EIGEN_GPUCC)
514 #pragma pop_macro("EIGEN_DEVICE_FUNC")
515 #endif
516 #endif // Emulate support for half floats
517 
518 // Division by an index. Do it in full float precision to avoid accuracy
519 // issues in converting the denominator to half.
520 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half operator / (const half& a, Index b) {
521  return half(static_cast<float>(a) / static_cast<float>(b));
522 }
523 
524 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half operator++(half& a) {
525  a += half(1);
526  return a;
527 }
528 
529 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half operator--(half& a) {
530  a -= half(1);
531  return a;
532 }
533 
534 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half operator++(half& a, int) {
535  half original_value = a;
536  ++a;
537  return original_value;
538 }
539 
540 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half operator--(half& a, int) {
541  half original_value = a;
542  --a;
543  return original_value;
544 }
545 
546 // Conversion routines, including fallbacks for the host or older CUDA.
547 // Note that newer Intel CPUs (Haswell or newer) have vectorized versions of
548 // these in hardware. If we need more performance on older/other CPUs, they are
549 // also possible to vectorize directly.
550 
552  // We cannot simply do a "return __half_raw(x)" here, because __half_raw is union type
553  // in the hip_fp16 header file, and that will trigger a compile error
554  // On the other hand, having anything but a return statement also triggers a compile error
555  // because this is constexpr function.
556  // Fortunately, since we need to disable EIGEN_CONSTEXPR for GPU anyway, we can get out
557  // of this catch22 by having separate bodies for GPU / non GPU
558 #if defined(EIGEN_HAS_GPU_FP16)
559  __half_raw h;
560  h.x = x;
561  return h;
562 #else
563  return __half_raw(x);
564 #endif
565 }
566 
568  // HIP/CUDA/Default have a member 'x' of type uint16_t.
569  // For ARM64 native half, the member 'x' is of type __fp16, so we need to bit-cast.
570  // For SYCL, cl::sycl::half is _Float16, so cast directly.
571 #if defined(EIGEN_HAS_ARM64_FP16_SCALAR_ARITHMETIC)
572  return numext::bit_cast<numext::uint16_t>(h.x);
573 #elif defined(SYCL_DEVICE_ONLY)
574  return numext::bit_cast<numext::uint16_t>(h);
575 #else
576  return h.x;
577 #endif
578 }
579 
581  unsigned int u;
582  float f;
583 };
584 
585 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC __half_raw float_to_half_rtne(float ff) {
586 #if (defined(EIGEN_HAS_CUDA_FP16) && defined(EIGEN_CUDA_ARCH) && EIGEN_CUDA_ARCH >= 300) || \
587  (defined(EIGEN_HAS_HIP_FP16) && defined(EIGEN_HIP_DEVICE_COMPILE))
588  __half tmp_ff = __float2half(ff);
589  return *(__half_raw*)&tmp_ff;
590 
591 #elif defined(EIGEN_HAS_FP16_C)
592  __half_raw h;
593  #if EIGEN_COMP_MSVC
594  // MSVC does not have scalar instructions.
595  h.x =_mm_extract_epi16(_mm_cvtps_ph(_mm_set_ss(ff), 0), 0);
596  #else
597  h.x = _cvtss_sh(ff, 0);
598  #endif
599  return h;
600 
601 #elif defined(EIGEN_HAS_ARM64_FP16_SCALAR_ARITHMETIC)
602  __half_raw h;
603  h.x = static_cast<__fp16>(ff);
604  return h;
605 
606 #else
607  float32_bits f; f.f = ff;
608 
609  const float32_bits f32infty = { 255 << 23 };
610  const float32_bits f16max = { (127 + 16) << 23 };
611  const float32_bits denorm_magic = { ((127 - 15) + (23 - 10) + 1) << 23 };
612  unsigned int sign_mask = 0x80000000u;
613  __half_raw o;
614  o.x = static_cast<numext::uint16_t>(0x0u);
615 
616  unsigned int sign = f.u & sign_mask;
617  f.u ^= sign;
618 
619  // NOTE all the integer compares in this function can be safely
620  // compiled into signed compares since all operands are below
621  // 0x80000000. Important if you want fast straight SSE2 code
622  // (since there's no unsigned PCMPGTD).
623 
624  if (f.u >= f16max.u) { // result is Inf or NaN (all exponent bits set)
625  o.x = (f.u > f32infty.u) ? 0x7e00 : 0x7c00; // NaN->qNaN and Inf->Inf
626  } else { // (De)normalized number or zero
627  if (f.u < (113 << 23)) { // resulting FP16 is subnormal or zero
628  // use a magic value to align our 10 mantissa bits at the bottom of
629  // the float. as long as FP addition is round-to-nearest-even this
630  // just works.
631  f.f += denorm_magic.f;
632 
633  // and one integer subtract of the bias later, we have our final float!
634  o.x = static_cast<numext::uint16_t>(f.u - denorm_magic.u);
635  } else {
636  unsigned int mant_odd = (f.u >> 13) & 1; // resulting mantissa is odd
637 
638  // update exponent, rounding bias part 1
639  // Equivalent to `f.u += ((unsigned int)(15 - 127) << 23) + 0xfff`, but
640  // without arithmetic overflow.
641  f.u += 0xc8000fffU;
642  // rounding bias part 2
643  f.u += mant_odd;
644  // take the bits!
645  o.x = static_cast<numext::uint16_t>(f.u >> 13);
646  }
647  }
648 
649  o.x |= static_cast<numext::uint16_t>(sign >> 16);
650  return o;
651 #endif
652 }
653 
654 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC float half_to_float(__half_raw h) {
655 #if (defined(EIGEN_HAS_CUDA_FP16) && defined(EIGEN_CUDA_ARCH) && EIGEN_CUDA_ARCH >= 300) || \
656  (defined(EIGEN_HAS_HIP_FP16) && defined(EIGEN_HIP_DEVICE_COMPILE))
657  return __half2float(h);
658 #elif defined(EIGEN_HAS_FP16_C)
659  #if EIGEN_COMP_MSVC
660  // MSVC does not have scalar instructions.
661  return _mm_cvtss_f32(_mm_cvtph_ps(_mm_set1_epi16(h.x)));
662  #else
663  return _cvtsh_ss(h.x);
664  #endif
665 #elif defined(EIGEN_HAS_ARM64_FP16_SCALAR_ARITHMETIC)
666  return static_cast<float>(h.x);
667 #else
668  const float32_bits magic = { 113 << 23 };
669  const unsigned int shifted_exp = 0x7c00 << 13; // exponent mask after shift
670  float32_bits o;
671 
672  o.u = (h.x & 0x7fff) << 13; // exponent/mantissa bits
673  unsigned int exp = shifted_exp & o.u; // just the exponent
674  o.u += (127 - 15) << 23; // exponent adjust
675 
676  // handle exponent special cases
677  if (exp == shifted_exp) { // Inf/NaN?
678  o.u += (128 - 16) << 23; // extra exp adjust
679  } else if (exp == 0) { // Zero/Denormal?
680  o.u += 1 << 23; // extra exp adjust
681  o.f -= magic.f; // renormalize
682  }
683 
684  o.u |= (h.x & 0x8000) << 16; // sign bit
685  return o.f;
686 #endif
687 }
688 
689 // --- standard functions ---
690 
691 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC bool (isinf)(const half& a) {
692 #ifdef EIGEN_HAS_ARM64_FP16_SCALAR_ARITHMETIC
693  return (numext::bit_cast<numext::uint16_t>(a.x) & 0x7fff) == 0x7c00;
694 #else
695  return (a.x & 0x7fff) == 0x7c00;
696 #endif
697 }
698 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC bool (isnan)(const half& a) {
699 #if (defined(EIGEN_HAS_CUDA_FP16) && defined(EIGEN_CUDA_ARCH) && EIGEN_CUDA_ARCH >= 530) || \
700  (defined(EIGEN_HAS_HIP_FP16) && defined(EIGEN_HIP_DEVICE_COMPILE))
701  return __hisnan(a);
702 #elif defined(EIGEN_HAS_ARM64_FP16_SCALAR_ARITHMETIC)
703  return (numext::bit_cast<numext::uint16_t>(a.x) & 0x7fff) > 0x7c00;
704 #else
705  return (a.x & 0x7fff) > 0x7c00;
706 #endif
707 }
708 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC bool (isfinite)(const half& a) {
709  return !(isinf EIGEN_NOT_A_MACRO (a)) && !(isnan EIGEN_NOT_A_MACRO (a));
710 }
711 
712 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half abs(const half& a) {
713 #if defined(EIGEN_HAS_ARM64_FP16_SCALAR_ARITHMETIC)
714  return half(vabsh_f16(a.x));
715 #else
716  half result;
717  result.x = a.x & 0x7FFF;
718  return result;
719 #endif
720 }
721 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half exp(const half& a) {
722 #if (EIGEN_CUDA_SDK_VER >= 80000 && defined EIGEN_CUDA_ARCH && EIGEN_CUDA_ARCH >= 530) || \
723  defined(EIGEN_HIP_DEVICE_COMPILE)
724  return half(hexp(a));
725 #else
726  return half(::expf(float(a)));
727 #endif
728 }
729 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half expm1(const half& a) {
730  return half(numext::expm1(float(a)));
731 }
732 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half log(const half& a) {
733 #if (defined(EIGEN_HAS_CUDA_FP16) && EIGEN_CUDA_SDK_VER >= 80000 && defined(EIGEN_CUDA_ARCH) && EIGEN_CUDA_ARCH >= 530) || \
734  (defined(EIGEN_HAS_HIP_FP16) && defined(EIGEN_HIP_DEVICE_COMPILE))
735  return half(::hlog(a));
736 #else
737  return half(::logf(float(a)));
738 #endif
739 }
740 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half log1p(const half& a) {
741  return half(numext::log1p(float(a)));
742 }
743 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half log10(const half& a) {
744  return half(::log10f(float(a)));
745 }
746 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half log2(const half& a) {
747  return half(static_cast<float>(EIGEN_LOG2E) * ::logf(float(a)));
748 }
749 
750 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half sqrt(const half& a) {
751 #if (EIGEN_CUDA_SDK_VER >= 80000 && defined EIGEN_CUDA_ARCH && EIGEN_CUDA_ARCH >= 530) || \
752  defined(EIGEN_HIP_DEVICE_COMPILE)
753  return half(hsqrt(a));
754 #else
755  return half(::sqrtf(float(a)));
756 #endif
757 }
758 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half pow(const half& a, const half& b) {
759  return half(::powf(float(a), float(b)));
760 }
761 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half atan2(const half& a, const half& b) {
762  return half(::atan2f(float(a), float(b)));
763 }
764 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half sin(const half& a) {
765  return half(::sinf(float(a)));
766 }
767 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half cos(const half& a) {
768  return half(::cosf(float(a)));
769 }
770 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half tan(const half& a) {
771  return half(::tanf(float(a)));
772 }
773 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half tanh(const half& a) {
774  return half(::tanhf(float(a)));
775 }
776 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half asin(const half& a) {
777  return half(::asinf(float(a)));
778 }
779 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half acos(const half& a) {
780  return half(::acosf(float(a)));
781 }
782 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half atan(const half& a) {
783  return half(::atanf(float(a)));
784 }
785 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half atanh(const half& a) {
786  return half(::atanhf(float(a)));
787 }
788 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half floor(const half& a) {
789 #if (EIGEN_CUDA_SDK_VER >= 80000 && defined EIGEN_CUDA_ARCH && EIGEN_CUDA_ARCH >= 300) || \
790  defined(EIGEN_HIP_DEVICE_COMPILE)
791  return half(hfloor(a));
792 #else
793  return half(::floorf(float(a)));
794 #endif
795 }
796 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half ceil(const half& a) {
797 #if (EIGEN_CUDA_SDK_VER >= 80000 && defined EIGEN_CUDA_ARCH && EIGEN_CUDA_ARCH >= 300) || \
798  defined(EIGEN_HIP_DEVICE_COMPILE)
799  return half(hceil(a));
800 #else
801  return half(::ceilf(float(a)));
802 #endif
803 }
804 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half rint(const half& a) {
805  return half(::rintf(float(a)));
806 }
807 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half round(const half& a) {
808  return half(::roundf(float(a)));
809 }
810 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half fmod(const half& a, const half& b) {
811  return half(::fmodf(float(a), float(b)));
812 }
813 
814 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half (min)(const half& a, const half& b) {
815 #if (defined(EIGEN_HAS_CUDA_FP16) && defined(EIGEN_CUDA_ARCH) && EIGEN_CUDA_ARCH >= 530) || \
816  (defined(EIGEN_HAS_HIP_FP16) && defined(EIGEN_HIP_DEVICE_COMPILE))
817  return __hlt(b, a) ? b : a;
818 #else
819  const float f1 = static_cast<float>(a);
820  const float f2 = static_cast<float>(b);
821  return f2 < f1 ? b : a;
822 #endif
823 }
824 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half (max)(const half& a, const half& b) {
825 #if (defined(EIGEN_HAS_CUDA_FP16) && defined(EIGEN_CUDA_ARCH) && EIGEN_CUDA_ARCH >= 530) || \
826  (defined(EIGEN_HAS_HIP_FP16) && defined(EIGEN_HIP_DEVICE_COMPILE))
827  return __hlt(a, b) ? b : a;
828 #else
829  const float f1 = static_cast<float>(a);
830  const float f2 = static_cast<float>(b);
831  return f1 < f2 ? b : a;
832 #endif
833 }
834 
835 #ifndef EIGEN_NO_IO
836 EIGEN_ALWAYS_INLINE std::ostream& operator << (std::ostream& os, const half& v) {
837  os << static_cast<float>(v);
838  return os;
839 }
840 #endif
841 
842 } // end namespace half_impl
843 
844 // import Eigen::half_impl::half into Eigen namespace
845 // using half_impl::half;
846 
847 namespace internal {
848 
849 template<>
850 struct random_default_impl<half, false, false>
851 {
852  static inline half run(const half& x, const half& y)
853  {
854  return x + (y-x) * half(float(std::rand()) / float(RAND_MAX));
855  }
856  static inline half run()
857  {
858  return run(half(-1.f), half(1.f));
859  }
860 };
861 
862 template<> struct is_arithmetic<half> { enum { value = true }; };
863 
864 } // end namespace internal
865 
866 template<> struct NumTraits<Eigen::half>
867  : GenericNumTraits<Eigen::half>
868 {
869  enum {
870  IsSigned = true,
871  IsInteger = false,
872  IsComplex = false,
873  RequireInitialization = false
874  };
875 
876  EIGEN_DEVICE_FUNC EIGEN_CONSTEXPR static EIGEN_STRONG_INLINE Eigen::half epsilon() {
877  return half_impl::raw_uint16_to_half(0x0800);
878  }
880  return half_impl::raw_uint16_to_half(0x211f); // Eigen::half(1e-2f);
881  }
882  EIGEN_DEVICE_FUNC EIGEN_CONSTEXPR static EIGEN_STRONG_INLINE Eigen::half highest() {
883  return half_impl::raw_uint16_to_half(0x7bff);
884  }
885  EIGEN_DEVICE_FUNC EIGEN_CONSTEXPR static EIGEN_STRONG_INLINE Eigen::half lowest() {
886  return half_impl::raw_uint16_to_half(0xfbff);
887  }
888  EIGEN_DEVICE_FUNC EIGEN_CONSTEXPR static EIGEN_STRONG_INLINE Eigen::half infinity() {
889  return half_impl::raw_uint16_to_half(0x7c00);
890  }
892  return half_impl::raw_uint16_to_half(0x7e00);
893  }
894 };
895 
896 } // end namespace Eigen
897 
898 #if defined(EIGEN_HAS_GPU_FP16) || defined(EIGEN_HAS_ARM64_FP16_SCALAR_ARITHMETIC)
899  #pragma pop_macro("EIGEN_CONSTEXPR")
900 #endif
901 
902 namespace Eigen {
903 namespace numext {
904 
905 #if defined(EIGEN_GPU_COMPILE_PHASE)
906 
907 template <>
909  return (half_impl::isnan)(h);
910 }
911 
912 template <>
914  return (half_impl::isinf)(h);
915 }
916 
917 template <>
919  return (half_impl::isfinite)(h);
920 }
921 
922 #endif
923 
924 template <>
925 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC Eigen::half bit_cast<Eigen::half, uint16_t>(const uint16_t& src) {
927 }
928 
929 template <>
930 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC uint16_t bit_cast<uint16_t, Eigen::half>(const Eigen::half& src) {
932 }
933 
934 } // namespace numext
935 } // namespace Eigen
936 
937 // Add the missing shfl* intrinsics.
938 // The __shfl* functions are only valid on HIP or _CUDA_ARCH_ >= 300.
939 // CUDA defines them for (__CUDA_ARCH__ >= 300 || !defined(__CUDA_ARCH__))
940 //
941 // HIP and CUDA prior to SDK 9.0 define
942 // __shfl, __shfl_up, __shfl_down, __shfl_xor for int and float
943 // CUDA since 9.0 deprecates those and instead defines
944 // __shfl_sync, __shfl_up_sync, __shfl_down_sync, __shfl_xor_sync,
945 // with native support for __half and __nv_bfloat16
946 //
947 // Note that the following are __device__ - only functions.
948 #if (defined(EIGEN_CUDACC) && (!defined(EIGEN_CUDA_ARCH) || EIGEN_CUDA_ARCH >= 300)) \
949  || defined(EIGEN_HIPCC)
950 
951 #if defined(EIGEN_HAS_CUDA_FP16) && EIGEN_CUDA_SDK_VER >= 90000
952 
953 __device__ EIGEN_STRONG_INLINE Eigen::half __shfl_sync(unsigned mask, Eigen::half var, int srcLane, int width=warpSize) {
954  const __half h = var;
955  return static_cast<Eigen::half>(__shfl_sync(mask, h, srcLane, width));
956 }
957 
958 __device__ EIGEN_STRONG_INLINE Eigen::half __shfl_up_sync(unsigned mask, Eigen::half var, unsigned int delta, int width=warpSize) {
959  const __half h = var;
960  return static_cast<Eigen::half>(__shfl_up_sync(mask, h, delta, width));
961 }
962 
963 __device__ EIGEN_STRONG_INLINE Eigen::half __shfl_down_sync(unsigned mask, Eigen::half var, unsigned int delta, int width=warpSize) {
964  const __half h = var;
965  return static_cast<Eigen::half>(__shfl_down_sync(mask, h, delta, width));
966 }
967 
968 __device__ EIGEN_STRONG_INLINE Eigen::half __shfl_xor_sync(unsigned mask, Eigen::half var, int laneMask, int width=warpSize) {
969  const __half h = var;
970  return static_cast<Eigen::half>(__shfl_xor_sync(mask, h, laneMask, width));
971 }
972 
973 #else // HIP or CUDA SDK < 9.0
974 
975 __device__ EIGEN_STRONG_INLINE Eigen::half __shfl(Eigen::half var, int srcLane, int width=warpSize) {
976  const int ivar = static_cast<int>(Eigen::numext::bit_cast<Eigen::numext::uint16_t>(var));
977  return Eigen::numext::bit_cast<Eigen::half>(static_cast<Eigen::numext::uint16_t>(__shfl(ivar, srcLane, width)));
978 }
979 
980 __device__ EIGEN_STRONG_INLINE Eigen::half __shfl_up(Eigen::half var, unsigned int delta, int width=warpSize) {
981  const int ivar = static_cast<int>(Eigen::numext::bit_cast<Eigen::numext::uint16_t>(var));
982  return Eigen::numext::bit_cast<Eigen::half>(static_cast<Eigen::numext::uint16_t>(__shfl_up(ivar, delta, width)));
983 }
984 
985 __device__ EIGEN_STRONG_INLINE Eigen::half __shfl_down(Eigen::half var, unsigned int delta, int width=warpSize) {
986  const int ivar = static_cast<int>(Eigen::numext::bit_cast<Eigen::numext::uint16_t>(var));
987  return Eigen::numext::bit_cast<Eigen::half>(static_cast<Eigen::numext::uint16_t>(__shfl_down(ivar, delta, width)));
988 }
989 
990 __device__ EIGEN_STRONG_INLINE Eigen::half __shfl_xor(Eigen::half var, int laneMask, int width=warpSize) {
991  const int ivar = static_cast<int>(Eigen::numext::bit_cast<Eigen::numext::uint16_t>(var));
992  return Eigen::numext::bit_cast<Eigen::half>(static_cast<Eigen::numext::uint16_t>(__shfl_xor(ivar, laneMask, width)));
993 }
994 
995 #endif // HIP vs CUDA
996 #endif // __shfl*
997 
998 // ldg() has an overload for __half_raw, but we also need one for Eigen::half.
999 #if (defined(EIGEN_CUDACC) && (!defined(EIGEN_CUDA_ARCH) || EIGEN_CUDA_ARCH >= 350)) \
1000  || defined(EIGEN_HIPCC)
1001 EIGEN_STRONG_INLINE __device__ Eigen::half __ldg(const Eigen::half* ptr) {
1002  return Eigen::half_impl::raw_uint16_to_half(__ldg(reinterpret_cast<const Eigen::numext::uint16_t*>(ptr)));
1003 }
1004 #endif // __ldg
1005 
1006 #if EIGEN_HAS_STD_HASH
1007 namespace std {
1008 template <>
1009 struct hash<Eigen::half> {
1010  EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE std::size_t operator()(const Eigen::half& a) const {
1011  return static_cast<std::size_t>(Eigen::numext::bit_cast<Eigen::numext::uint16_t>(a));
1012  }
1013 };
1014 } // end namespace std
1015 #endif
1016 
1017 namespace Eigen {
1018 namespace internal {
1019 
1020 template <>
1021 struct cast_impl<float, half> {
1023  static inline half run(const float& a) {
1024 #if (defined(EIGEN_HAS_CUDA_FP16) && defined(EIGEN_CUDA_ARCH) && EIGEN_CUDA_ARCH >= 300) || \
1025  (defined(EIGEN_HAS_HIP_FP16) && defined(EIGEN_HIP_DEVICE_COMPILE))
1026  return __float2half(a);
1027 #else
1028  return half(a);
1029 #endif
1030  }
1031 };
1032 
1033 template <>
1034 struct cast_impl<int, half> {
1036  static inline half run(const int& a) {
1037 #if (defined(EIGEN_HAS_CUDA_FP16) && defined(EIGEN_CUDA_ARCH) && EIGEN_CUDA_ARCH >= 300) || \
1038  (defined(EIGEN_HAS_HIP_FP16) && defined(EIGEN_HIP_DEVICE_COMPILE))
1039  return __float2half(static_cast<float>(a));
1040 #else
1041  return half(static_cast<float>(a));
1042 #endif
1043  }
1044 };
1045 
1046 template <>
1047 struct cast_impl<half, float> {
1049  static inline float run(const half& a) {
1050 #if (defined(EIGEN_HAS_CUDA_FP16) && defined(EIGEN_CUDA_ARCH) && EIGEN_CUDA_ARCH >= 300) || \
1051  (defined(EIGEN_HAS_HIP_FP16) && defined(EIGEN_HIP_DEVICE_COMPILE))
1052  return __half2float(a);
1053 #else
1054  return static_cast<float>(a);
1055 #endif
1056  }
1057 };
1058 
1059 } // namespace internal
1060 } // namespace Eigen
1061 
1062 #endif // EIGEN_HALF_H
const Log1pReturnType log1p() const
const Expm1ReturnType expm1() const
Array< int, Dynamic, 1 > v
Array< int, 3, 1 > b
Array33i c
IndexedView_or_Block operator()(const RowIndices &rowIndices, const ColIndices &colIndices)
#define EIGEN_ALWAYS_INLINE
Definition: Macros.h:836
#define EIGEN_CONSTEXPR
Definition: Macros.h:747
#define EIGEN_DEVICE_FUNC
Definition: Macros.h:883
#define EIGEN_NOT_A_MACRO
Definition: Macros.h:804
#define EIGEN_LOG2E
Definition: MathFunctions.h:17
half() max(const half &a, const half &b)
Definition: Half.h:824
half atan(const half &a)
Definition: Half.h:782
half operator--(half &a)
Definition: Half.h:529
EIGEN_ALWAYS_INLINE std::ostream & operator<<(std::ostream &os, const half &v)
Definition: Half.h:836
half atanh(const half &a)
Definition: Half.h:785
half & operator/=(half &a, const half &b)
Definition: Half.h:490
half round(const half &a)
Definition: Half.h:807
bool operator>(const half &a, const half &b)
Definition: Half.h:506
float half_to_float(__half_raw h)
Definition: Half.h:654
bool() isnan(const half &a)
Definition: Half.h:698
half floor(const half &a)
Definition: Half.h:788
half exp(const half &a)
Definition: Half.h:721
half operator*(const half &a, const half &b)
Definition: Half.h:464
bool operator>=(const half &a, const half &b)
Definition: Half.h:509
half tan(const half &a)
Definition: Half.h:770
half acos(const half &a)
Definition: Half.h:779
half rint(const half &a)
Definition: Half.h:804
half abs(const half &a)
Definition: Half.h:712
half() min(const half &a, const half &b)
Definition: Half.h:814
half log(const half &a)
Definition: Half.h:732
half operator++(half &a)
Definition: Half.h:524
numext::uint16_t raw_half_as_uint16(const __half_raw &h)
Definition: Half.h:567
half operator-(const half &a, const half &b)
Definition: Half.h:467
half atan2(const half &a, const half &b)
Definition: Half.h:761
half log10(const half &a)
Definition: Half.h:743
bool operator<=(const half &a, const half &b)
Definition: Half.h:503
half & operator*=(half &a, const half &b)
Definition: Half.h:482
half log2(const half &a)
Definition: Half.h:746
half & operator-=(half &a, const half &b)
Definition: Half.h:486
bool() isinf(const half &a)
Definition: Half.h:691
half fmod(const half &a, const half &b)
Definition: Half.h:810
half & operator+=(half &a, const half &b)
Definition: Half.h:478
half log1p(const half &a)
Definition: Half.h:740
EIGEN_CONSTEXPR __half_raw raw_uint16_to_half(numext::uint16_t x)
Definition: Half.h:551
half asin(const half &a)
Definition: Half.h:776
bool operator!=(const half &a, const half &b)
Definition: Half.h:497
half tanh(const half &a)
Definition: Half.h:773
half ceil(const half &a)
Definition: Half.h:796
__half_raw float_to_half_rtne(float ff)
Definition: Half.h:585
half cos(const half &a)
Definition: Half.h:767
half operator+(const half &a, const half &b)
Definition: Half.h:461
half sqrt(const half &a)
Definition: Half.h:750
bool operator<(const half &a, const half &b)
Definition: Half.h:500
half operator/(const half &a, const half &b)
Definition: Half.h:470
bool() isfinite(const half &a)
Definition: Half.h:708
half pow(const half &a, const half &b)
Definition: Half.h:758
half expm1(const half &a)
Definition: Half.h:729
bool operator==(const half &a, const half &b)
Definition: Half.h:494
half sin(const half &a)
Definition: Half.h:764
const Scalar & y
bool equal_strict(const X &x, const Y &y)
Definition: Meta.h:460
bool not_equal_strict(const X &x, const Y &y)
Definition: Meta.h:485
EIGEN_ALWAYS_INLINE bool() isinf(const Eigen::bfloat16 &h)
Definition: BFloat16.h:784
std::uint16_t uint16_t
Definition: Meta.h:37
Tgt bit_cast(const Src &src)
Definition: NumTraits.h:87
EIGEN_ALWAYS_INLINE bool() isnan(const Eigen::bfloat16 &h)
Definition: BFloat16.h:778
EIGEN_ALWAYS_INLINE bool() isfinite(const Eigen::bfloat16 &h)
Definition: BFloat16.h:790
: InteropHeaders
Definition: Core:139
EIGEN_DEFAULT_DENSE_INDEX_TYPE Index
The Index type as used for the API.
Definition: Meta.h:82
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_real_op< typename Derived::Scalar >, const Derived > real(const Eigen::ArrayBase< Derived > &x)
Definition: BFloat16.h:222
static EIGEN_CONSTEXPR Eigen::half infinity()
Definition: Half.h:888
static EIGEN_CONSTEXPR Eigen::half dummy_precision()
Definition: Half.h:879
static EIGEN_CONSTEXPR Eigen::half epsilon()
Definition: Half.h:876
static EIGEN_CONSTEXPR Eigen::half lowest()
Definition: Half.h:885
static EIGEN_CONSTEXPR Eigen::half highest()
Definition: Half.h:882
static EIGEN_CONSTEXPR Eigen::half quiet_NaN()
Definition: Half.h:891
Holds information about the various numeric (i.e. scalar) types allowed by Eigen.
Definition: NumTraits.h:231
EIGEN_CONSTEXPR __half_raw()
Definition: Half.h:96
numext::uint16_t x
Definition: Half.h:104
EIGEN_CONSTEXPR __half_raw(numext::uint16_t raw)
Definition: Half.h:103
EIGEN_CONSTEXPR half_base()
Definition: Half.h:125
EIGEN_CONSTEXPR half_base(const __half_raw &h)
Definition: Half.h:126
static EIGEN_CONSTEXPR const bool has_infinity
Definition: Half.h:214
static EIGEN_CONSTEXPR Eigen::half round_error()
Definition: Half.h:242
static EIGEN_CONSTEXPR const int radix
Definition: Half.h:228
static EIGEN_CONSTEXPR const int digits10
Definition: Half.h:226
static EIGEN_CONSTEXPR Eigen::half() min()
Definition: Half.h:238
static EIGEN_CONSTEXPR Eigen::half infinity()
Definition: Half.h:243
static EIGEN_CONSTEXPR const int min_exponent
Definition: Half.h:229
static EIGEN_CONSTEXPR const int max_exponent10
Definition: Half.h:232
static EIGEN_CONSTEXPR Eigen::half epsilon()
Definition: Half.h:241
static EIGEN_CONSTEXPR Eigen::half() max()
Definition: Half.h:240
static EIGEN_CONSTEXPR const bool has_quiet_NaN
Definition: Half.h:215
static EIGEN_CONSTEXPR const bool is_signed
Definition: Half.h:211
static EIGEN_CONSTEXPR const int max_digits10
Definition: Half.h:227
static EIGEN_CONSTEXPR const bool is_iec559
Definition: Half.h:220
static EIGEN_CONSTEXPR const bool has_signaling_NaN
Definition: Half.h:216
static EIGEN_CONSTEXPR const bool has_denorm_loss
Definition: Half.h:218
static EIGEN_CONSTEXPR const bool is_integer
Definition: Half.h:212
static EIGEN_CONSTEXPR const int max_exponent
Definition: Half.h:231
static EIGEN_CONSTEXPR const bool is_specialized
Definition: Half.h:210
static EIGEN_CONSTEXPR const bool tinyness_before
Definition: Half.h:236
static EIGEN_CONSTEXPR Eigen::half quiet_NaN()
Definition: Half.h:244
static EIGEN_CONSTEXPR const bool traps
Definition: Half.h:233
static EIGEN_CONSTEXPR const std::float_denorm_style has_denorm
Definition: Half.h:217
static EIGEN_CONSTEXPR const std::float_round_style round_style
Definition: Half.h:219
static EIGEN_CONSTEXPR const bool is_modulo
Definition: Half.h:224
static EIGEN_CONSTEXPR const bool is_exact
Definition: Half.h:213
static EIGEN_CONSTEXPR const int min_exponent10
Definition: Half.h:230
static EIGEN_CONSTEXPR Eigen::half denorm_min()
Definition: Half.h:246
static EIGEN_CONSTEXPR Eigen::half lowest()
Definition: Half.h:239
static EIGEN_CONSTEXPR Eigen::half signaling_NaN()
Definition: Half.h:245
static EIGEN_CONSTEXPR const int digits
Definition: Half.h:225
static EIGEN_CONSTEXPR const bool is_bounded
Definition: Half.h:223
EIGEN_CONSTEXPR half(bool b)
Definition: Half.h:178
half_impl::__half_raw __half_raw
Definition: Half.h:150
EIGEN_CONSTEXPR half()
Definition: Half.h:163
EIGEN_CONSTEXPR half(const __half_raw &h)
Definition: Half.h:165
half(T val)
Definition: Half.h:181
half(float f)
Definition: Half.h:183
half(std::complex< RealScalar > c)
Definition: Half.h:189