From 94898488a6fe3096a7a44d0bb108e514f0e44699 Mon Sep 17 00:00:00 2001 From: Deven Desai Date: Mon, 1 Oct 2018 14:28:37 +0000 Subject: This commit contains the following (HIP specific) updates: - unsupported/Eigen/CXX11/src/Tensor/TensorReductionGpu.h Changing "pass-by-reference" argument to be "pass-by-value" instead (in a __global__ function decl). "pass-by-reference" arguments to __global__ functions are unwise, and will be explicitly flagged as errors by the newer versions of HIP. - Eigen/src/Core/util/Memory.h - unsupported/Eigen/CXX11/src/Tensor/TensorContraction.h Changes introduced in recent commits breaks the HIP compile. Adding EIGEN_DEVICE_FUNC attribute to some functions and calling ::malloc/free instead of the corresponding std:: versions to get the HIP compile working again - unsupported/Eigen/CXX11/src/Tensor/TensorReduction.h Change introduced a recent commit breaks the HIP compile (link stage errors out due to failure to inline a function). Disabling the recently introduced code (only for HIP compile), to get the eigen nightly testing going again. Will submit another PR once we have te proper fix. - Eigen/src/Core/util/ConfigureVectorization.h Enabling GPU VECTOR support when HIP compiler is in use (for both the host and device compile phases) --- unsupported/Eigen/CXX11/src/Tensor/TensorContraction.h | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) (limited to 'unsupported/Eigen/CXX11/src/Tensor/TensorContraction.h') diff --git a/unsupported/Eigen/CXX11/src/Tensor/TensorContraction.h b/unsupported/Eigen/CXX11/src/Tensor/TensorContraction.h index b92753c44..6fc1e4a6e 100644 --- a/unsupported/Eigen/CXX11/src/Tensor/TensorContraction.h +++ b/unsupported/Eigen/CXX11/src/Tensor/TensorContraction.h @@ -186,21 +186,21 @@ struct TensorContractionKernel { /*ConjugateLhs*/ false, /*ConjugateRhs*/ false> GebpKernel; - EIGEN_DONT_INLINE + EIGEN_DEVICE_FUNC EIGEN_DONT_INLINE static void packLhs(LhsScalar* lhsBlock, const typename LhsMapper::SubMapper& data_mapper, const StorageIndex depth, const StorageIndex rows) { LhsPacker()(lhsBlock, data_mapper, depth, rows, /*stride*/ 0, /*offset*/ 0); } - EIGEN_DONT_INLINE + EIGEN_DEVICE_FUNC EIGEN_DONT_INLINE static void packRhs(RhsScalar* rhsBlock, const typename RhsMapper::SubMapper& data_mapper, const StorageIndex depth, const StorageIndex cols) { RhsPacker()(rhsBlock, data_mapper, depth, cols); } - EIGEN_DONT_INLINE + EIGEN_DEVICE_FUNC EIGEN_DONT_INLINE static void invoke(const OutputMapper& output_mapper, const LhsScalar* lhsBlock, const RhsScalar* rhsBlock, const StorageIndex rows, const StorageIndex depth, -- cgit v1.2.3