diff options
| author |  Benoit Steiner <benoit.steiner.goog@gmail.com> | 2016-03-22 09:30:23 -0700 |
|---|---|---|
| committer | Benoit Steiner <benoit.steiner.goog@gmail.com> | 2016-03-22 09:30:23 -0700 |
| commit | f9ad25e4d8453c4265a5fd6d4962a76a386564df (patch) | |
| tree | 72ebdde98fd7560c477596bd5b37bda30c2b6f30 /Eigen/src | |
| parent | 8ef3181f15a9be76ac783bedd2926ee6f4c69a2f (diff) | |
Fixed contractions of 16 bit floats
Diffstat (limited to 'Eigen/src')
| -rw-r--r-- | Eigen/src/Core/arch/CUDA/Half.h | 20 |
1 files changed, 10 insertions, 10 deletions
diff --git a/Eigen/src/Core/arch/CUDA/Half.h b/Eigen/src/Core/arch/CUDA/Half.h index c385b882a..921c5bcb2 100644 --- a/Eigen/src/Core/arch/CUDA/Half.h +++ b/Eigen/src/Core/arch/CUDA/Half.h @@ -73,8 +73,6 @@ struct half : public __half { : __half(internal::raw_uint16_to_half(b ? 0x3c00 : 0)) {} EIGEN_DEVICE_FUNC half(const __half& h) : __half(h) {} EIGEN_DEVICE_FUNC half(const half& h) : __half(h) {} - EIGEN_DEVICE_FUNC half(const volatile half& h) - : __half(internal::raw_uint16_to_half(h.x)) {} EIGEN_DEVICE_FUNC EIGEN_EXPLICIT_CAST(float) const { return internal::half_to_float(*this); @@ -87,14 +85,6 @@ struct half : public __half { x = other.x; return *this; } - EIGEN_DEVICE_FUNC half& operator=(const volatile half& other) { - x = other.x; - return *this; - } - EIGEN_DEVICE_FUNC volatile half& operator=(const half& other) volatile { - x = other.x; - return *this; - } }; #if defined(EIGEN_HAS_CUDA_FP16) && defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 530 @@ -341,4 +331,14 @@ static inline EIGEN_DEVICE_FUNC Eigen::half log(const Eigen::half& a) { } // end namespace std + +// Add the missing shfl_xor intrinsic +#if defined(EIGEN_HAS_CUDA_FP16) && defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 300 +__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 + + #endif // EIGEN_HALF_CUDA_H |