From 309190cf02366fa7ad3d442f51d2dec91bba3ab9 Mon Sep 17 00:00:00 2001 From: Benoit Steiner Date: Tue, 13 Sep 2016 12:42:13 -0700 Subject: [PATCH 1/5] Suppress message 1222 when compiling with nvcc: this ensures that we don't warnings about unknown warning messages when compiling with older versions of nvcc --- Eigen/src/Core/util/DisableStupidWarnings.h | 1 + 1 file changed, 1 insertion(+) diff --git a/Eigen/src/Core/util/DisableStupidWarnings.h b/Eigen/src/Core/util/DisableStupidWarnings.h index dd44c7cbc..23f16ba83 100755 --- a/Eigen/src/Core/util/DisableStupidWarnings.h +++ b/Eigen/src/Core/util/DisableStupidWarnings.h @@ -57,6 +57,7 @@ // Disable the "dynamic initialization in unreachable code" message #pragma diag_suppress initialization_not_reachable // Disable the "calling a __host__ function from a __host__ __device__ function is not allowed" messages (yes, there are 4 of them) + #pragma diag_suppress 1222 #pragma diag_suppress 2651 #pragma diag_suppress 2653 #pragma diag_suppress 2668 From ff47717f25aeede4878f65b214cdce264b8314e8 Mon Sep 17 00:00:00 2001 From: Benoit Steiner Date: Tue, 13 Sep 2016 12:49:40 -0700 Subject: [PATCH 2/5] Suppress warning 2527 and 2529, which correspond to the "calling a __host__ function from a __host__ __device__ function is not allowed" message in nvcc 6.5. --- Eigen/src/Core/util/DisableStupidWarnings.h | 5 ++++- 1 file changed, 4 insertions(+), 1 deletion(-) diff --git a/Eigen/src/Core/util/DisableStupidWarnings.h b/Eigen/src/Core/util/DisableStupidWarnings.h index 23f16ba83..b13e5da25 100755 --- a/Eigen/src/Core/util/DisableStupidWarnings.h +++ b/Eigen/src/Core/util/DisableStupidWarnings.h @@ -56,8 +56,11 @@ #pragma diag_suppress code_is_unreachable // Disable the "dynamic initialization in unreachable code" message #pragma diag_suppress initialization_not_reachable - // Disable the "calling a __host__ function from a __host__ __device__ function is not allowed" messages (yes, there are 4 of them) + // Disable the "invalid error number" message that we get with older versions of nvcc #pragma diag_suppress 1222 + // Disable the "calling a __host__ function from a __host__ __device__ function is not allowed" messages (yes, there are many of them and they seem to change with every version of the compiler) + #pragma diag_suppress 2527 + #pragma diag_suppress 2529 #pragma diag_suppress 2651 #pragma diag_suppress 2653 #pragma diag_suppress 2668 From 779faaaeba8b4d6fa9b2cc62906cccb0be3edf03 Mon Sep 17 00:00:00 2001 From: Benoit Steiner Date: Wed, 14 Sep 2016 09:56:11 -0700 Subject: [PATCH 3/5] Fixed compilation warnings generated by nvcc 6.5 (and below) when compiling the EIGEN_THROW macro --- Eigen/src/Core/util/Macros.h | 4 ++-- Eigen/src/Core/util/Memory.h | 3 +++ 2 files changed, 5 insertions(+), 2 deletions(-) diff --git a/Eigen/src/Core/util/Macros.h b/Eigen/src/Core/util/Macros.h index a9db2f4c7..c1049f7b2 100644 --- a/Eigen/src/Core/util/Macros.h +++ b/Eigen/src/Core/util/Macros.h @@ -954,8 +954,8 @@ namespace Eigen { # define EIGEN_CATCH(X) catch (X) #else # ifdef __CUDA_ARCH__ -# define EIGEN_THROW_X(X) asm("trap;") return {} -# define EIGEN_THROW asm("trap;"); return {} +# define EIGEN_THROW_X(X) asm("trap;") +# define EIGEN_THROW asm("trap;") # else # define EIGEN_THROW_X(X) std::abort() # define EIGEN_THROW std::abort() diff --git a/Eigen/src/Core/util/Memory.h b/Eigen/src/Core/util/Memory.h index 8601c8321..0439655ca 100644 --- a/Eigen/src/Core/util/Memory.h +++ b/Eigen/src/Core/util/Memory.h @@ -275,6 +275,7 @@ template EIGEN_DEVICE_FUNC inline T* construct_elements_of_array(T * destruct_elements_of_array(ptr, i); EIGEN_THROW; } + return NULL; } /***************************************************************************** @@ -305,6 +306,7 @@ template EIGEN_DEVICE_FUNC inline T* aligned_new(size_t size) aligned_free(result); EIGEN_THROW; } + return result; } template EIGEN_DEVICE_FUNC inline T* conditional_aligned_new(size_t size) @@ -320,6 +322,7 @@ template EIGEN_DEVICE_FUNC inline T* conditional_aligned conditional_aligned_free(result); EIGEN_THROW; } + return result; } /** \internal Deletes objects constructed with aligned_new From 488ad7dd1b2806ad57f435ebb2461b6ab57443b9 Mon Sep 17 00:00:00 2001 From: Benoit Steiner Date: Wed, 14 Sep 2016 13:35:00 -0700 Subject: [PATCH 4/5] Added missing EIGEN_DEVICE_FUNC qualifiers --- unsupported/Eigen/CXX11/src/Tensor/TensorCostModel.h | 3 +++ 1 file changed, 3 insertions(+) diff --git a/unsupported/Eigen/CXX11/src/Tensor/TensorCostModel.h b/unsupported/Eigen/CXX11/src/Tensor/TensorCostModel.h index d66e45d50..83c449cf1 100644 --- a/unsupported/Eigen/CXX11/src/Tensor/TensorCostModel.h +++ b/unsupported/Eigen/CXX11/src/Tensor/TensorCostModel.h @@ -51,12 +51,15 @@ class TensorOpCost { internal::scalar_cast_op >::Cost; } + EIGEN_DEVICE_FUNC TensorOpCost() : bytes_loaded_(0), bytes_stored_(0), compute_cycles_(0) {} + EIGEN_DEVICE_FUNC TensorOpCost(double bytes_loaded, double bytes_stored, double compute_cycles) : bytes_loaded_(bytes_loaded), bytes_stored_(bytes_stored), compute_cycles_(compute_cycles) {} + EIGEN_DEVICE_FUNC TensorOpCost(double bytes_loaded, double bytes_stored, double compute_cycles, bool vectorized, double packet_size) : bytes_loaded_(bytes_loaded), From c0d56a543e170d221e58b177b245c0e57ce1993a Mon Sep 17 00:00:00 2001 From: Benoit Steiner Date: Wed, 14 Sep 2016 14:06:21 -0700 Subject: [PATCH 5/5] Added several missing EIGEN_DEVICE_FUNC qualifiers --- Eigen/src/Core/Inverse.h | 2 +- Eigen/src/Core/MathFunctions.h | 3 ++- Eigen/src/Core/MatrixBase.h | 2 +- Eigen/src/LU/FullPivLU.h | 8 ++++---- Eigen/src/LU/InverseImpl.h | 2 +- Eigen/src/LU/PartialPivLU.h | 4 ++-- 6 files changed, 11 insertions(+), 10 deletions(-) diff --git a/Eigen/src/Core/Inverse.h b/Eigen/src/Core/Inverse.h index f3ec84990..f303aebf9 100644 --- a/Eigen/src/Core/Inverse.h +++ b/Eigen/src/Core/Inverse.h @@ -50,7 +50,7 @@ public: typedef typename internal::ref_selector::type Nested; typedef typename internal::remove_all::type NestedExpression; - explicit Inverse(const XprType &xpr) + explicit EIGEN_DEVICE_FUNC Inverse(const XprType &xpr) : m_xpr(xpr) {} diff --git a/Eigen/src/Core/MathFunctions.h b/Eigen/src/Core/MathFunctions.h index bf3044b96..fa322aca7 100644 --- a/Eigen/src/Core/MathFunctions.h +++ b/Eigen/src/Core/MathFunctions.h @@ -1312,11 +1312,12 @@ template struct scalar_fuzzy_default_impl { typedef typename NumTraits::Real RealScalar; - template + template EIGEN_DEVICE_FUNC static inline bool isMuchSmallerThan(const Scalar& x, const OtherScalar& y, const RealScalar& prec) { return numext::abs2(x) <= numext::abs2(y) * prec * prec; } + EIGEN_DEVICE_FUNC static inline bool isApprox(const Scalar& x, const Scalar& y, const RealScalar& prec) { return numext::abs2(x - y) <= numext::mini(numext::abs2(x), numext::abs2(y)) * prec * prec; diff --git a/Eigen/src/Core/MatrixBase.h b/Eigen/src/Core/MatrixBase.h index 334a4d71e..976b130e3 100644 --- a/Eigen/src/Core/MatrixBase.h +++ b/Eigen/src/Core/MatrixBase.h @@ -98,7 +98,7 @@ template class MatrixBase /** \returns the size of the main diagonal, which is min(rows(),cols()). * \sa rows(), cols(), SizeAtCompileTime. */ EIGEN_DEVICE_FUNC - inline Index diagonalSize() const { return (std::min)(rows(),cols()); } + inline Index diagonalSize() const { return (numext::mini)(rows(),cols()); } typedef typename Base::PlainObject PlainObject; diff --git a/Eigen/src/LU/FullPivLU.h b/Eigen/src/LU/FullPivLU.h index 2b30fc146..ebcd5c208 100644 --- a/Eigen/src/LU/FullPivLU.h +++ b/Eigen/src/LU/FullPivLU.h @@ -156,7 +156,7 @@ template class FullPivLU * * \sa permutationQ() */ - inline const PermutationPType& permutationP() const + EIGEN_DEVICE_FUNC inline const PermutationPType& permutationP() const { eigen_assert(m_isInitialized && "LU is not initialized."); return m_p; @@ -406,8 +406,8 @@ template class FullPivLU MatrixType reconstructedMatrix() const; - inline Index rows() const { return m_lu.rows(); } - inline Index cols() const { return m_lu.cols(); } + EIGEN_DEVICE_FUNC inline Index rows() const { return m_lu.rows(); } + EIGEN_DEVICE_FUNC inline Index cols() const { return m_lu.cols(); } #ifndef EIGEN_PARSED_BY_DOXYGEN template @@ -879,7 +879,7 @@ struct Assignment >, internal::assign_ * * \sa class FullPivLU */ -template +template EIGEN_DEVICE_FUNC inline const FullPivLU::PlainObject> MatrixBase::fullPivLu() const { diff --git a/Eigen/src/LU/InverseImpl.h b/Eigen/src/LU/InverseImpl.h index 3134632e1..147f9496c 100644 --- a/Eigen/src/LU/InverseImpl.h +++ b/Eigen/src/LU/InverseImpl.h @@ -327,7 +327,7 @@ struct Assignment, internal::assign_op +template EIGEN_DEVICE_FUNC inline const Inverse MatrixBase::inverse() const { EIGEN_STATIC_ASSERT(!NumTraits::IsInteger,THIS_FUNCTION_IS_NOT_FOR_INTEGER_NUMERIC_TYPES) diff --git a/Eigen/src/LU/PartialPivLU.h b/Eigen/src/LU/PartialPivLU.h index d43961887..13394fffa 100644 --- a/Eigen/src/LU/PartialPivLU.h +++ b/Eigen/src/LU/PartialPivLU.h @@ -584,7 +584,7 @@ struct Assignment >, internal::assi * * \sa class PartialPivLU */ -template +template EIGEN_DEVICE_FUNC inline const PartialPivLU::PlainObject> MatrixBase::partialPivLu() const { @@ -599,7 +599,7 @@ MatrixBase::partialPivLu() const * * \sa class PartialPivLU */ -template +template EIGEN_DEVICE_FUNC inline const PartialPivLU::PlainObject> MatrixBase::lu() const {