diff options
| author |  Benoit Steiner <benoit.steiner.goog@gmail.com> | 2016-05-11 20:11:14 -0700 |
|---|---|---|
| committer | Benoit Steiner <benoit.steiner.goog@gmail.com> | 2016-05-11 20:11:14 -0700 |
| commit | 518149e8685cd022fe6c2eac549b3c70284409a9 (patch) | |
| tree | ed59d0c17ac411312b6d65d403e7db49d197ab72 /Eigen/src/Core/arch/CUDA | |
| parent | 56a1757d7409d1c0392aaf33ee0086af897600ca (diff) | |
Misc fixes for fp16
Diffstat (limited to 'Eigen/src/Core/arch/CUDA')
| -rw-r--r-- | Eigen/src/Core/arch/CUDA/Half.h | 4 |
1 files changed, 2 insertions, 2 deletions
diff --git a/Eigen/src/Core/arch/CUDA/Half.h b/Eigen/src/Core/arch/CUDA/Half.h index 2cf49c97b..3ead82829 100644 --- a/Eigen/src/Core/arch/CUDA/Half.h +++ b/Eigen/src/Core/arch/CUDA/Half.h @@ -551,14 +551,14 @@ struct hash<Eigen::half> { // Add the missing shfl_xor intrinsic -#if defined(EIGEN_HAS_CUDA_FP16) && defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 300 +#if defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 300 __device__ EIGEN_STRONG_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 +#if defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 350 static EIGEN_STRONG_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))); |