41 #if defined(EIGEN_HAS_GPU_FP16) || defined(EIGEN_HAS_ARM64_FP16_SCALAR_ARITHMETIC) 48 #pragma push_macro("EIGEN_CONSTEXPR") 49 #undef EIGEN_CONSTEXPR 50 #define EIGEN_CONSTEXPR 53 #define F16_PACKET_FUNCTION(PACKET_F, PACKET_F16, METHOD) \ 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))); \ 85 #if !defined(EIGEN_HAS_GPU_FP16) || !defined(EIGEN_GPU_COMPILE_PHASE) 88 #if (defined(EIGEN_HAS_GPU_FP16) && !defined(EIGEN_GPU_COMPILE_PHASE)) 94 EIGEN_DEVICE_FUNC __half_raw() {}
96 EIGEN_DEVICE_FUNC EIGEN_CONSTEXPR __half_raw() : x(0) {}
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)) {
103 explicit EIGEN_DEVICE_FUNC EIGEN_CONSTEXPR __half_raw(numext::uint16_t raw) : x(raw) {}
108 #elif defined(EIGEN_HAS_HIP_FP16) 111 #elif defined(EIGEN_HAS_CUDA_FP16) 112 #if EIGEN_CUDA_SDK_VER < 90000 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;
120 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC EIGEN_CONSTEXPR __half_raw raw_uint16_to_half(numext::uint16_t x);
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);
124 struct half_base :
public __half_raw {
125 EIGEN_DEVICE_FUNC EIGEN_CONSTEXPR half_base() {}
126 EIGEN_DEVICE_FUNC EIGEN_CONSTEXPR half_base(
const __half_raw& h) : __half_raw(h) {}
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 133 EIGEN_DEVICE_FUNC EIGEN_CONSTEXPR half_base(
const __half& h) : __half_raw(*(__half_raw*)&h) {}
142 struct half :
public half_impl::half_base {
146 #if !defined(EIGEN_HAS_GPU_FP16) || !defined(EIGEN_GPU_COMPILE_PHASE) 150 typedef half_impl::__half_raw __half_raw;
151 #elif defined(EIGEN_HAS_HIP_FP16) 154 #elif defined(EIGEN_HAS_CUDA_FP16) 158 #if defined(EIGEN_CUDA_SDK_VER) && EIGEN_CUDA_SDK_VER < 90000 159 typedef half_impl::__half_raw __half_raw;
163 EIGEN_DEVICE_FUNC EIGEN_CONSTEXPR half() {}
165 EIGEN_DEVICE_FUNC EIGEN_CONSTEXPR half(
const __half_raw& h) : half_impl::half_base(h) {}
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) {}
178 explicit EIGEN_DEVICE_FUNC EIGEN_CONSTEXPR half(
bool b)
179 : half_impl::half_base(half_impl::raw_uint16_to_half(b ? 0x3c00 : 0)) {}
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)) {}
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()))) {}
192 EIGEN_DEVICE_FUNC
operator float()
const {
193 return half_impl::half_to_float(*
this);
196 #if defined(EIGEN_HAS_GPU_FP16) && !defined(EIGEN_GPU_COMPILE_PHASE) 197 EIGEN_DEVICE_FUNC
operator __half()
const {
209 struct numeric_limits<
Eigen::half> {
210 static const bool is_specialized =
true;
211 static const bool is_signed =
true;
212 static const bool is_integer =
false;
213 static const bool is_exact =
false;
214 static const bool has_infinity =
true;
215 static const bool has_quiet_NaN =
true;
216 static const bool has_signaling_NaN =
true;
217 static const float_denorm_style has_denorm = denorm_present;
218 static const bool has_denorm_loss =
false;
219 static const std::float_round_style round_style = std::round_to_nearest;
220 static const bool is_iec559 =
false;
221 static const bool is_bounded =
false;
222 static const bool is_modulo =
false;
223 static const int digits = 11;
224 static const int digits10 = 3;
225 static const int max_digits10 = 5;
226 static const int radix = 2;
227 static const int min_exponent = -13;
228 static const int min_exponent10 = -4;
229 static const int max_exponent = 16;
230 static const int max_exponent10 = 4;
231 static const bool traps =
true;
232 static const bool tinyness_before =
false;
234 static Eigen::half (min)() {
return Eigen::half_impl::raw_uint16_to_half(0x400); }
235 static Eigen::half lowest() {
return Eigen::half_impl::raw_uint16_to_half(0xfbff); }
236 static Eigen::half (max)() {
return Eigen::half_impl::raw_uint16_to_half(0x7bff); }
237 static Eigen::half epsilon() {
return Eigen::half_impl::raw_uint16_to_half(0x0800); }
238 static Eigen::half round_error() {
return Eigen::half(0.5); }
239 static Eigen::half infinity() {
return Eigen::half_impl::raw_uint16_to_half(0x7c00); }
240 static Eigen::half quiet_NaN() {
return Eigen::half_impl::raw_uint16_to_half(0x7e00); }
241 static Eigen::half signaling_NaN() {
return Eigen::half_impl::raw_uint16_to_half(0x7d00); }
242 static Eigen::half denorm_min() {
return Eigen::half_impl::raw_uint16_to_half(0x1); }
250 struct numeric_limits<const
Eigen::half> : numeric_limits<Eigen::half> {};
252 struct numeric_limits<volatile Eigen::half> : numeric_limits<Eigen::half> {};
254 struct numeric_limits<const volatile Eigen::half> : numeric_limits<Eigen::half> {};
259 namespace half_impl {
261 #if (defined(EIGEN_HAS_CUDA_FP16) && defined(EIGEN_CUDA_ARCH) && \ 262 EIGEN_CUDA_ARCH >= 530) || \ 263 (defined(EIGEN_HAS_HIP_FP16) && defined(HIP_DEVICE_COMPILE)) 267 #define EIGEN_HAS_NATIVE_FP16 275 #if defined(EIGEN_HAS_NATIVE_FP16) 276 EIGEN_STRONG_INLINE __device__ half operator + (
const half& a,
const half& b) {
277 #if defined(EIGEN_CUDA_SDK_VER) && EIGEN_CUDA_SDK_VER >= 90000 278 return __hadd(::__half(a), ::__half(b));
283 EIGEN_STRONG_INLINE __device__ half
operator * (
const half& a,
const half& b) {
286 EIGEN_STRONG_INLINE __device__ half operator - (
const half& a,
const half& b) {
289 EIGEN_STRONG_INLINE __device__ half operator / (
const half& a,
const half& b) {
290 #if defined(EIGEN_CUDA_SDK_VER) && EIGEN_CUDA_SDK_VER >= 90000 293 float num = __half2float(a);
294 float denom = __half2float(b);
295 return __float2half(num / denom);
298 EIGEN_STRONG_INLINE __device__ half operator - (
const half& a) {
301 EIGEN_STRONG_INLINE __device__ half& operator += (half& a,
const half& b) {
305 EIGEN_STRONG_INLINE __device__ half& operator *= (half& a,
const half& b) {
309 EIGEN_STRONG_INLINE __device__ half& operator -= (half& a,
const half& b) {
313 EIGEN_STRONG_INLINE __device__ half& operator /= (half& a,
const half& b) {
317 EIGEN_STRONG_INLINE __device__
bool operator == (
const half& a,
const half& b) {
320 EIGEN_STRONG_INLINE __device__
bool operator != (
const half& a,
const half& b) {
323 EIGEN_STRONG_INLINE __device__
bool operator < (
const half& a,
const half& b) {
326 EIGEN_STRONG_INLINE __device__
bool operator <= (
const half& a,
const half& b) {
329 EIGEN_STRONG_INLINE __device__
bool operator > (
const half& a,
const half& b) {
332 EIGEN_STRONG_INLINE __device__
bool operator >= (
const half& a,
const half& b) {
337 #if defined(EIGEN_HAS_ARM64_FP16_SCALAR_ARITHMETIC) 338 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half operator + (
const half& a,
const half& b) {
339 return half(vaddh_f16(a.x, b.x));
341 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half
operator * (
const half& a,
const half& b) {
342 return half(vmulh_f16(a.x, b.x));
344 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half operator - (
const half& a,
const half& b) {
345 return half(vsubh_f16(a.x, b.x));
347 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half operator / (
const half& a,
const half& b) {
348 return half(vdivh_f16(a.x, b.x));
350 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half operator - (
const half& a) {
351 return half(vnegh_f16(a.x));
353 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half& operator += (half& a,
const half& b) {
354 a = half(vaddh_f16(a.x, b.x));
357 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half& operator *= (half& a,
const half& b) {
358 a = half(vmulh_f16(a.x, b.x));
361 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half& operator -= (half& a,
const half& b) {
362 a = half(vsubh_f16(a.x, b.x));
365 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half& operator /= (half& a,
const half& b) {
366 a = half(vdivh_f16(a.x, b.x));
369 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC
bool operator == (
const half& a,
const half& b) {
370 return vceqh_f16(a.x, b.x);
372 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC
bool operator != (
const half& a,
const half& b) {
373 return !vceqh_f16(a.x, b.x);
375 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC
bool operator < (
const half& a,
const half& b) {
376 return vclth_f16(a.x, b.x);
378 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC
bool operator <= (
const half& a,
const half& b) {
379 return vcleh_f16(a.x, b.x);
381 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC
bool operator > (
const half& a,
const half& b) {
382 return vcgth_f16(a.x, b.x);
384 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC
bool operator >= (
const half& a,
const half& b) {
385 return vcgeh_f16(a.x, b.x);
390 #elif !defined(EIGEN_HAS_NATIVE_FP16) || (EIGEN_COMP_CLANG && !EIGEN_COMP_NVCC) // Emulate support for half floats 392 #if EIGEN_COMP_CLANG && defined(EIGEN_CUDACC) 394 #pragma push_macro("EIGEN_DEVICE_FUNC") 395 #undef EIGEN_DEVICE_FUNC 396 #if defined(EIGEN_HAS_CUDA_FP16) && defined(EIGEN_HAS_NATIVE_FP16) 397 #define EIGEN_DEVICE_FUNC __host__ 398 #else // both host and device need emulated ops. 399 #define EIGEN_DEVICE_FUNC __host__ __device__ 405 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half operator + (
const half& a,
const half& b) {
406 return half(
float(a) +
float(b));
408 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half
operator * (
const half& a,
const half& b) {
409 return half(
float(a) *
float(b));
411 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half operator - (
const half& a,
const half& b) {
412 return half(
float(a) -
float(b));
414 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half operator / (
const half& a,
const half& b) {
415 return half(
float(a) /
float(b));
417 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half operator - (
const half& a) {
419 result.x = a.x ^ 0x8000;
422 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half& operator += (half& a,
const half& b) {
423 a = half(
float(a) +
float(b));
426 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half& operator *= (half& a,
const half& b) {
427 a = half(
float(a) *
float(b));
430 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half& operator -= (half& a,
const half& b) {
431 a = half(
float(a) -
float(b));
434 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half& operator /= (half& a,
const half& b) {
435 a = half(
float(a) /
float(b));
438 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC
bool operator == (
const half& a,
const half& b) {
439 return numext::equal_strict(
float(a),
float(b));
441 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC
bool operator != (
const half& a,
const half& b) {
442 return numext::not_equal_strict(
float(a),
float(b));
444 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC
bool operator < (
const half& a,
const half& b) {
445 return float(a) < float(b);
447 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC
bool operator <= (
const half& a,
const half& b) {
448 return float(a) <= float(b);
450 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC
bool operator > (
const half& a,
const half& b) {
451 return float(a) > float(b);
453 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC
bool operator >= (
const half& a,
const half& b) {
454 return float(a) >= float(b);
457 #if defined(__clang__) && defined(__CUDA__) 458 #pragma pop_macro("EIGEN_DEVICE_FUNC") 460 #endif // Emulate support for half floats 464 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half operator / (
const half& a,
Index b) {
465 return half(static_cast<float>(a) / static_cast<float>(b));
468 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half operator++(half& a) {
473 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half operator--(half& a) {
478 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half operator++(half& a,
int) {
479 half original_value = a;
481 return original_value;
484 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half operator--(half& a,
int) {
485 half original_value = a;
487 return original_value;
495 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC EIGEN_CONSTEXPR __half_raw raw_uint16_to_half(numext::uint16_t x) {
502 #if defined(EIGEN_HAS_GPU_FP16) 507 return __half_raw(x);
511 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC numext::uint16_t raw_half_as_uint16(
const __half_raw& h) {
515 #if defined(EIGEN_HAS_ARM64_FP16_SCALAR_ARITHMETIC) 516 return numext::bit_cast<numext::uint16_t>(h.x);
517 #elif defined(SYCL_DEVICE_ONLY) 518 return numext::bit_cast<numext::uint16_t>(h);
529 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC __half_raw float_to_half_rtne(
float ff) {
530 #if (defined(EIGEN_HAS_CUDA_FP16) && defined(EIGEN_CUDA_ARCH) && EIGEN_CUDA_ARCH >= 300) || \ 531 (defined(EIGEN_HAS_HIP_FP16) && defined(EIGEN_HIP_DEVICE_COMPILE)) 532 __half tmp_ff = __float2half(ff);
533 return *(__half_raw*)&tmp_ff;
535 #elif defined(EIGEN_HAS_FP16_C) 537 h.x = _cvtss_sh(ff, 0);
540 #elif defined(EIGEN_HAS_ARM64_FP16_SCALAR_ARITHMETIC) 542 h.x =
static_cast<__fp16
>(ff);
546 float32_bits f; f.f = ff;
548 const float32_bits f32infty = { 255 << 23 };
549 const float32_bits f16max = { (127 + 16) << 23 };
550 const float32_bits denorm_magic = { ((127 - 15) + (23 - 10) + 1) << 23 };
551 unsigned int sign_mask = 0x80000000u;
553 o.x =
static_cast<numext::uint16_t
>(0x0u);
555 unsigned int sign = f.u & sign_mask;
563 if (f.u >= f16max.u) {
564 o.x = (f.u > f32infty.u) ? 0x7e00 : 0x7c00;
566 if (f.u < (113 << 23)) {
570 f.f += denorm_magic.f;
573 o.x =
static_cast<numext::uint16_t
>(f.u - denorm_magic.u);
575 unsigned int mant_odd = (f.u >> 13) & 1;
584 o.x =
static_cast<numext::uint16_t
>(f.u >> 13);
588 o.x |=
static_cast<numext::uint16_t
>(sign >> 16);
593 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC
float half_to_float(__half_raw h) {
594 #if (defined(EIGEN_HAS_CUDA_FP16) && defined(EIGEN_CUDA_ARCH) && EIGEN_CUDA_ARCH >= 300) || \ 595 (defined(EIGEN_HAS_HIP_FP16) && defined(EIGEN_HIP_DEVICE_COMPILE)) 596 return __half2float(h);
597 #elif defined(EIGEN_HAS_FP16_C) 598 return _cvtsh_ss(h.x);
599 #elif defined(EIGEN_HAS_ARM64_FP16_SCALAR_ARITHMETIC) 600 return static_cast<float>(h.x);
602 const float32_bits magic = { 113 << 23 };
603 const unsigned int shifted_exp = 0x7c00 << 13;
606 o.u = (h.x & 0x7fff) << 13;
607 unsigned int exp = shifted_exp & o.u;
608 o.u += (127 - 15) << 23;
611 if (exp == shifted_exp) {
612 o.u += (128 - 16) << 23;
613 }
else if (exp == 0) {
618 o.u |= (h.x & 0x8000) << 16;
625 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC bool (
isinf)(
const half& a) {
626 #ifdef EIGEN_HAS_ARM64_FP16_SCALAR_ARITHMETIC 627 return (numext::bit_cast<numext::uint16_t>(a.x) & 0x7fff) == 0x7c00;
629 return (a.x & 0x7fff) == 0x7c00;
632 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC bool (
isnan)(
const half& a) {
633 #if (defined(EIGEN_HAS_CUDA_FP16) && defined(EIGEN_CUDA_ARCH) && EIGEN_CUDA_ARCH >= 530) || \ 634 (defined(EIGEN_HAS_HIP_FP16) && defined(EIGEN_HIP_DEVICE_COMPILE)) 636 #elif defined(EIGEN_HAS_ARM64_FP16_SCALAR_ARITHMETIC) 637 return (numext::bit_cast<numext::uint16_t>(a.x) & 0x7fff) > 0x7c00;
639 return (a.x & 0x7fff) > 0x7c00;
642 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC bool (
isfinite)(
const half& a) {
643 return !(
isinf EIGEN_NOT_A_MACRO (a)) && !(
isnan EIGEN_NOT_A_MACRO (a));
646 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half
abs(
const half& a) {
647 #if defined(EIGEN_HAS_ARM64_FP16_SCALAR_ARITHMETIC) 648 return half(vabsh_f16(a.x));
651 result.x = a.x & 0x7FFF;
655 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half
exp(
const half& a) {
656 #if (EIGEN_CUDA_SDK_VER >= 80000 && defined EIGEN_CUDA_ARCH && EIGEN_CUDA_ARCH >= 530) || \ 657 defined(EIGEN_HIP_DEVICE_COMPILE) 658 return half(hexp(a));
660 return half(::expf(
float(a)));
663 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half
expm1(
const half& a) {
664 return half(numext::expm1(
float(a)));
666 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half
log(
const half& a) {
667 #if (defined(EIGEN_HAS_CUDA_FP16) && EIGEN_CUDA_SDK_VER >= 80000 && defined(EIGEN_CUDA_ARCH) && EIGEN_CUDA_ARCH >= 530) || \ 668 (defined(EIGEN_HAS_HIP_FP16) && defined(EIGEN_HIP_DEVICE_COMPILE)) 669 return half(::hlog(a));
671 return half(::logf(
float(a)));
674 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half
log1p(
const half& a) {
675 return half(numext::log1p(
float(a)));
677 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half
log10(
const half& a) {
678 return half(::log10f(
float(a)));
680 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half
log2(
const half& a) {
681 return half(static_cast<float>(EIGEN_LOG2E) * ::logf(
float(a)));
684 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half
sqrt(
const half& a) {
685 #if (EIGEN_CUDA_SDK_VER >= 80000 && defined EIGEN_CUDA_ARCH && EIGEN_CUDA_ARCH >= 530) || \ 686 defined(EIGEN_HIP_DEVICE_COMPILE) 687 return half(hsqrt(a));
689 return half(::sqrtf(
float(a)));
692 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half pow(
const half& a,
const half& b) {
693 return half(::powf(
float(a),
float(b)));
695 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half
sin(
const half& a) {
696 return half(::sinf(
float(a)));
698 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half
cos(
const half& a) {
699 return half(::cosf(
float(a)));
701 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half
tan(
const half& a) {
702 return half(::tanf(
float(a)));
704 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half
tanh(
const half& a) {
705 return half(::tanhf(
float(a)));
707 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half
asin(
const half& a) {
708 return half(::asinf(
float(a)));
710 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half
acos(
const half& a) {
711 return half(::acosf(
float(a)));
713 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half
floor(
const half& a) {
714 #if (EIGEN_CUDA_SDK_VER >= 80000 && defined EIGEN_CUDA_ARCH && EIGEN_CUDA_ARCH >= 300) || \ 715 defined(EIGEN_HIP_DEVICE_COMPILE) 716 return half(hfloor(a));
718 return half(::floorf(
float(a)));
721 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half
ceil(
const half& a) {
722 #if (EIGEN_CUDA_SDK_VER >= 80000 && defined EIGEN_CUDA_ARCH && EIGEN_CUDA_ARCH >= 300) || \ 723 defined(EIGEN_HIP_DEVICE_COMPILE) 724 return half(hceil(a));
726 return half(::ceilf(
float(a)));
729 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half
rint(
const half& a) {
730 return half(::rintf(
float(a)));
732 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half
round(
const half& a) {
733 return half(::roundf(
float(a)));
735 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half fmod(
const half& a,
const half& b) {
736 return half(::fmodf(
float(a),
float(b)));
739 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half (min)(
const half& a,
const half& b) {
740 #if (defined(EIGEN_HAS_CUDA_FP16) && defined(EIGEN_CUDA_ARCH) && EIGEN_CUDA_ARCH >= 530) || \ 741 (defined(EIGEN_HAS_HIP_FP16) && defined(EIGEN_HIP_DEVICE_COMPILE)) 742 return __hlt(b, a) ? b : a;
744 const float f1 =
static_cast<float>(a);
745 const float f2 =
static_cast<float>(b);
746 return f2 < f1 ? b : a;
749 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half (max)(
const half& a,
const half& b) {
750 #if (defined(EIGEN_HAS_CUDA_FP16) && defined(EIGEN_CUDA_ARCH) && EIGEN_CUDA_ARCH >= 530) || \ 751 (defined(EIGEN_HAS_HIP_FP16) && defined(EIGEN_HIP_DEVICE_COMPILE)) 752 return __hlt(a, b) ? b : a;
754 const float f1 =
static_cast<float>(a);
755 const float f2 =
static_cast<float>(b);
756 return f1 < f2 ? b : a;
761 EIGEN_ALWAYS_INLINE std::ostream& operator << (std::ostream& os,
const half& v) {
762 os << static_cast<float>(v);
775 struct random_default_impl<half, false, false>
777 static inline half run(
const half& x,
const half& y)
779 return x + (y-x) * half(
float(std::rand()) / float(RAND_MAX));
781 static inline half run()
783 return run(half(-1.f), half(1.f));
787 template<>
struct is_arithmetic<half> {
enum { value =
true }; };
791 template<>
struct NumTraits<
Eigen::half>
792 : GenericNumTraits<Eigen::half>
798 RequireInitialization =
false 801 EIGEN_DEVICE_FUNC EIGEN_CONSTEXPR
static EIGEN_STRONG_INLINE Eigen::half epsilon() {
802 return half_impl::raw_uint16_to_half(0x0800);
804 EIGEN_DEVICE_FUNC EIGEN_CONSTEXPR
static EIGEN_STRONG_INLINE Eigen::half dummy_precision() {
805 return half_impl::raw_uint16_to_half(0x211f);
807 EIGEN_DEVICE_FUNC EIGEN_CONSTEXPR
static EIGEN_STRONG_INLINE Eigen::half highest() {
808 return half_impl::raw_uint16_to_half(0x7bff);
810 EIGEN_DEVICE_FUNC EIGEN_CONSTEXPR
static EIGEN_STRONG_INLINE Eigen::half lowest() {
811 return half_impl::raw_uint16_to_half(0xfbff);
813 EIGEN_DEVICE_FUNC EIGEN_CONSTEXPR
static EIGEN_STRONG_INLINE Eigen::half infinity() {
814 return half_impl::raw_uint16_to_half(0x7c00);
816 EIGEN_DEVICE_FUNC EIGEN_CONSTEXPR
static EIGEN_STRONG_INLINE Eigen::half quiet_NaN() {
817 return half_impl::raw_uint16_to_half(0x7e00);
823 #if defined(EIGEN_HAS_GPU_FP16) || defined(EIGEN_HAS_ARM64_FP16_SCALAR_ARITHMETIC) 824 #pragma pop_macro("EIGEN_CONSTEXPR") 830 #if defined(EIGEN_GPU_COMPILE_PHASE) 833 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE bool(
isnan)(
const Eigen::half& h) {
834 return (half_impl::isnan)(h);
838 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE bool(
isinf)(
const Eigen::half& h) {
839 return (half_impl::isinf)(h);
843 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE bool(
isfinite)(
const Eigen::half& h) {
844 return (half_impl::isfinite)(h);
850 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC Eigen::half bit_cast<Eigen::half, uint16_t>(
const uint16_t& src) {
851 return Eigen::half(Eigen::half_impl::raw_uint16_to_half(src));
855 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC uint16_t bit_cast<uint16_t, Eigen::half>(
const Eigen::half& src) {
856 return Eigen::half_impl::raw_half_as_uint16(src);
873 #if (defined(EIGEN_CUDACC) && (!defined(EIGEN_CUDA_ARCH) || EIGEN_CUDA_ARCH >= 300)) \ 874 || defined(EIGEN_HIPCC) 876 #if defined(EIGEN_HAS_CUDA_FP16) && EIGEN_CUDA_SDK_VER >= 90000 878 __device__ EIGEN_STRONG_INLINE Eigen::half __shfl_sync(
unsigned mask, Eigen::half var,
int srcLane,
int width=warpSize) {
879 const __half h = var;
880 return static_cast<Eigen::half
>(__shfl_sync(mask, h, srcLane, width));
883 __device__ EIGEN_STRONG_INLINE Eigen::half __shfl_up_sync(
unsigned mask, Eigen::half var,
unsigned int delta,
int width=warpSize) {
884 const __half h = var;
885 return static_cast<Eigen::half
>(__shfl_up_sync(mask, h, delta, width));
888 __device__ EIGEN_STRONG_INLINE Eigen::half __shfl_down_sync(
unsigned mask, Eigen::half var,
unsigned int delta,
int width=warpSize) {
889 const __half h = var;
890 return static_cast<Eigen::half
>(__shfl_down_sync(mask, h, delta, width));
893 __device__ EIGEN_STRONG_INLINE Eigen::half __shfl_xor_sync(
unsigned mask, Eigen::half var,
int laneMask,
int width=warpSize) {
894 const __half h = var;
895 return static_cast<Eigen::half
>(__shfl_xor_sync(mask, h, laneMask, width));
898 #else // HIP or CUDA SDK < 9.0 900 __device__ EIGEN_STRONG_INLINE Eigen::half __shfl(Eigen::half var,
int srcLane,
int width=warpSize) {
901 const int ivar =
static_cast<int>(Eigen::numext::bit_cast<Eigen::numext::uint16_t>(var));
902 return Eigen::numext::bit_cast<Eigen::half>(
static_cast<Eigen::numext::uint16_t
>(__shfl(ivar, srcLane, width)));
905 __device__ EIGEN_STRONG_INLINE Eigen::half __shfl_up(Eigen::half var,
unsigned int delta,
int width=warpSize) {
906 const int ivar =
static_cast<int>(Eigen::numext::bit_cast<Eigen::numext::uint16_t>(var));
907 return Eigen::numext::bit_cast<Eigen::half>(
static_cast<Eigen::numext::uint16_t
>(__shfl_up(ivar, delta, width)));
910 __device__ EIGEN_STRONG_INLINE Eigen::half __shfl_down(Eigen::half var,
unsigned int delta,
int width=warpSize) {
911 const int ivar =
static_cast<int>(Eigen::numext::bit_cast<Eigen::numext::uint16_t>(var));
912 return Eigen::numext::bit_cast<Eigen::half>(
static_cast<Eigen::numext::uint16_t
>(__shfl_down(ivar, delta, width)));
915 __device__ EIGEN_STRONG_INLINE Eigen::half __shfl_xor(Eigen::half var,
int laneMask,
int width=warpSize) {
916 const int ivar =
static_cast<int>(Eigen::numext::bit_cast<Eigen::numext::uint16_t>(var));
917 return Eigen::numext::bit_cast<Eigen::half>(
static_cast<Eigen::numext::uint16_t
>(__shfl_xor(ivar, laneMask, width)));
920 #endif // HIP vs CUDA 924 #if (defined(EIGEN_CUDACC) && (!defined(EIGEN_CUDA_ARCH) || EIGEN_CUDA_ARCH >= 350)) \ 925 || defined(EIGEN_HIPCC) 926 EIGEN_STRONG_INLINE __device__ Eigen::half __ldg(
const Eigen::half* ptr) {
927 return Eigen::half_impl::raw_uint16_to_half(__ldg(reinterpret_cast<const Eigen::numext::uint16_t*>(ptr)));
931 #if EIGEN_HAS_STD_HASH 934 struct hash<
Eigen::half> {
935 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE std::size_t operator()(
const Eigen::half& a)
const {
936 return static_cast<std::size_t
>(Eigen::numext::bit_cast<Eigen::numext::uint16_t>(a));
942 #endif // EIGEN_HALF_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: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)
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_log2_op< typename Derived::Scalar >, const Derived > log2(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)