diff options
| -rw-r--r-- | Eigen/src/Core/arch/CUDA/Half.h | 20 | ||||
| -rw-r--r-- | unsupported/Eigen/CXX11/src/Tensor/TensorContractionCuda.h | 10 |
2 files changed, 15 insertions, 15 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 diff --git a/unsupported/Eigen/CXX11/src/Tensor/TensorContractionCuda.h b/unsupported/Eigen/CXX11/src/Tensor/TensorContractionCuda.h index a4a06ab5f..dbff660a9 100644 --- a/unsupported/Eigen/CXX11/src/Tensor/TensorContractionCuda.h +++ b/unsupported/Eigen/CXX11/src/Tensor/TensorContractionCuda.h @@ -20,7 +20,7 @@ template<typename Scalar, typename Index, typename LhsMapper, typename RhsMapper, typename OutputMapper, bool needs_edge_check> __device__ EIGEN_STRONG_INLINE void EigenContractionKernelInternal(const LhsMapper lhs, const RhsMapper rhs, - const OutputMapper output, volatile Scalar* lhs_shmem, volatile Scalar* rhs_shmem, + const OutputMapper output, Scalar* lhs_shmem, Scalar* rhs_shmem, const Index m_size, const Index n_size, const Index k_size) { const Index m_block_idx = blockIdx.x; @@ -319,8 +319,8 @@ EigenContractionKernelInternal(const LhsMapper lhs, const RhsMapper rhs, Scalar rrow(7); // Now x corresponds to k, y to m, and z to n - const volatile Scalar* lhs_block = &lhs_shmem[threadIdx.x + 9 * threadIdx.y]; - const volatile Scalar* rhs_block = &rhs_shmem[threadIdx.x + 8 * threadIdx.z]; + const Scalar* lhs_block = &lhs_shmem[threadIdx.x + 9 * threadIdx.y]; + const Scalar* rhs_block = &rhs_shmem[threadIdx.x + 8 * threadIdx.z]; #define lhs_element(i, j) lhs_block[72 * ((i) + 8 * (j))] #define rhs_element(i, j) rhs_block[72 * ((i) + 8 * (j))] @@ -503,8 +503,8 @@ __launch_bounds__(512) EigenContractionKernel(const LhsMapper lhs, const RhsMapper rhs, const OutputMapper output, const Index m_size, const Index n_size, const Index k_size) { - __shared__ volatile Scalar lhs_shmem[72 * 64]; - __shared__ volatile Scalar rhs_shmem[72 * 64]; + __shared__ Scalar lhs_shmem[72 * 64]; + __shared__ Scalar rhs_shmem[72 * 64]; const Index m_block_idx = blockIdx.x; const Index n_block_idx = blockIdx.y; |