diff options
author | Jonas Adler <jonasadl@kth.se> | 2015-07-22 12:29:18 +0200 |
---|---|---|
committer | Jonas Adler <jonasadl@kth.se> | 2015-07-22 12:29:18 +0200 |
commit | 815fa0dbf63f112f98fe3ac38483e0248caf9eec (patch) | |
tree | 45cff143ff8ffc55edde6caff72d667af6251b8f | |
parent | d259b719d1703e9271612d944e432edd72be1051 (diff) |
Fixed some compiler bugs in NVCC, now compiles with CUDA.
(chtz: Manually joined sevaral commits to keep the history clean)
-rw-r--r-- | Eigen/Core | 8 | ||||
-rw-r--r-- | Eigen/src/Core/DenseBase.h | 49 | ||||
-rw-r--r-- | Eigen/src/Core/DenseStorage.h | 78 | ||||
-rw-r--r-- | Eigen/src/Core/Replicate.h | 17 | ||||
-rw-r--r-- | Eigen/src/Core/Reverse.h | 9 | ||||
-rw-r--r-- | Eigen/src/Core/VectorwiseOp.h | 77 | ||||
-rw-r--r-- | Eigen/src/Core/util/Macros.h | 9 | ||||
-rw-r--r-- | Eigen/src/Core/util/Memory.h | 6 | ||||
-rw-r--r-- | Eigen/src/Core/util/Meta.h | 4 | ||||
-rw-r--r-- | doc/snippets/compile_snippet.cpp.in | 5 |
10 files changed, 139 insertions, 123 deletions
diff --git a/Eigen/Core b/Eigen/Core index de94b5b75..91713a43e 100644 --- a/Eigen/Core +++ b/Eigen/Core @@ -24,9 +24,15 @@ #ifdef EIGEN_INTERNAL_DEBUGGING #undef EIGEN_INTERNAL_DEBUGGING #endif - + // Do not try to vectorize on CUDA! + #ifndef EIGEN_DONT_VECTORIZE #define EIGEN_DONT_VECTORIZE + #endif + + #ifdef EIGEN_EXCEPTIONS + #undef EIGEN_EXCEPTIONS + #endif // All functions callable from CUDA code must be qualified with __device__ #define EIGEN_DEVICE_FUNC __host__ __device__ diff --git a/Eigen/src/Core/DenseBase.h b/Eigen/src/Core/DenseBase.h index c603b5a6d..dc22a128e 100644 --- a/Eigen/src/Core/DenseBase.h +++ b/Eigen/src/Core/DenseBase.h @@ -491,9 +491,29 @@ template<typename Derived> class DenseBase typedef VectorwiseOp<Derived, Vertical> ColwiseReturnType; typedef const VectorwiseOp<const Derived, Vertical> ConstColwiseReturnType; - EIGEN_DEVICE_FUNC ConstRowwiseReturnType rowwise() const; + /** \returns a VectorwiseOp wrapper of *this providing additional partial reduction operations + * + * Example: \include MatrixBase_rowwise.cpp + * Output: \verbinclude MatrixBase_rowwise.out + * + * \sa colwise(), class VectorwiseOp, \ref TutorialReductionsVisitorsBroadcasting + */ + //Code moved here due to a CUDA compiler bug + EIGEN_DEVICE_FUNC inline ConstRowwiseReturnType rowwise() const { + return ConstRowwiseReturnType(derived()); + } EIGEN_DEVICE_FUNC RowwiseReturnType rowwise(); - EIGEN_DEVICE_FUNC ConstColwiseReturnType colwise() const; + + /** \returns a VectorwiseOp wrapper of *this providing additional partial reduction operations + * + * Example: \include MatrixBase_colwise.cpp + * Output: \verbinclude MatrixBase_colwise.out + * + * \sa rowwise(), class VectorwiseOp, \ref TutorialReductionsVisitorsBroadcasting + */ + EIGEN_DEVICE_FUNC inline ConstColwiseReturnType colwise() const { + return ConstColwiseReturnType(derived()); + } EIGEN_DEVICE_FUNC ColwiseReturnType colwise(); typedef CwiseNullaryOp<internal::scalar_random_op<Scalar>,PlainObject> RandomReturnType; @@ -519,14 +539,31 @@ template<typename Derived> class DenseBase template<int RowFactor, int ColFactor> EIGEN_DEVICE_FUNC const Replicate<Derived,RowFactor,ColFactor> replicate() const; + /** + * \return an expression of the replication of \c *this + * + * Example: \include MatrixBase_replicate_int_int.cpp + * Output: \verbinclude MatrixBase_replicate_int_int.out + * + * \sa VectorwiseOp::replicate(), DenseBase::replicate<int,int>(), class Replicate + */ + //Code moved here due to a CUDA compiler bug EIGEN_DEVICE_FUNC - const Replicate<Derived,Dynamic,Dynamic> replicate(Index rowFacor,Index colFactor) const; + const Replicate<Derived, Dynamic, Dynamic> replicate(Index rowFactor, Index colFactor) const + { + return Replicate<Derived, Dynamic, Dynamic>(derived(), rowFactor, colFactor); + } typedef Reverse<Derived, BothDirections> ReverseReturnType; typedef const Reverse<const Derived, BothDirections> ConstReverseReturnType; - ReverseReturnType reverse(); - ConstReverseReturnType reverse() const; - void reverseInPlace(); + EIGEN_DEVICE_FUNC ReverseReturnType reverse(); + /** This is the const version of reverse(). */ + //Code moved here due to a CUDA compiler bug + EIGEN_DEVICE_FUNC ConstReverseReturnType reverse() const + { + return ConstReverseReturnType(derived()); + } + EIGEN_DEVICE_FUNC void reverseInPlace(); #define EIGEN_CURRENT_STORAGE_BASE_CLASS Eigen::DenseBase # include "../plugins/BlockMethods.h" diff --git a/Eigen/src/Core/DenseStorage.h b/Eigen/src/Core/DenseStorage.h index 80c4c6e8e..5eb434c6d 100644 --- a/Eigen/src/Core/DenseStorage.h +++ b/Eigen/src/Core/DenseStorage.h @@ -270,10 +270,10 @@ template<typename T, int Size, int _Options> class DenseStorage<T, Size, Dynamic Index m_cols; public: EIGEN_DEVICE_FUNC DenseStorage() : m_rows(0), m_cols(0) {} - explicit DenseStorage(internal::constructor_without_unaligned_array_assert) + EIGEN_DEVICE_FUNC explicit DenseStorage(internal::constructor_without_unaligned_array_assert) : m_data(internal::constructor_without_unaligned_array_assert()), m_rows(0), m_cols(0) {} - DenseStorage(const DenseStorage& other) : m_data(other.m_data), m_rows(other.m_rows), m_cols(other.m_cols) {} - DenseStorage& operator=(const DenseStorage& other) + EIGEN_DEVICE_FUNC DenseStorage(const DenseStorage& other) : m_data(other.m_data), m_rows(other.m_rows), m_cols(other.m_cols) {} + EIGEN_DEVICE_FUNC DenseStorage& operator=(const DenseStorage& other) { if (this != &other) { @@ -283,13 +283,13 @@ template<typename T, int Size, int _Options> class DenseStorage<T, Size, Dynamic } return *this; } - DenseStorage(Index, Index rows, Index cols) : m_rows(rows), m_cols(cols) {} - void swap(DenseStorage& other) + EIGEN_DEVICE_FUNC DenseStorage(Index, Index rows, Index cols) : m_rows(rows), m_cols(cols) {} + EIGEN_DEVICE_FUNC void swap(DenseStorage& other) { std::swap(m_data,other.m_data); std::swap(m_rows,other.m_rows); std::swap(m_cols,other.m_cols); } EIGEN_DEVICE_FUNC Index rows() const {return m_rows;} EIGEN_DEVICE_FUNC Index cols() const {return m_cols;} - void conservativeResize(Index, Index rows, Index cols) { m_rows = rows; m_cols = cols; } - void resize(Index, Index rows, Index cols) { m_rows = rows; m_cols = cols; } + EIGEN_DEVICE_FUNC void conservativeResize(Index, Index rows, Index cols) { m_rows = rows; m_cols = cols; } + EIGEN_DEVICE_FUNC void resize(Index, Index rows, Index cols) { m_rows = rows; m_cols = cols; } EIGEN_DEVICE_FUNC const T *data() const { return m_data.array; } EIGEN_DEVICE_FUNC T *data() { return m_data.array; } }; @@ -301,10 +301,10 @@ template<typename T, int Size, int _Cols, int _Options> class DenseStorage<T, Si Index m_rows; public: EIGEN_DEVICE_FUNC DenseStorage() : m_rows(0) {} - explicit DenseStorage(internal::constructor_without_unaligned_array_assert) + EIGEN_DEVICE_FUNC explicit DenseStorage(internal::constructor_without_unaligned_array_assert) : m_data(internal::constructor_without_unaligned_array_assert()), m_rows(0) {} - DenseStorage(const DenseStorage& other) : m_data(other.m_data), m_rows(other.m_rows) {} - DenseStorage& operator=(const DenseStorage& other) + EIGEN_DEVICE_FUNC DenseStorage(const DenseStorage& other) : m_data(other.m_data), m_rows(other.m_rows) {} + EIGEN_DEVICE_FUNC DenseStorage& operator=(const DenseStorage& other) { if (this != &other) { @@ -313,12 +313,12 @@ template<typename T, int Size, int _Cols, int _Options> class DenseStorage<T, Si } return *this; } - DenseStorage(Index, Index rows, Index) : m_rows(rows) {} - void swap(DenseStorage& other) { std::swap(m_data,other.m_data); std::swap(m_rows,other.m_rows); } + EIGEN_DEVICE_FUNC DenseStorage(Index, Index rows, Index) : m_rows(rows) {} + EIGEN_DEVICE_FUNC void swap(DenseStorage& other) { std::swap(m_data,other.m_data); std::swap(m_rows,other.m_rows); } EIGEN_DEVICE_FUNC Index rows(void) const {return m_rows;} EIGEN_DEVICE_FUNC Index cols(void) const {return _Cols;} - void conservativeResize(Index, Index rows, Index) { m_rows = rows; } - void resize(Index, Index rows, Index) { m_rows = rows; } + EIGEN_DEVICE_FUNC void conservativeResize(Index, Index rows, Index) { m_rows = rows; } + EIGEN_DEVICE_FUNC void resize(Index, Index rows, Index) { m_rows = rows; } EIGEN_DEVICE_FUNC const T *data() const { return m_data.array; } EIGEN_DEVICE_FUNC T *data() { return m_data.array; } }; @@ -330,10 +330,10 @@ template<typename T, int Size, int _Rows, int _Options> class DenseStorage<T, Si Index m_cols; public: EIGEN_DEVICE_FUNC DenseStorage() : m_cols(0) {} - explicit DenseStorage(internal::constructor_without_unaligned_array_assert) + EIGEN_DEVICE_FUNC explicit DenseStorage(internal::constructor_without_unaligned_array_assert) : m_data(internal::constructor_without_unaligned_array_assert()), m_cols(0) {} - DenseStorage(const DenseStorage& other) : m_data(other.m_data), m_cols(other.m_cols) {} - DenseStorage& operator=(const DenseStorage& other) + EIGEN_DEVICE_FUNC DenseStorage(const DenseStorage& other) : m_data(other.m_data), m_cols(other.m_cols) {} + EIGEN_DEVICE_FUNC DenseStorage& operator=(const DenseStorage& other) { if (this != &other) { @@ -342,8 +342,8 @@ template<typename T, int Size, int _Rows, int _Options> class DenseStorage<T, Si } return *this; } - DenseStorage(Index, Index, Index cols) : m_cols(cols) {} - void swap(DenseStorage& other) { std::swap(m_data,other.m_data); std::swap(m_cols,other.m_cols); } + EIGEN_DEVICE_FUNC DenseStorage(Index, Index, Index cols) : m_cols(cols) {} + EIGEN_DEVICE_FUNC void swap(DenseStorage& other) { std::swap(m_data,other.m_data); std::swap(m_cols,other.m_cols); } EIGEN_DEVICE_FUNC Index rows(void) const {return _Rows;} EIGEN_DEVICE_FUNC Index cols(void) const {return m_cols;} void conservativeResize(Index, Index, Index cols) { m_cols = cols; } @@ -360,22 +360,22 @@ template<typename T, int _Options> class DenseStorage<T, Dynamic, Dynamic, Dynam Index m_cols; public: EIGEN_DEVICE_FUNC DenseStorage() : m_data(0), m_rows(0), m_cols(0) {} - explicit DenseStorage(internal::constructor_without_unaligned_array_assert) + EIGEN_DEVICE_FUNC explicit DenseStorage(internal::constructor_without_unaligned_array_assert) : m_data(0), m_rows(0), m_cols(0) {} - DenseStorage(Index size, Index rows, Index cols) + EIGEN_DEVICE_FUNC DenseStorage(Index size, Index rows, Index cols) : m_data(internal::conditional_aligned_new_auto<T,(_Options&DontAlign)==0>(size)), m_rows(rows), m_cols(cols) { EIGEN_INTERNAL_DENSE_STORAGE_CTOR_PLUGIN eigen_internal_assert(size==rows*cols && rows>=0 && cols >=0); } - DenseStorage(const DenseStorage& other) + EIGEN_DEVICE_FUNC DenseStorage(const DenseStorage& other) : m_data(internal::conditional_aligned_new_auto<T,(_Options&DontAlign)==0>(other.m_rows*other.m_cols)) , m_rows(other.m_rows) , m_cols(other.m_cols) { internal::smart_copy(other.m_data, other.m_data+other.m_rows*other.m_cols, m_data); } - DenseStorage& operator=(const DenseStorage& other) + EIGEN_DEVICE_FUNC DenseStorage& operator=(const DenseStorage& other) { if (this != &other) { @@ -405,8 +405,8 @@ template<typename T, int _Options> class DenseStorage<T, Dynamic, Dynamic, Dynam return *this; } #endif - ~DenseStorage() { internal::conditional_aligned_delete_auto<T,(_Options&DontAlign)==0>(m_data, m_rows*m_cols); } - void swap(DenseStorage& other) + EIGEN_DEVICE_FUNC ~DenseStorage() { internal::conditional_aligned_delete_auto<T,(_Options&DontAlign)==0>(m_data, m_rows*m_cols); } + EIGEN_DEVICE_FUNC void swap(DenseStorage& other) { std::swap(m_data,other.m_data); std::swap(m_rows,other.m_rows); std::swap(m_cols,other.m_cols); } EIGEN_DEVICE_FUNC Index rows(void) const {return m_rows;} EIGEN_DEVICE_FUNC Index cols(void) const {return m_cols;} @@ -416,7 +416,7 @@ template<typename T, int _Options> class DenseStorage<T, Dynamic, Dynamic, Dynam m_rows = rows; m_cols = cols; } - void resize(Index size, Index rows, Index cols) + EIGEN_DEVICE_FUNC void resize(Index size, Index rows, Index cols) { if(size != m_rows*m_cols) { @@ -442,19 +442,19 @@ template<typename T, int _Rows, int _Options> class DenseStorage<T, Dynamic, _Ro public: EIGEN_DEVICE_FUNC DenseStorage() : m_data(0), m_cols(0) {} explicit DenseStorage(internal::constructor_without_unaligned_array_assert) : m_data(0), m_cols(0) {} - DenseStorage(Index size, Index rows, Index cols) : m_data(internal::conditional_aligned_new_auto<T,(_Options&DontAlign)==0>(size)), m_cols(cols) + EIGEN_DEVICE_FUNC DenseStorage(Index size, Index rows, Index cols) : m_data(internal::conditional_aligned_new_auto<T,(_Options&DontAlign)==0>(size)), m_cols(cols) { EIGEN_INTERNAL_DENSE_STORAGE_CTOR_PLUGIN eigen_internal_assert(size==rows*cols && rows==_Rows && cols >=0); EIGEN_UNUSED_VARIABLE(rows); } - DenseStorage(const DenseStorage& other) + EIGEN_DEVICE_FUNC DenseStorage(const DenseStorage& other) : m_data(internal::conditional_aligned_new_auto<T,(_Options&DontAlign)==0>(_Rows*other.m_cols)) , m_cols(other.m_cols) { internal::smart_copy(other.m_data, other.m_data+_Rows*m_cols, m_data); } - DenseStorage& operator=(const DenseStorage& other) + EIGEN_DEVICE_FUNC DenseStorage& operator=(const DenseStorage& other) { if (this != &other) { @@ -481,16 +481,16 @@ template<typename T, int _Rows, int _Options> class DenseStorage<T, Dynamic, _Ro return *this; } #endif - ~DenseStorage() { internal::conditional_aligned_delete_auto<T,(_Options&DontAlign)==0>(m_data, _Rows*m_cols); } - void swap(DenseStorage& other) { std::swap(m_data,other.m_data); std::swap(m_cols,other.m_cols); } + EIGEN_DEVICE_FUNC ~DenseStorage() { internal::conditional_aligned_delete_auto<T,(_Options&DontAlign)==0>(m_data, _Rows*m_cols); } + EIGEN_DEVICE_FUNC void swap(DenseStorage& other) { std::swap(m_data,other.m_data); std::swap(m_cols,other.m_cols); } EIGEN_DEVICE_FUNC static Index rows(void) {return _Rows;} EIGEN_DEVICE_FUNC Index cols(void) const {return m_cols;} - void conservativeResize(Index size, Index, Index cols) + EIGEN_DEVICE_FUNC void conservativeResize(Index size, Index, Index cols) { m_data = internal::conditional_aligned_realloc_new_auto<T,(_Options&DontAlign)==0>(m_data, size, _Rows*m_cols); m_cols = cols; } - EIGEN_STRONG_INLINE void resize(Index size, Index, Index cols) + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void resize(Index size, Index, Index cols) { if(size != _Rows*m_cols) { @@ -515,19 +515,19 @@ template<typename T, int _Cols, int _Options> class DenseStorage<T, Dynamic, Dyn public: EIGEN_DEVICE_FUNC DenseStorage() : m_data(0), m_rows(0) {} explicit DenseStorage(internal::constructor_without_unaligned_array_assert) : m_data(0), m_rows(0) {} - DenseStorage(Index size, Index rows, Index cols) : m_data(internal::conditional_aligned_new_auto<T,(_Options&DontAlign)==0>(size)), m_rows(rows) + EIGEN_DEVICE_FUNC DenseStorage(Index size, Index rows, Index cols) : m_data(internal::conditional_aligned_new_auto<T,(_Options&DontAlign)==0>(size)), m_rows(rows) { EIGEN_INTERNAL_DENSE_STORAGE_CTOR_PLUGIN eigen_internal_assert(size==rows*cols && rows>=0 && cols == _Cols); EIGEN_UNUSED_VARIABLE(cols); } - DenseStorage(const DenseStorage& other) + EIGEN_DEVICE_FUNC DenseStorage(const DenseStorage& other) : m_data(internal::conditional_aligned_new_auto<T,(_Options&DontAlign)==0>(other.m_rows*_Cols)) , m_rows(other.m_rows) { internal::smart_copy(other.m_data, other.m_data+other.m_rows*_Cols, m_data); } - DenseStorage& operator=(const DenseStorage& other) + EIGEN_DEVICE_FUNC DenseStorage& operator=(const DenseStorage& other) { if (this != &other) { @@ -554,8 +554,8 @@ template<typename T, int _Cols, int _Options> class DenseStorage<T, Dynamic, Dyn return *this; } #endif - ~DenseStorage() { internal::conditional_aligned_delete_auto<T,(_Options&DontAlign)==0>(m_data, _Cols*m_rows); } - void swap(DenseStorage& other) { std::swap(m_data,other.m_data); std::swap(m_rows,other.m_rows); } + EIGEN_DEVICE_FUNC ~DenseStorage() { internal::conditional_aligned_delete_auto<T,(_Options&DontAlign)==0>(m_data, _Cols*m_rows); } + EIGEN_DEVICE_FUNC void swap(DenseStorage& other) { std::swap(m_data,other.m_data); std::swap(m_rows,other.m_rows); } EIGEN_DEVICE_FUNC Index rows(void) const {return m_rows;} EIGEN_DEVICE_FUNC static Index cols(void) {return _Cols;} void conservativeResize(Index size, Index rows, Index) @@ -563,7 +563,7 @@ template<typename T, int _Cols, int _Options> class DenseStorage<T, Dynamic, Dyn m_data = internal::conditional_aligned_realloc_new_auto<T,(_Options&DontAlign)==0>(m_data, size, m_rows*_Cols); m_rows = rows; } - EIGEN_STRONG_INLINE void resize(Index size, Index rows, Index) + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void resize(Index size, Index rows, Index) { if(size != m_rows*_Cols) { diff --git a/Eigen/src/Core/Replicate.h b/Eigen/src/Core/Replicate.h index 518c52e57..bec598310 100644 --- a/Eigen/src/Core/Replicate.h +++ b/Eigen/src/Core/Replicate.h @@ -114,28 +114,13 @@ template<typename MatrixType,int RowFactor,int ColFactor> class Replicate */ template<typename Derived> template<int RowFactor, int ColFactor> -inline const Replicate<Derived,RowFactor,ColFactor> +const Replicate<Derived,RowFactor,ColFactor> DenseBase<Derived>::replicate() const { return Replicate<Derived,RowFactor,ColFactor>(derived()); } /** - * \return an expression of the replication of \c *this - * - * Example: \include MatrixBase_replicate_int_int.cpp - * Output: \verbinclude MatrixBase_replicate_int_int.out - * - * \sa VectorwiseOp::replicate(), DenseBase::replicate<int,int>(), class Replicate - */ -template<typename Derived> -inline const Replicate<Derived,Dynamic,Dynamic> -DenseBase<Derived>::replicate(Index rowFactor,Index colFactor) const -{ - return Replicate<Derived,Dynamic,Dynamic>(derived(),rowFactor,colFactor); -} - -/** * \return an expression of the replication of each column (or row) of \c *this * * Example: \include DirectionWise_replicate_int.cpp diff --git a/Eigen/src/Core/Reverse.h b/Eigen/src/Core/Reverse.h index ef301e66d..8e7f6e927 100644 --- a/Eigen/src/Core/Reverse.h +++ b/Eigen/src/Core/Reverse.h @@ -120,13 +120,8 @@ DenseBase<Derived>::reverse() return ReverseReturnType(derived()); } -/** This is the const version of reverse(). */ -template<typename Derived> -inline const typename DenseBase<Derived>::ConstReverseReturnType -DenseBase<Derived>::reverse() const -{ - return ConstReverseReturnType(derived()); -} + +//reverse const overload moved DenseBase.h due to a CUDA compiler bug /** This is the "in place" version of reverse: it reverses \c *this. * diff --git a/Eigen/src/Core/VectorwiseOp.h b/Eigen/src/Core/VectorwiseOp.h index 0cc5eff16..fd0ea412d 100644 --- a/Eigen/src/Core/VectorwiseOp.h +++ b/Eigen/src/Core/VectorwiseOp.h @@ -82,7 +82,7 @@ class PartialReduxExpr : public internal::dense_xpr_base< PartialReduxExpr<Matri return m_functor(m_matrix.row(i)); } - const Scalar coeff(Index index) const + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE const Scalar coeff(Index index) const { if (Direction==Vertical) return m_functor(m_matrix.col(index)); @@ -185,15 +185,15 @@ template<typename ExpressionType, int Direction> class VectorwiseOp }; enum { - IsVertical = (Direction==Vertical) ? 1 : 0, - IsHorizontal = (Direction==Horizontal) ? 1 : 0 + isVertical = (Direction==Vertical) ? 1 : 0, + isHorizontal = (Direction==Horizontal) ? 1 : 0 }; protected: /** \internal * \returns the i-th subvector according to the \c Direction */ - typedef typename internal::conditional<Direction==Vertical, + typedef typename internal::conditional<isVertical, typename ExpressionType::ColXpr, typename ExpressionType::RowXpr>::type SubVector; EIGEN_DEVICE_FUNC @@ -206,12 +206,12 @@ template<typename ExpressionType, int Direction> class VectorwiseOp * \returns the number of subvectors in the direction \c Direction */ EIGEN_DEVICE_FUNC Index subVectors() const - { return Direction==Vertical?m_matrix.cols():m_matrix.rows(); } + { return isVertical?m_matrix.cols():m_matrix.rows(); } template<typename OtherDerived> struct ExtendedType { typedef Replicate<OtherDerived, - Direction==Vertical ? 1 : ExpressionType::RowsAtCompileTime, - Direction==Horizontal ? 1 : ExpressionType::ColsAtCompileTime> Type; + isVertical ? 1 : ExpressionType::RowsAtCompileTime, + isHorizontal ? 1 : ExpressionType::ColsAtCompileTime> Type; }; /** \internal @@ -221,20 +221,20 @@ template<typename ExpressionType, int Direction> class VectorwiseOp typename ExtendedType<OtherDerived>::Type extendedTo(const DenseBase<OtherDerived>& other) const { - EIGEN_STATIC_ASSERT(EIGEN_IMPLIES(Direction==Vertical, OtherDerived::MaxColsAtCompileTime==1), + EIGEN_STATIC_ASSERT(EIGEN_IMPLIES(isVertical, OtherDerived::MaxColsAtCompileTime==1), YOU_PASSED_A_ROW_VECTOR_BUT_A_COLUMN_VECTOR_WAS_EXPECTED) - EIGEN_STATIC_ASSERT(EIGEN_IMPLIES(Direction==Horizontal, OtherDerived::MaxRowsAtCompileTime==1), + EIGEN_STATIC_ASSERT(EIGEN_IMPLIES(isHorizontal, OtherDerived::MaxRowsAtCompileTime==1), YOU_PASSED_A_COLUMN_VECTOR_BUT_A_ROW_VECTOR_WAS_EXPECTED) return typename ExtendedType<OtherDerived>::Type (other.derived(), - Direction==Vertical ? 1 : m_matrix.rows(), - Direction==Horizontal ? 1 : m_matrix.cols()); + isVertical ? 1 : m_matrix.rows(), + isHorizontal ? 1 : m_matrix.cols()); } template<typename OtherDerived> struct OppositeExtendedType { typedef Replicate<OtherDerived, - Direction==Horizontal ? 1 : ExpressionType::RowsAtCompileTime, - Direction==Vertical ? 1 : ExpressionType::ColsAtCompileTime> Type; + isHorizontal ? 1 : ExpressionType::RowsAtCompileTime, + isVertical ? 1 : ExpressionType::ColsAtCompileTime> Type; }; /** \internal @@ -244,18 +244,17 @@ template<typename ExpressionType, int Direction> class VectorwiseOp typename OppositeExtendedType<OtherDerived>::Type extendedToOpposite(const DenseBase<OtherDerived>& other) const { - EIGEN_STATIC_ASSERT(EIGEN_IMPLIES(Direction==Horizontal, OtherDerived::MaxColsAtCompileTime==1), + EIGEN_STATIC_ASSERT(EIGEN_IMPLIES(isHorizontal, OtherDerived::MaxColsAtCompileTime==1), YOU_PASSED_A_ROW_VECTOR_BUT_A_COLUMN_VECTOR_WAS_EXPECTED) - EIGEN_STATIC_ASSERT(EIGEN_IMPLIES(Direction==Vertical, OtherDerived::MaxRowsAtCompileTime==1), + EIGEN_STATIC_ASSERT(EIGEN_IMPLIES(isVertical, OtherDerived::MaxRowsAtCompileTime==1), YOU_PASSED_A_COLUMN_VECTOR_BUT_A_ROW_VECTOR_WAS_EXPECTED) return typename OppositeExtendedType<OtherDerived>::Type (other.derived(), - Direction==Horizontal ? 1 : m_matrix.rows(), - Direction==Vertical ? 1 : m_matrix.cols()); + isHorizontal ? 1 : m_matrix.rows(), + isVertical ? 1 : m_matrix.cols()); } public: - EIGEN_DEVICE_FUNC explicit inline VectorwiseOp(ExpressionType& matrix) : m_matrix(matrix) {} @@ -271,8 +270,8 @@ template<typename ExpressionType, int Direction> class VectorwiseOp * \sa class VectorwiseOp, DenseBase::colwise(), DenseBase::rowwise() */ template<typename BinaryOp> - const typename ReduxReturnType<BinaryOp>::Type EIGEN_DEVICE_FUNC + const typename ReduxReturnType<BinaryOp>::Type redux(const BinaryOp& func = BinaryOp()) const { return typename ReduxReturnType<BinaryOp>::Type(_expression(), internal::member_redux<BinaryOp,Scalar>(func)); } @@ -447,7 +446,7 @@ template<typename ExpressionType, int Direction> class VectorwiseOp const ReverseReturnType reverse() const { return ReverseReturnType( _expression() ); } - typedef Replicate<ExpressionType,Direction==Vertical?Dynamic:1,Direction==Horizontal?Dynamic:1> ReplicateReturnType; + typedef Replicate<ExpressionType,(isVertical?Dynamic:1),(isHorizontal?Dynamic:1)> ReplicateReturnType; EIGEN_DEVICE_FUNC const ReplicateReturnType replicate(Index factor) const; @@ -460,12 +459,13 @@ template<typename ExpressionType, int Direction> class VectorwiseOp * \sa VectorwiseOp::replicate(Index), DenseBase::replicate(), class Replicate */ // NOTE implemented here because of sunstudio's compilation errors - template<int Factor> const Replicate<ExpressionType,(IsVertical?Factor:1),(IsHorizontal?Factor:1)> + // isVertical*Factor+isHorizontal instead of (isVertical?Factor:1) to handle CUDA bug with ternary operator + template<int Factor> const Replicate<ExpressionType,isVertical*Factor+isHorizontal,isHorizontal*Factor+isVertical> EIGEN_DEVICE_FUNC replicate(Index factor = Factor) const { - return Replicate<ExpressionType,Direction==Vertical?Factor:1,Direction==Horizontal?Factor:1> - (_expression(),Direction==Vertical?factor:1,Direction==Horizontal?factor:1); + return Replicate<ExpressionType,(isVertical?Factor:1),(isHorizontal?Factor:1)> + (_expression(),isVertical?factor:1,isHorizontal?factor:1); } /////////// Artithmetic operators /////////// @@ -556,6 +556,7 @@ template<typename ExpressionType, int Direction> class VectorwiseOp CwiseBinaryOp<internal::scalar_product_op<Scalar>, const ExpressionTypeNestedCleaned, const typename ExtendedType<OtherDerived>::Type> + EIGEN_DEVICE_FUNC operator*(const DenseBase<OtherDerived>& other) const { EIGEN_STATIC_ASSERT_VECTOR_ONLY(OtherDerived) @@ -637,19 +638,8 @@ template<typename ExpressionType, int Direction> class VectorwiseOp ExpressionTypeNested m_matrix; }; -/** \returns a VectorwiseOp wrapper of *this providing additional partial reduction operations - * - * Example: \include MatrixBase_colwise.cpp - * Output: \verbinclude MatrixBase_colwise.out - * - * \sa rowwise(), class VectorwiseOp, \ref TutorialReductionsVisitorsBroadcasting - */ -template<typename Derived> -inline const typename DenseBase<Derived>::ConstColwiseReturnType -DenseBase<Derived>::colwise() const -{ - return ConstColwiseReturnType(derived()); -} +//const colwise moved to DenseBase.h due to CUDA compiler bug + /** \returns a writable VectorwiseOp wrapper of *this providing additional partial reduction operations * @@ -662,19 +652,8 @@ DenseBase<Derived>::colwise() return ColwiseReturnType(derived()); } -/** \returns a VectorwiseOp wrapper of *this providing additional partial reduction operations - * - * Example: \include MatrixBase_rowwise.cpp - * Output: \verbinclude MatrixBase_rowwise.out - * - * \sa colwise(), class VectorwiseOp, \ref TutorialReductionsVisitorsBroadcasting - */ -template<typename Derived> -inline const typename DenseBase<Derived>::ConstRowwiseReturnType -DenseBase<Derived>::rowwise() const -{ - return ConstRowwiseReturnType(derived()); -} +//const rowwise moved to DenseBase.h due to CUDA compiler bug + /** \returns a writable VectorwiseOp wrapper of *this providing additional partial reduction operations * diff --git a/Eigen/src/Core/util/Macros.h b/Eigen/src/Core/util/Macros.h index d187257a6..b90c88ed4 100644 --- a/Eigen/src/Core/util/Macros.h +++ b/Eigen/src/Core/util/Macros.h @@ -760,8 +760,13 @@ namespace Eigen { # define EIGEN_TRY try # define EIGEN_CATCH(X) catch (X) #else -# define EIGEN_THROW_X(X) std::abort() -# define EIGEN_THROW std::abort() +# ifdef __CUDA_ARCH__ +# define EIGEN_THROW_X(X) asm("trap;") return {} +# define EIGEN_THROW asm("trap;"); return {} +# else +# define EIGEN_THROW_X(X) std::abort() +# define EIGEN_THROW std::abort() +# endif # define EIGEN_TRY if (true) # define EIGEN_CATCH(X) else #endif diff --git a/Eigen/src/Core/util/Memory.h b/Eigen/src/Core/util/Memory.h index 62f329984..73287d6ca 100644 --- a/Eigen/src/Core/util/Memory.h +++ b/Eigen/src/Core/util/Memory.h @@ -559,18 +559,18 @@ inline Index first_multiple(Index size, Index base) // use memcpy on trivial types, i.e., on types that does not require an initialization ctor. template<typename T, bool UseMemcpy> struct smart_copy_helper; -template<typename T> void smart_copy(const T* start, const T* end, T* target) +template<typename T> EIGEN_DEVICE_FUNC void smart_copy(const T* start, const T* end, T* target) { smart_copy_helper<T,!NumTraits<T>::RequireInitialization>::run(start, end, target); } template<typename T> struct smart_copy_helper<T,true> { - static inline void run(const T* start, const T* end, T* target) + EIGEN_DEVICE_FUNC static inline void run(const T* start, const T* end, T* target) { memcpy(target, start, std::ptrdiff_t(end)-std::ptrdiff_t(start)); } }; template<typename T> struct smart_copy_helper<T,false> { - static inline void run(const T* start, const T* end, T* target) + EIGEN_DEVICE_FUNC static inline void run(const T* start, const T* end, T* target) { std::copy(start, end, target); } }; diff --git a/Eigen/src/Core/util/Meta.h b/Eigen/src/Core/util/Meta.h index 974f11516..7c8932511 100644 --- a/Eigen/src/Core/util/Meta.h +++ b/Eigen/src/Core/util/Meta.h @@ -117,6 +117,10 @@ template<typename T> struct enable_if<true,T> { typedef T type; }; #if defined(__CUDA_ARCH__) +#if !defined(__FLT_EPSILON__) +#define __FLT_EPSILON__ FLT_EPSILON +#define __DBL_EPSILON__ DBL_EPSILON +#endif namespace device { diff --git a/doc/snippets/compile_snippet.cpp.in b/doc/snippets/compile_snippet.cpp.in index 82ae89162..fdae39bcf 100644 --- a/doc/snippets/compile_snippet.cpp.in +++ b/doc/snippets/compile_snippet.cpp.in @@ -1,6 +1,11 @@ #include <Eigen/Eigen> #include <iostream> +#ifndef M_PI +#define M_PI 3.1415926535897932384626433832795 +#endif + + using namespace Eigen; using namespace std; |