aboutsummaryrefslogtreecommitdiffhomepage
path: root/Eigen
diff options
context:
space:
mode:
authorGravatar Gael Guennebaud <g.gael@free.fr>2016-08-04 00:00:43 +0200
committerGravatar Gael Guennebaud <g.gael@free.fr>2016-08-04 00:00:43 +0200
commit17b9a55d98cf9fba419d1b43c8723bcbeb7e58f0 (patch)
tree69bd5e52adbfff2c19ffac482a0e78c48093a0bb /Eigen
parentca2cee27397e04357756c2d2b03751899bc0de00 (diff)
Move Eigen::half_impl::half to Eigen::half while preserving the free functions to the Eigen::half_impl namespace together with ADL
Diffstat (limited to 'Eigen')
-rw-r--r--Eigen/src/Core/arch/CUDA/Half.h72
1 files changed, 44 insertions, 28 deletions
diff --git a/Eigen/src/Core/arch/CUDA/Half.h b/Eigen/src/Core/arch/CUDA/Half.h
index 9df7c2f69..6ae2c53c5 100644
--- a/Eigen/src/Core/arch/CUDA/Half.h
+++ b/Eigen/src/Core/arch/CUDA/Half.h
@@ -45,6 +45,8 @@
namespace Eigen {
+struct half;
+
namespace half_impl {
#if !defined(EIGEN_HAS_CUDA_FP16)
@@ -62,60 +64,72 @@ EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC __half raw_uint16_to_half(unsigned short x
EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC __half float_to_half_rtne(float ff);
EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC float half_to_float(__half h);
+struct half_base : public __half {
+ explicit EIGEN_DEVICE_FUNC half_base(unsigned short raw) : __half(raw) {}
+ EIGEN_DEVICE_FUNC half_base() {}
+ EIGEN_DEVICE_FUNC half_base(const __half& h) : __half(h) {}
+};
+
+} // namespace half_impl
+
// Class definition.
-struct half : public __half {
+struct half : public half_impl::half_base {
+ #if !defined(EIGEN_HAS_CUDA_FP16)
+ typedef half_impl::__half __half;
+ #endif
+
EIGEN_DEVICE_FUNC half() {}
- EIGEN_DEVICE_FUNC half(const __half& h) : __half(h) {}
- EIGEN_DEVICE_FUNC half(const half& h) : __half(h) {}
+ EIGEN_DEVICE_FUNC half(const __half& h) : half_impl::half_base(h) {}
+ EIGEN_DEVICE_FUNC half(const half& h) : half_impl::half_base(h) {}
explicit EIGEN_DEVICE_FUNC half(bool b)
- : __half(raw_uint16_to_half(b ? 0x3c00 : 0)) {}
+ : half_impl::half_base(half_impl::raw_uint16_to_half(b ? 0x3c00 : 0)) {}
template<class T>
explicit EIGEN_DEVICE_FUNC half(const T& val)
- : __half(float_to_half_rtne(static_cast<float>(val))) {}
+ : half_impl::half_base(half_impl::float_to_half_rtne(static_cast<float>(val))) {}
explicit EIGEN_DEVICE_FUNC half(float f)
- : __half(float_to_half_rtne(f)) {}
+ : half_impl::half_base(half_impl::float_to_half_rtne(f)) {}
EIGEN_DEVICE_FUNC EIGEN_EXPLICIT_CAST(bool) const {
// +0.0 and -0.0 become false, everything else becomes true.
return (x & 0x7fff) != 0;
}
EIGEN_DEVICE_FUNC EIGEN_EXPLICIT_CAST(signed char) const {
- return static_cast<signed char>(half_to_float(*this));
+ return static_cast<signed char>(half_impl::half_to_float(*this));
}
EIGEN_DEVICE_FUNC EIGEN_EXPLICIT_CAST(unsigned char) const {
- return static_cast<unsigned char>(half_to_float(*this));
+ return static_cast<unsigned char>(half_impl::half_to_float(*this));
}
EIGEN_DEVICE_FUNC EIGEN_EXPLICIT_CAST(short) const {
- return static_cast<short>(half_to_float(*this));
+ return static_cast<short>(half_impl::half_to_float(*this));
}
EIGEN_DEVICE_FUNC EIGEN_EXPLICIT_CAST(unsigned short) const {
- return static_cast<unsigned short>(half_to_float(*this));
+ return static_cast<unsigned short>(half_impl::half_to_float(*this));
}
EIGEN_DEVICE_FUNC EIGEN_EXPLICIT_CAST(int) const {
- return static_cast<int>(half_to_float(*this));
+ return static_cast<int>(half_impl::half_to_float(*this));
}
EIGEN_DEVICE_FUNC EIGEN_EXPLICIT_CAST(unsigned int) const {
- return static_cast<unsigned int>(half_to_float(*this));
+ return static_cast<unsigned int>(half_impl::half_to_float(*this));
}
EIGEN_DEVICE_FUNC EIGEN_EXPLICIT_CAST(long) const {
- return static_cast<long>(half_to_float(*this));
+ return static_cast<long>(half_impl::half_to_float(*this));
}
EIGEN_DEVICE_FUNC EIGEN_EXPLICIT_CAST(unsigned long) const {
- return static_cast<unsigned long>(half_to_float(*this));
+ return static_cast<unsigned long>(half_impl::half_to_float(*this));
}
EIGEN_DEVICE_FUNC EIGEN_EXPLICIT_CAST(long long) const {
- return static_cast<long long>(half_to_float(*this));
+ return static_cast<long long>(half_impl::half_to_float(*this));
}
EIGEN_DEVICE_FUNC EIGEN_EXPLICIT_CAST(unsigned long long) const {
return static_cast<unsigned long long>(half_to_float(*this));
}
EIGEN_DEVICE_FUNC EIGEN_EXPLICIT_CAST(float) const {
- return half_to_float(*this);
+ return half_impl::half_to_float(*this);
}
EIGEN_DEVICE_FUNC EIGEN_EXPLICIT_CAST(double) const {
- return static_cast<double>(half_to_float(*this));
+ return static_cast<double>(half_impl::half_to_float(*this));
}
EIGEN_DEVICE_FUNC half& operator=(const half& other) {
@@ -124,6 +138,8 @@ struct half : public __half {
}
};
+namespace half_impl {
+
#if defined(EIGEN_HAS_CUDA_FP16) && defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 530
// Intrinsics for native fp16 support. Note that on current hardware,
@@ -430,12 +446,12 @@ EIGEN_ALWAYS_INLINE std::ostream& operator << (std::ostream& os, const half& v)
} // end namespace half_impl
// import Eigen::half_impl::half into Eigen namespace
-using half_impl::half;
+// using half_impl::half;
namespace internal {
template<>
-struct random_default_impl<half_impl::half, false, false>
+struct random_default_impl<half, false, false>
{
static inline half run(const half& x, const half& y)
{
@@ -447,27 +463,27 @@ struct random_default_impl<half_impl::half, false, false>
}
};
-template<> struct is_arithmetic<half_impl::half> { enum { value = true }; };
+template<> struct is_arithmetic<half> { enum { value = true }; };
} // end namespace internal
-template<> struct NumTraits<Eigen::half_impl::half>
- : GenericNumTraits<Eigen::half_impl::half>
+template<> struct NumTraits<Eigen::half>
+ : GenericNumTraits<Eigen::half>
{
- EIGEN_DEVICE_FUNC static EIGEN_STRONG_INLINE Eigen::half_impl::half epsilon() {
+ EIGEN_DEVICE_FUNC static EIGEN_STRONG_INLINE Eigen::half epsilon() {
return half_impl::raw_uint16_to_half(0x0800);
}
- EIGEN_DEVICE_FUNC static EIGEN_STRONG_INLINE Eigen::half_impl::half dummy_precision() { return half_impl::half(1e-2f); }
- EIGEN_DEVICE_FUNC static EIGEN_STRONG_INLINE Eigen::half_impl::half highest() {
+ EIGEN_DEVICE_FUNC static EIGEN_STRONG_INLINE Eigen::half dummy_precision() { return Eigen::half(1e-2f); }
+ EIGEN_DEVICE_FUNC static EIGEN_STRONG_INLINE Eigen::half highest() {
return half_impl::raw_uint16_to_half(0x7bff);
}
- EIGEN_DEVICE_FUNC static EIGEN_STRONG_INLINE Eigen::half_impl::half lowest() {
+ EIGEN_DEVICE_FUNC static EIGEN_STRONG_INLINE Eigen::half lowest() {
return half_impl::raw_uint16_to_half(0xfbff);
}
- EIGEN_DEVICE_FUNC static EIGEN_STRONG_INLINE Eigen::half_impl::half infinity() {
+ EIGEN_DEVICE_FUNC static EIGEN_STRONG_INLINE Eigen::half infinity() {
return half_impl::raw_uint16_to_half(0x7c00);
}
- EIGEN_DEVICE_FUNC static EIGEN_STRONG_INLINE Eigen::half_impl::half quiet_NaN() {
+ EIGEN_DEVICE_FUNC static EIGEN_STRONG_INLINE Eigen::half quiet_NaN() {
return half_impl::raw_uint16_to_half(0x7c01);
}
};