39 #include "../../InternalHeaderCheck.h"
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))
98 #if defined(EIGEN_HAS_ARM64_FP16_SCALAR_ARITHMETIC)
108 #elif defined(EIGEN_HAS_HIP_FP16)
111 #elif defined(EIGEN_HAS_CUDA_FP16)
112 #if EIGEN_CUDA_SDK_VER < 90000
116 #elif defined(SYCL_DEVICE_ONLY)
128 #if defined(EIGEN_HAS_GPU_FP16)
129 #if defined(EIGEN_HAS_HIP_FP16)
131 #elif defined(EIGEN_HAS_CUDA_FP16)
132 #if EIGEN_CUDA_SDK_VER >= 90000
146 #if !defined(EIGEN_HAS_GPU_FP16) || !defined(EIGEN_GPU_COMPILE_PHASE)
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
167 #if defined(EIGEN_HAS_GPU_FP16)
168 #if defined(EIGEN_HAS_HIP_FP16)
170 #elif defined(EIGEN_HAS_CUDA_FP16)
171 #if defined(EIGEN_CUDA_SDK_VER) && EIGEN_CUDA_SDK_VER >= 90000
188 template<
typename RealScalar>
196 #if defined(EIGEN_HAS_GPU_FP16) && !defined(EIGEN_GPU_COMPILE_PHASE)
207 namespace half_impl {
208 template <
typename =
void>
306 class numeric_limits<const
Eigen::half> :
public numeric_limits<Eigen::half> {};
308 class numeric_limits<volatile
Eigen::half> :
public numeric_limits<Eigen::half> {};
310 class numeric_limits<const volatile
Eigen::half> :
public numeric_limits<Eigen::half> {};
315 namespace half_impl {
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))
323 #define EIGEN_HAS_NATIVE_FP16
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));
339 EIGEN_STRONG_INLINE __device__ half
operator * (
const half&
a,
const half&
b) {
342 EIGEN_STRONG_INLINE __device__ half
operator - (
const half&
a,
const half&
b) {
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
349 float num = __half2float(
a);
350 float denom = __half2float(
b);
351 return __float2half(num / denom);
354 EIGEN_STRONG_INLINE __device__ half
operator - (
const half&
a) {
357 EIGEN_STRONG_INLINE __device__ half&
operator += (half&
a,
const half&
b) {
361 EIGEN_STRONG_INLINE __device__ half&
operator *= (half&
a,
const half&
b) {
365 EIGEN_STRONG_INLINE __device__ half&
operator -= (half&
a,
const half&
b) {
369 EIGEN_STRONG_INLINE __device__ half&
operator /= (half&
a,
const half&
b) {
373 EIGEN_STRONG_INLINE __device__
bool operator == (
const half&
a,
const half&
b) {
376 EIGEN_STRONG_INLINE __device__
bool operator != (
const half&
a,
const half&
b) {
379 EIGEN_STRONG_INLINE __device__
bool operator < (
const half&
a,
const half&
b) {
382 EIGEN_STRONG_INLINE __device__
bool operator <= (
const half&
a,
const half&
b) {
385 EIGEN_STRONG_INLINE __device__
bool operator > (
const half&
a,
const half&
b) {
388 EIGEN_STRONG_INLINE __device__
bool operator >= (
const half&
a,
const half&
b) {
393 #if defined(EIGEN_HAS_ARM64_FP16_SCALAR_ARITHMETIC) && !defined(EIGEN_GPU_COMPILE_PHASE)
395 return half(vaddh_f16(
a.x,
b.x));
398 return half(vmulh_f16(
a.x,
b.x));
401 return half(vsubh_f16(
a.x,
b.x));
404 return half(vdivh_f16(
a.x,
b.x));
407 return half(vnegh_f16(
a.x));
410 a = half(vaddh_f16(
a.x,
b.x));
414 a = half(vmulh_f16(
a.x,
b.x));
418 a = half(vsubh_f16(
a.x,
b.x));
422 a = half(vdivh_f16(
a.x,
b.x));
426 return vceqh_f16(
a.x,
b.x);
429 return !vceqh_f16(
a.x,
b.x);
432 return vclth_f16(
a.x,
b.x);
435 return vcleh_f16(
a.x,
b.x);
438 return vcgth_f16(
a.x,
b.x);
441 return vcgeh_f16(
a.x,
b.x);
446 #elif !defined(EIGEN_HAS_NATIVE_FP16) || (EIGEN_COMP_CLANG && !EIGEN_COMP_NVCC)
448 #if EIGEN_COMP_CLANG && defined(EIGEN_GPUCC)
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__
455 #define EIGEN_DEVICE_FUNC __host__ __device__
462 return half(
float(
a) +
float(
b));
465 return half(
float(
a) *
float(
b));
468 return half(
float(
a) -
float(
b));
471 return half(
float(
a) /
float(
b));
475 result.
x =
a.x ^ 0x8000;
501 return float(
a) < float(
b);
504 return float(
a) <= float(
b);
507 return float(
a) > float(
b);
510 return float(
a) >= float(
b);
513 #if EIGEN_COMP_CLANG && defined(EIGEN_GPUCC)
514 #pragma pop_macro("EIGEN_DEVICE_FUNC")
521 return half(
static_cast<float>(
a) /
static_cast<float>(
b));
535 half original_value =
a;
537 return original_value;
541 half original_value =
a;
543 return original_value;
558 #if defined(EIGEN_HAS_GPU_FP16)
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);
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);
591 #elif defined(EIGEN_HAS_FP16_C)
595 h.
x =_mm_extract_epi16(_mm_cvtps_ph(_mm_set_ss(ff), 0), 0);
597 h.
x = _cvtss_sh(ff, 0);
601 #elif defined(EIGEN_HAS_ARM64_FP16_SCALAR_ARITHMETIC)
603 h.
x =
static_cast<__fp16
>(ff);
611 const float32_bits denorm_magic = { ((127 - 15) + (23 - 10) + 1) << 23 };
612 unsigned int sign_mask = 0x80000000u;
616 unsigned int sign = f.
u & sign_mask;
624 if (f.
u >= f16max.
u) {
625 o.
x = (f.
u > f32infty.
u) ? 0x7e00 : 0x7c00;
627 if (f.
u < (113 << 23)) {
631 f.
f += denorm_magic.
f;
636 unsigned int mant_odd = (f.
u >> 13) & 1;
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)
661 return _mm_cvtss_f32(_mm_cvtph_ps(_mm_set1_epi16(h.
x)));
663 return _cvtsh_ss(h.
x);
665 #elif defined(EIGEN_HAS_ARM64_FP16_SCALAR_ARITHMETIC)
666 return static_cast<float>(h.
x);
669 const unsigned int shifted_exp = 0x7c00 << 13;
672 o.
u = (h.
x & 0x7fff) << 13;
673 unsigned int exp = shifted_exp & o.
u;
674 o.
u += (127 - 15) << 23;
677 if (
exp == shifted_exp) {
678 o.
u += (128 - 16) << 23;
679 }
else if (
exp == 0) {
684 o.
u |= (h.
x & 0x8000) << 16;
692 #ifdef EIGEN_HAS_ARM64_FP16_SCALAR_ARITHMETIC
693 return (numext::bit_cast<numext::uint16_t>(
a.x) & 0x7fff) == 0x7c00;
695 return (
a.x & 0x7fff) == 0x7c00;
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))
702 #elif defined(EIGEN_HAS_ARM64_FP16_SCALAR_ARITHMETIC)
703 return (numext::bit_cast<numext::uint16_t>(
a.x) & 0x7fff) > 0x7c00;
705 return (
a.x & 0x7fff) > 0x7c00;
713 #if defined(EIGEN_HAS_ARM64_FP16_SCALAR_ARITHMETIC)
714 return half(vabsh_f16(
a.x));
717 result.
x =
a.x & 0x7FFF;
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));
726 return half(::expf(
float(
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));
737 return half(::logf(
float(
a)));
744 return half(::log10f(
float(
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));
755 return half(::sqrtf(
float(
a)));
759 return half(::powf(
float(
a),
float(
b)));
762 return half(::atan2f(
float(
a),
float(
b)));
765 return half(::sinf(
float(
a)));
768 return half(::cosf(
float(
a)));
771 return half(::tanf(
float(
a)));
774 return half(::tanhf(
float(
a)));
777 return half(::asinf(
float(
a)));
780 return half(::acosf(
float(
a)));
783 return half(::atanf(
float(
a)));
786 return half(::atanhf(
float(
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));
793 return half(::floorf(
float(
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));
801 return half(::ceilf(
float(
a)));
805 return half(::rintf(
float(
a)));
808 return half(::roundf(
float(
a)));
811 return half(::fmodf(
float(
a),
float(
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;
819 const float f1 =
static_cast<float>(
a);
820 const float f2 =
static_cast<float>(
b);
821 return f2 < f1 ?
b :
a;
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;
829 const float f1 =
static_cast<float>(
a);
830 const float f2 =
static_cast<float>(
b);
831 return f1 < f2 ?
b :
a;
837 os << static_cast<float>(
v);
850 struct random_default_impl<half, false, false>
852 static inline half run(
const half&
x,
const half&
y)
854 return x + (
y-
x) * half(
float(std::rand()) / float(RAND_MAX));
856 static inline half run()
858 return run(half(-1.f), half(1.f));
862 template<>
struct is_arithmetic<half> {
enum { value =
true }; };
898 #if defined(EIGEN_HAS_GPU_FP16) || defined(EIGEN_HAS_ARM64_FP16_SCALAR_ARITHMETIC)
899 #pragma pop_macro("EIGEN_CONSTEXPR")
905 #if defined(EIGEN_GPU_COMPILE_PHASE)
948 #if (defined(EIGEN_CUDACC) && (!defined(EIGEN_CUDA_ARCH) || EIGEN_CUDA_ARCH >= 300)) \
949 || defined(EIGEN_HIPCC)
951 #if defined(EIGEN_HAS_CUDA_FP16) && EIGEN_CUDA_SDK_VER >= 90000
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));
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));
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));
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));
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)));
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)));
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)));
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)));
999 #if (defined(EIGEN_CUDACC) && (!defined(EIGEN_CUDA_ARCH) || EIGEN_CUDA_ARCH >= 350)) \
1000 || defined(EIGEN_HIPCC)
1006 #if EIGEN_HAS_STD_HASH
1009 struct hash<
Eigen::half> {
1011 return static_cast<std::size_t
>(Eigen::numext::bit_cast<Eigen::numext::uint16_t>(
a));
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);
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));
1041 return half(
static_cast<float>(
a));
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);
1054 return static_cast<float>(
a);
const Log1pReturnType log1p() const
const Expm1ReturnType expm1() const
Array< int, Dynamic, 1 > v
IndexedView_or_Block operator()(const RowIndices &rowIndices, const ColIndices &colIndices)
#define EIGEN_ALWAYS_INLINE
#define EIGEN_DEVICE_FUNC
#define EIGEN_NOT_A_MACRO
half() max(const half &a, const half &b)
EIGEN_ALWAYS_INLINE std::ostream & operator<<(std::ostream &os, const half &v)
half atanh(const half &a)
half & operator/=(half &a, const half &b)
half round(const half &a)
bool operator>(const half &a, const half &b)
float half_to_float(__half_raw h)
bool() isnan(const half &a)
half floor(const half &a)
half operator*(const half &a, const half &b)
bool operator>=(const half &a, const half &b)
half() min(const half &a, const half &b)
numext::uint16_t raw_half_as_uint16(const __half_raw &h)
half operator-(const half &a, const half &b)
half atan2(const half &a, const half &b)
half log10(const half &a)
bool operator<=(const half &a, const half &b)
half & operator*=(half &a, const half &b)
half & operator-=(half &a, const half &b)
bool() isinf(const half &a)
half fmod(const half &a, const half &b)
half & operator+=(half &a, const half &b)
half log1p(const half &a)
EIGEN_CONSTEXPR __half_raw raw_uint16_to_half(numext::uint16_t x)
bool operator!=(const half &a, const half &b)
__half_raw float_to_half_rtne(float ff)
half operator+(const half &a, const half &b)
bool operator<(const half &a, const half &b)
half operator/(const half &a, const half &b)
bool() isfinite(const half &a)
half pow(const half &a, const half &b)
half expm1(const half &a)
bool operator==(const half &a, const half &b)
bool equal_strict(const X &x, const Y &y)
bool not_equal_strict(const X &x, const Y &y)
EIGEN_ALWAYS_INLINE bool() isinf(const Eigen::bfloat16 &h)
Tgt bit_cast(const Src &src)
EIGEN_ALWAYS_INLINE bool() isnan(const Eigen::bfloat16 &h)
EIGEN_ALWAYS_INLINE bool() isfinite(const Eigen::bfloat16 &h)
EIGEN_DEFAULT_DENSE_INDEX_TYPE Index
The Index type as used for the API.
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)
static EIGEN_CONSTEXPR Eigen::half infinity()
static EIGEN_CONSTEXPR Eigen::half dummy_precision()
static EIGEN_CONSTEXPR Eigen::half epsilon()
static EIGEN_CONSTEXPR Eigen::half lowest()
static EIGEN_CONSTEXPR Eigen::half highest()
static EIGEN_CONSTEXPR Eigen::half quiet_NaN()
Holds information about the various numeric (i.e. scalar) types allowed by Eigen.
EIGEN_CONSTEXPR __half_raw()
EIGEN_CONSTEXPR __half_raw(numext::uint16_t raw)
EIGEN_CONSTEXPR half_base()
EIGEN_CONSTEXPR half_base(const __half_raw &h)
static EIGEN_CONSTEXPR const bool has_infinity
static EIGEN_CONSTEXPR Eigen::half round_error()
static EIGEN_CONSTEXPR const int radix
static EIGEN_CONSTEXPR const int digits10
static EIGEN_CONSTEXPR Eigen::half() min()
static EIGEN_CONSTEXPR Eigen::half infinity()
static EIGEN_CONSTEXPR const int min_exponent
static EIGEN_CONSTEXPR const int max_exponent10
static EIGEN_CONSTEXPR Eigen::half epsilon()
static EIGEN_CONSTEXPR Eigen::half() max()
static EIGEN_CONSTEXPR const bool has_quiet_NaN
static EIGEN_CONSTEXPR const bool is_signed
static EIGEN_CONSTEXPR const int max_digits10
static EIGEN_CONSTEXPR const bool is_iec559
static EIGEN_CONSTEXPR const bool has_signaling_NaN
static EIGEN_CONSTEXPR const bool has_denorm_loss
static EIGEN_CONSTEXPR const bool is_integer
static EIGEN_CONSTEXPR const int max_exponent
static EIGEN_CONSTEXPR const bool is_specialized
static EIGEN_CONSTEXPR const bool tinyness_before
static EIGEN_CONSTEXPR Eigen::half quiet_NaN()
static EIGEN_CONSTEXPR const bool traps
static EIGEN_CONSTEXPR const std::float_denorm_style has_denorm
static EIGEN_CONSTEXPR const std::float_round_style round_style
static EIGEN_CONSTEXPR const bool is_modulo
static EIGEN_CONSTEXPR const bool is_exact
static EIGEN_CONSTEXPR const int min_exponent10
static EIGEN_CONSTEXPR Eigen::half denorm_min()
static EIGEN_CONSTEXPR Eigen::half lowest()
static EIGEN_CONSTEXPR Eigen::half signaling_NaN()
static EIGEN_CONSTEXPR const int digits
static EIGEN_CONSTEXPR const bool is_bounded
EIGEN_CONSTEXPR half(bool b)
half_impl::__half_raw __half_raw
EIGEN_CONSTEXPR half(const __half_raw &h)
half(std::complex< RealScalar > c)