diff options
| author |  Rasmus Munk Larsen <rmlarsen@google.com> | 2019-12-03 12:08:22 -0800 |
|---|---|---|
| committer | Rasmus Munk Larsen <rmlarsen@google.com> | 2019-12-03 12:08:22 -0800 |
| commit | ead81559c85d6ec7903a47d0c34867f9a87781d7 (patch) | |
| tree | 94f811133885c26df000acae7b1ea240cc03e3a2 /unsupported/Eigen/CXX11 | |
| parent | 6358599ecbe23a8ba1829dea21e532aa85aab932 (diff) | |
Use EIGEN_DEVICE_FUNC macro instead of __device__.
Diffstat (limited to 'unsupported/Eigen/CXX11')
| -rw-r--r-- | unsupported/Eigen/CXX11/src/Tensor/TensorContractionGpu.h | 6 | ||||
| -rw-r--r-- | unsupported/Eigen/CXX11/src/Tensor/TensorExecutor.h | 4 |
2 files changed, 5 insertions, 5 deletions
diff --git a/unsupported/Eigen/CXX11/src/Tensor/TensorContractionGpu.h b/unsupported/Eigen/CXX11/src/Tensor/TensorContractionGpu.h index 3471d1056..941ed0317 100644 --- a/unsupported/Eigen/CXX11/src/Tensor/TensorContractionGpu.h +++ b/unsupported/Eigen/CXX11/src/Tensor/TensorContractionGpu.h @@ -18,7 +18,7 @@ namespace Eigen { template<typename Scalar, typename Index, typename LhsMapper, typename RhsMapper, typename OutputMapper, bool needs_edge_check> -__device__ EIGEN_STRONG_INLINE void +EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void EigenContractionKernelInternal(const LhsMapper lhs, const RhsMapper rhs, const OutputMapper output, Scalar* lhs_shmem, Scalar* rhs_shmem, const Index m_size, const Index n_size, const Index k_size) { @@ -531,7 +531,7 @@ EigenContractionKernel(const LhsMapper lhs, const RhsMapper rhs, template<typename Index, typename LhsMapper, typename RhsMapper, typename OutputMapper, bool CHECK_LHS_BOUNDARY, bool CHECK_RHS_BOUNDARY> -__device__ EIGEN_STRONG_INLINE void +EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void EigenFloatContractionKernelInternal16x16(const LhsMapper lhs, const RhsMapper rhs, const OutputMapper output, float2 lhs_shmem2[][16], float2 rhs_shmem2[][8], const Index m_size, @@ -771,7 +771,7 @@ EigenFloatContractionKernelInternal16x16(const LhsMapper lhs, const RhsMapper rh template<typename Index, typename LhsMapper, typename RhsMapper, typename OutputMapper, bool CHECK_LHS_BOUNDARY, bool CHECK_RHS_BOUNDARY> -__device__ EIGEN_STRONG_INLINE void +EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void EigenFloatContractionKernelInternal(const LhsMapper lhs, const RhsMapper rhs, const OutputMapper output, float2 lhs_shmem2[][32], float2 rhs_shmem2[][8], const Index m_size, diff --git a/unsupported/Eigen/CXX11/src/Tensor/TensorExecutor.h b/unsupported/Eigen/CXX11/src/Tensor/TensorExecutor.h index b83174ab7..ca056e96e 100644 --- a/unsupported/Eigen/CXX11/src/Tensor/TensorExecutor.h +++ b/unsupported/Eigen/CXX11/src/Tensor/TensorExecutor.h @@ -582,7 +582,7 @@ class TensorExecutor<Expression, GpuDevice, Vectorizable, Tiling> { #if defined(EIGEN_GPUCC) template <typename Evaluator, typename StorageIndex, bool Vectorizable> struct EigenMetaKernelEval { - static __device__ EIGEN_ALWAYS_INLINE + static EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE void run(Evaluator& eval, StorageIndex firstIdx, StorageIndex lastIdx, StorageIndex step_size) { for (StorageIndex i = firstIdx; i < lastIdx; i += step_size) { eval.evalScalar(i); @@ -592,7 +592,7 @@ struct EigenMetaKernelEval { template <typename Evaluator, typename StorageIndex> struct EigenMetaKernelEval<Evaluator, StorageIndex, true> { - static __device__ EIGEN_ALWAYS_INLINE + static EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE void run(Evaluator& eval, StorageIndex firstIdx, StorageIndex lastIdx, StorageIndex step_size) { const StorageIndex PacketSize = unpacket_traits<typename Evaluator::PacketReturnType>::size; const StorageIndex vectorized_size = (lastIdx / PacketSize) * PacketSize; |