diff options
author | tillahoffmann <tillahoffmann@gmail.com> | 2016-04-01 14:36:15 +0100 |
---|---|---|
committer | tillahoffmann <tillahoffmann@gmail.com> | 2016-04-01 14:36:15 +0100 |
commit | 49960adbdd78dceab342cf9ea6277ed40fa3b1f8 (patch) | |
tree | c57e9317d19c8dfa93be927a1c5b577c49bf76fa /Eigen | |
parent | 57239f4a8149dbd603ad376e90a0a4574b846710 (diff) | |
parent | 3da495e6b9a9e8def7914b53a8698a09b1998037 (diff) |
Merged eigen/eigen into default
Diffstat (limited to 'Eigen')
-rw-r--r-- | Eigen/src/Core/NumTraits.h | 8 | ||||
-rw-r--r-- | Eigen/src/Core/arch/CUDA/Half.h | 30 |
2 files changed, 35 insertions, 3 deletions
diff --git a/Eigen/src/Core/NumTraits.h b/Eigen/src/Core/NumTraits.h index b7b5e7d22..e065fa714 100644 --- a/Eigen/src/Core/NumTraits.h +++ b/Eigen/src/Core/NumTraits.h @@ -153,7 +153,9 @@ template<typename _Real> struct NumTraits<std::complex<_Real> > MulCost = 4 * NumTraits<Real>::MulCost + 2 * NumTraits<Real>::AddCost }; + EIGEN_DEVICE_FUNC static inline Real epsilon() { return NumTraits<Real>::epsilon(); } + EIGEN_DEVICE_FUNC static inline Real dummy_precision() { return NumTraits<Real>::dummy_precision(); } }; @@ -166,7 +168,7 @@ struct NumTraits<Array<Scalar, Rows, Cols, Options, MaxRows, MaxCols> > typedef typename NumTraits<Scalar>::NonInteger NonIntegerScalar; typedef Array<NonIntegerScalar, Rows, Cols, Options, MaxRows, MaxCols> NonInteger; typedef ArrayType & Nested; - + enum { IsComplex = NumTraits<Scalar>::IsComplex, IsInteger = NumTraits<Scalar>::IsInteger, @@ -176,8 +178,10 @@ struct NumTraits<Array<Scalar, Rows, Cols, Options, MaxRows, MaxCols> > AddCost = ArrayType::SizeAtCompileTime==Dynamic ? HugeCost : ArrayType::SizeAtCompileTime * NumTraits<Scalar>::AddCost, MulCost = ArrayType::SizeAtCompileTime==Dynamic ? HugeCost : ArrayType::SizeAtCompileTime * NumTraits<Scalar>::MulCost }; - + + EIGEN_DEVICE_FUNC static inline RealScalar epsilon() { return NumTraits<RealScalar>::epsilon(); } + EIGEN_DEVICE_FUNC static inline RealScalar dummy_precision() { return NumTraits<RealScalar>::dummy_precision(); } }; diff --git a/Eigen/src/Core/arch/CUDA/Half.h b/Eigen/src/Core/arch/CUDA/Half.h index 6c412159c..212aa0d5d 100644 --- a/Eigen/src/Core/arch/CUDA/Half.h +++ b/Eigen/src/Core/arch/CUDA/Half.h @@ -341,6 +341,18 @@ template<> struct is_arithmetic<half> { enum { value = true }; }; } // end namespace internal +template<> struct NumTraits<Eigen::half> + : GenericNumTraits<Eigen::half> +{ + EIGEN_DEVICE_FUNC static inline float dummy_precision() { return 1e-3f; } + EIGEN_DEVICE_FUNC static inline Eigen::half highest() { + return internal::raw_uint16_to_half(0x7bff); + } + EIGEN_DEVICE_FUNC static inline Eigen::half lowest() { + return internal::raw_uint16_to_half(0xfbff); + } +}; + // Infinity/NaN checks. namespace numext { @@ -348,7 +360,7 @@ namespace numext { static inline EIGEN_DEVICE_FUNC bool (isinf)(const Eigen::half& a) { return (a.x & 0x7fff) == 0x7c00; } -static inline EIGEN_HALF_CUDA_H bool (isnan)(const Eigen::half& a) { +static inline EIGEN_DEVICE_FUNC bool (isnan)(const Eigen::half& a) { #if defined(EIGEN_HAS_CUDA_FP16) && defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 530 return __hisnan(a); #else @@ -403,6 +415,15 @@ using ::sqrt; using ::floor; using ::ceil; +#if __cplusplus > 199711L +template <> +struct hash<Eigen::half> { + size_t operator()(const Eigen::half& a) const { + return std::hash<unsigned short>()(a.x); + } +}; +#endif + } // end namespace std @@ -411,7 +432,14 @@ using ::ceil; __device__ inline Eigen::half __shfl_xor(Eigen::half var, int laneMask, int width=warpSize) { return static_cast<Eigen::half>(__shfl_xor(static_cast<float>(var), laneMask, width)); } +#endif +// ldg() has an overload for __half, but we also need one for Eigen::half. +#if defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 320 +static inline EIGEN_DEVICE_FUNC Eigen::half __ldg(const Eigen::half* ptr) { + return Eigen::internal::raw_uint16_to_half( + __ldg(reinterpret_cast<const unsigned short*>(ptr))); +} #endif |