From 5adcc6c7b48b7a213af91bc123a02ab87751228e Mon Sep 17 00:00:00 2001 From: Gael Guennebaud Date: Thu, 7 Feb 2013 19:06:14 +0100 Subject: Add support for NVCC5: most of the Core and part of LU are callable from CUDA code. Still a lot to do. --- Eigen/Core | 10 ++++ Eigen/src/Core/Array.h | 19 +++++- Eigen/src/Core/Assign.h | 16 +++++ Eigen/src/Core/Block.h | 44 ++++++++++---- Eigen/src/Core/CommaInitializer.h | 6 ++ Eigen/src/Core/CwiseBinaryOp.h | 8 +++ Eigen/src/Core/CwiseUnaryOp.h | 9 +++ Eigen/src/Core/DenseBase.h | 49 +++++++++++----- Eigen/src/Core/DenseCoeffsBase.h | 33 +++++++++++ Eigen/src/Core/DenseStorage.h | 32 +++++----- Eigen/src/Core/EigenBase.h | 25 ++++++-- Eigen/src/Core/Functors.h | 91 +++++++++++++++-------------- Eigen/src/Core/GeneralProduct.h | 3 +- Eigen/src/Core/Map.h | 5 ++ Eigen/src/Core/MapBase.h | 21 +++++-- Eigen/src/Core/Matrix.h | 22 ++++++- Eigen/src/Core/MatrixBase.h | 30 ++++++++-- Eigen/src/Core/NoAlias.h | 7 +++ Eigen/src/Core/PlainObjectBase.h | 30 ++++++++++ Eigen/src/Core/Redux.h | 4 ++ Eigen/src/Core/ReturnByValue.h | 5 +- Eigen/src/Core/Stride.h | 13 +++-- Eigen/src/Core/Transpose.h | 17 ++++-- Eigen/src/Core/VectorBlock.h | 2 + Eigen/src/Core/products/CoeffBasedProduct.h | 19 ++++-- Eigen/src/Core/util/Macros.h | 2 +- Eigen/src/Core/util/XprHelper.h | 33 ++++++----- Eigen/src/LU/FullPivLU.h | 2 + Eigen/src/LU/Inverse.h | 23 +++++++- Eigen/src/LU/PartialPivLU.h | 5 ++ Eigen/src/plugins/ArrayCwiseUnaryOps.h | 14 +++++ Eigen/src/plugins/BlockMethods.h | 60 +++++++++++++++++++ Eigen/src/plugins/CommonCwiseBinaryOps.h | 1 + Eigen/src/plugins/CommonCwiseUnaryOps.h | 14 +++++ Eigen/src/plugins/MatrixCwiseBinaryOps.h | 8 +++ Eigen/src/plugins/MatrixCwiseUnaryOps.h | 5 ++ 36 files changed, 550 insertions(+), 137 deletions(-) diff --git a/Eigen/Core b/Eigen/Core index d5b286e53..d8f62c825 100644 --- a/Eigen/Core +++ b/Eigen/Core @@ -49,6 +49,16 @@ #endif #endif +// Handle NVCC/CUDA +#ifdef __CUDACC__ + // Do not try to vectorize on CUDA! + #define EIGEN_DONT_VECTORIZE + // Do not try asserts on CUDA! + #define EIGEN_NO_DEBUG + // All functions callable from CUDA code must be qualified with __device__ + #define EIGEN_DEVICE_FUNC __host__ __device__ +#endif + #ifndef EIGEN_DONT_VECTORIZE #if defined (EIGEN_SSE2_ON_NON_MSVC_BUT_NOT_OLD_GCC) || defined(EIGEN_SSE2_ON_MSVC_2008_OR_LATER) diff --git a/Eigen/src/Core/Array.h b/Eigen/src/Core/Array.h index 539e1d22b..707a9d7f2 100644 --- a/Eigen/src/Core/Array.h +++ b/Eigen/src/Core/Array.h @@ -69,6 +69,7 @@ class Array * the usage of 'using'. This should be done only for operator=. */ template + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Array& operator=(const EigenBase &other) { return Base::operator=(other); @@ -84,6 +85,7 @@ class Array * remain row-vectors and vectors remain vectors. */ template + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Array& operator=(const ArrayBase& other) { return Base::_set(other); @@ -92,6 +94,7 @@ class Array /** This is a special case of the templated operator=. Its purpose is to * prevent a default operator= from hiding the templated operator=. */ + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Array& operator=(const Array& other) { return Base::_set(other); @@ -107,6 +110,7 @@ class Array * * \sa resize(Index,Index) */ + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE explicit Array() : Base() { Base::_check_template_params(); @@ -116,6 +120,7 @@ class Array #ifndef EIGEN_PARSED_BY_DOXYGEN // FIXME is it still needed ?? /** \internal */ + EIGEN_DEVICE_FUNC Array(internal::constructor_without_unaligned_array_assert) : Base(internal::constructor_without_unaligned_array_assert()) { @@ -130,6 +135,7 @@ class Array * it is redundant to pass the dimension here, so it makes more sense to use the default * constructor Matrix() instead. */ + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE explicit Array(Index dim) : Base(dim, RowsAtCompileTime == 1 ? 1 : dim, ColsAtCompileTime == 1 ? 1 : dim) { @@ -142,6 +148,7 @@ class Array #ifndef EIGEN_PARSED_BY_DOXYGEN template + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Array(const T0& val0, const T1& val1) { Base::_check_template_params(); @@ -159,6 +166,7 @@ class Array #endif /** constructs an initialized 3D vector with given coefficients */ + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Array(const Scalar& val0, const Scalar& val1, const Scalar& val2) { Base::_check_template_params(); @@ -168,6 +176,7 @@ class Array m_storage.data()[2] = val2; } /** constructs an initialized 4D vector with given coefficients */ + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Array(const Scalar& val0, const Scalar& val1, const Scalar& val2, const Scalar& val3) { Base::_check_template_params(); @@ -178,10 +187,11 @@ class Array m_storage.data()[3] = val3; } - explicit Array(const Scalar *data); + EIGEN_DEVICE_FUNC explicit Array(const Scalar *data); /** Constructor copying the value of the expression \a other */ template + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Array(const ArrayBase& other) : Base(other.rows() * other.cols(), other.rows(), other.cols()) { @@ -189,6 +199,7 @@ class Array Base::_set_noalias(other); } /** Copy constructor */ + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Array(const Array& other) : Base(other.rows() * other.cols(), other.rows(), other.cols()) { @@ -197,6 +208,7 @@ class Array } /** Copy constructor with in-place evaluation */ template + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Array(const ReturnByValue& other) { Base::_check_template_params(); @@ -206,6 +218,7 @@ class Array /** \sa MatrixBase::operator=(const EigenBase&) */ template + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Array(const EigenBase &other) : Base(other.derived().rows() * other.derived().cols(), other.derived().rows(), other.derived().cols()) { @@ -221,8 +234,8 @@ class Array void swap(ArrayBase const & other) { this->_swap(other.derived()); } - inline Index innerStride() const { return 1; } - inline Index outerStride() const { return this->innerSize(); } + EIGEN_DEVICE_FUNC inline Index innerStride() const { return 1; } + EIGEN_DEVICE_FUNC inline Index outerStride() const { return this->innerSize(); } #ifdef EIGEN_ARRAY_PLUGIN #include EIGEN_ARRAY_PLUGIN diff --git a/Eigen/src/Core/Assign.h b/Eigen/src/Core/Assign.h index cd29a88f0..dc9b55fa4 100644 --- a/Eigen/src/Core/Assign.h +++ b/Eigen/src/Core/Assign.h @@ -139,6 +139,7 @@ struct assign_DefaultTraversal_CompleteUnrolling inner = Index % Derived1::InnerSizeAtCompileTime }; + EIGEN_DEVICE_FUNC static EIGEN_STRONG_INLINE void run(Derived1 &dst, const Derived2 &src) { dst.copyCoeffByOuterInner(outer, inner, src); @@ -149,12 +150,14 @@ struct assign_DefaultTraversal_CompleteUnrolling template struct assign_DefaultTraversal_CompleteUnrolling { + EIGEN_DEVICE_FUNC static EIGEN_STRONG_INLINE void run(Derived1 &, const Derived2 &) {} }; template struct assign_DefaultTraversal_InnerUnrolling { + EIGEN_DEVICE_FUNC static EIGEN_STRONG_INLINE void run(Derived1 &dst, const Derived2 &src, int outer) { dst.copyCoeffByOuterInner(outer, Index, src); @@ -165,6 +168,7 @@ struct assign_DefaultTraversal_InnerUnrolling template struct assign_DefaultTraversal_InnerUnrolling { + EIGEN_DEVICE_FUNC static EIGEN_STRONG_INLINE void run(Derived1 &, const Derived2 &, int) {} }; @@ -175,6 +179,7 @@ struct assign_DefaultTraversal_InnerUnrolling template struct assign_LinearTraversal_CompleteUnrolling { + EIGEN_DEVICE_FUNC static EIGEN_STRONG_INLINE void run(Derived1 &dst, const Derived2 &src) { dst.copyCoeff(Index, src); @@ -185,6 +190,7 @@ struct assign_LinearTraversal_CompleteUnrolling template struct assign_LinearTraversal_CompleteUnrolling { + EIGEN_DEVICE_FUNC static EIGEN_STRONG_INLINE void run(Derived1 &, const Derived2 &) {} }; @@ -249,6 +255,7 @@ struct assign_impl; template struct assign_impl { + EIGEN_DEVICE_FUNC static inline void run(Derived1 &, const Derived2 &) { } }; @@ -256,6 +263,7 @@ template struct assign_impl { typedef typename Derived1::Index Index; + EIGEN_DEVICE_FUNC static inline void run(Derived1 &dst, const Derived2 &src) { const Index innerSize = dst.innerSize(); @@ -269,6 +277,7 @@ struct assign_impl template struct assign_impl { + EIGEN_DEVICE_FUNC static EIGEN_STRONG_INLINE void run(Derived1 &dst, const Derived2 &src) { assign_DefaultTraversal_CompleteUnrolling @@ -280,6 +289,7 @@ template struct assign_impl { typedef typename Derived1::Index Index; + EIGEN_DEVICE_FUNC static EIGEN_STRONG_INLINE void run(Derived1 &dst, const Derived2 &src) { const Index outerSize = dst.outerSize(); @@ -297,6 +307,7 @@ template struct assign_impl { typedef typename Derived1::Index Index; + EIGEN_DEVICE_FUNC static inline void run(Derived1 &dst, const Derived2 &src) { const Index size = dst.size(); @@ -308,6 +319,7 @@ struct assign_impl template struct assign_impl { + EIGEN_DEVICE_FUNC static EIGEN_STRONG_INLINE void run(Derived1 &dst, const Derived2 &src) { assign_LinearTraversal_CompleteUnrolling @@ -519,18 +531,22 @@ struct assign_selector; template struct assign_selector { + EIGEN_DEVICE_FUNC static EIGEN_STRONG_INLINE Derived& run(Derived& dst, const OtherDerived& other) { return dst.lazyAssign(other.derived()); } }; template struct assign_selector { + EIGEN_DEVICE_FUNC static EIGEN_STRONG_INLINE Derived& run(Derived& dst, const OtherDerived& other) { return dst.lazyAssign(other.eval()); } }; template struct assign_selector { + EIGEN_DEVICE_FUNC static EIGEN_STRONG_INLINE Derived& run(Derived& dst, const OtherDerived& other) { return dst.lazyAssign(other.transpose()); } }; template struct assign_selector { + EIGEN_DEVICE_FUNC static EIGEN_STRONG_INLINE Derived& run(Derived& dst, const OtherDerived& other) { return dst.lazyAssign(other.transpose().eval()); } }; diff --git a/Eigen/src/Core/Block.h b/Eigen/src/Core/Block.h index 5f6426517..fbc2cf227 100644 --- a/Eigen/src/Core/Block.h +++ b/Eigen/src/Core/Block.h @@ -111,6 +111,7 @@ template class /** Column or Row constructor */ + EIGEN_DEVICE_FUNC inline Block(XprType& xpr, Index i) : Impl(xpr,i) { eigen_assert( (i>=0) && ( @@ -120,6 +121,7 @@ template class /** Fixed-size constructor */ + EIGEN_DEVICE_FUNC inline Block(XprType& xpr, Index a_startRow, Index a_startCol) : Impl(xpr, a_startRow, a_startCol) { @@ -130,6 +132,7 @@ template class /** Dynamic-size constructor */ + EIGEN_DEVICE_FUNC inline Block(XprType& xpr, Index a_startRow, Index a_startCol, Index blockRows, Index blockCols) @@ -153,8 +156,9 @@ class BlockImpl public: typedef Impl Base; EIGEN_INHERIT_ASSIGNMENT_OPERATORS(BlockImpl) - inline BlockImpl(XprType& xpr, Index i) : Impl(xpr,i) {} - inline BlockImpl(XprType& xpr, Index a_startRow, Index a_startCol) : Impl(xpr, a_startRow, a_startCol) {} + EIGEN_DEVICE_FUNC inline BlockImpl(XprType& xpr, Index i) : Impl(xpr,i) {} + EIGEN_DEVICE_FUNC inline BlockImpl(XprType& xpr, Index a_startRow, Index a_startCol) : Impl(xpr, a_startRow, a_startCol) {} + EIGEN_DEVICE_FUNC inline BlockImpl(XprType& xpr, Index a_startRow, Index a_startCol, Index blockRows, Index blockCols) : Impl(xpr, a_startRow, a_startCol, blockRows, blockCols) {} }; @@ -176,6 +180,7 @@ template::type& nestedExpression() const + EIGEN_DEVICE_FUNC + const typename internal::remove_all::type& nestedExpression() const { return m_xpr; } - Index startRow() const + EIGEN_DEVICE_FUNC + Index startRow() const { return m_startRow.value(); } + EIGEN_DEVICE_FUNC Index startCol() const { return m_startCol.value(); @@ -322,6 +338,7 @@ class BlockImpl_dense /** Column or Row constructor */ + EIGEN_DEVICE_FUNC inline BlockImpl_dense(XprType& xpr, Index i) : Base(internal::const_cast_ptr(&xpr.coeffRef( (BlockRows==1) && (BlockCols==XprType::ColsAtCompileTime) ? i : 0, @@ -335,6 +352,7 @@ class BlockImpl_dense /** Fixed-size constructor */ + EIGEN_DEVICE_FUNC inline BlockImpl_dense(XprType& xpr, Index startRow, Index startCol) : Base(internal::const_cast_ptr(&xpr.coeffRef(startRow,startCol))), m_xpr(xpr) { @@ -343,6 +361,7 @@ class BlockImpl_dense /** Dynamic-size constructor */ + EIGEN_DEVICE_FUNC inline BlockImpl_dense(XprType& xpr, Index startRow, Index startCol, Index blockRows, Index blockCols) @@ -352,12 +371,14 @@ class BlockImpl_dense init(); } - const typename internal::remove_all::type& nestedExpression() const + EIGEN_DEVICE_FUNC + const typename internal::remove_all::type& nestedExpression() const { return m_xpr; } /** \sa MapBase::innerStride() */ + EIGEN_DEVICE_FUNC inline Index innerStride() const { return internal::traits::HasSameStorageOrderAsXprType @@ -366,6 +387,7 @@ class BlockImpl_dense } /** \sa MapBase::outerStride() */ + EIGEN_DEVICE_FUNC inline Index outerStride() const { return m_outerStride; @@ -379,6 +401,7 @@ class BlockImpl_dense #ifndef EIGEN_PARSED_BY_DOXYGEN /** \internal used by allowAligned() */ + EIGEN_DEVICE_FUNC inline BlockImpl_dense(XprType& xpr, const Scalar* data, Index blockRows, Index blockCols) : Base(data, blockRows, blockCols), m_xpr(xpr) { @@ -387,6 +410,7 @@ class BlockImpl_dense #endif protected: + EIGEN_DEVICE_FUNC void init() { m_outerStride = internal::traits::HasSameStorageOrderAsXprType diff --git a/Eigen/src/Core/CommaInitializer.h b/Eigen/src/Core/CommaInitializer.h index f20c1774c..1f801e2a0 100644 --- a/Eigen/src/Core/CommaInitializer.h +++ b/Eigen/src/Core/CommaInitializer.h @@ -30,6 +30,7 @@ struct CommaInitializer typedef typename XprType::Scalar Scalar; typedef typename XprType::Index Index; + EIGEN_DEVICE_FUNC inline CommaInitializer(XprType& xpr, const Scalar& s) : m_xpr(xpr), m_row(0), m_col(1), m_currentBlockRows(1) { @@ -37,6 +38,7 @@ struct CommaInitializer } template + EIGEN_DEVICE_FUNC inline CommaInitializer(XprType& xpr, const DenseBase& other) : m_xpr(xpr), m_row(0), m_col(other.cols()), m_currentBlockRows(other.rows()) { @@ -44,6 +46,7 @@ struct CommaInitializer } /* inserts a scalar value in the target matrix */ + EIGEN_DEVICE_FUNC CommaInitializer& operator,(const Scalar& s) { if (m_col==m_xpr.cols()) @@ -63,6 +66,7 @@ struct CommaInitializer /* inserts a matrix expression in the target matrix */ template + EIGEN_DEVICE_FUNC CommaInitializer& operator,(const DenseBase& other) { if(other.cols()==0 || other.rows()==0) @@ -88,6 +92,7 @@ struct CommaInitializer return *this; } + EIGEN_DEVICE_FUNC inline ~CommaInitializer() { eigen_assert((m_row+m_currentBlockRows) == m_xpr.rows() @@ -102,6 +107,7 @@ struct CommaInitializer * quaternion.fromRotationMatrix((Matrix3f() << axis0, axis1, axis2).finished()); * \endcode */ + EIGEN_DEVICE_FUNC inline XprType& finished() { return m_xpr; } XprType& m_xpr; // target expression diff --git a/Eigen/src/Core/CwiseBinaryOp.h b/Eigen/src/Core/CwiseBinaryOp.h index 686c2afa3..532b2b96e 100644 --- a/Eigen/src/Core/CwiseBinaryOp.h +++ b/Eigen/src/Core/CwiseBinaryOp.h @@ -122,6 +122,7 @@ class CwiseBinaryOp : internal::no_assignment_operator, typedef typename internal::remove_reference::type _LhsNested; typedef typename internal::remove_reference::type _RhsNested; + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE CwiseBinaryOp(const Lhs& aLhs, const Rhs& aRhs, const BinaryOp& func = BinaryOp()) : m_lhs(aLhs), m_rhs(aRhs), m_functor(func) { @@ -131,6 +132,7 @@ class CwiseBinaryOp : internal::no_assignment_operator, eigen_assert(aLhs.rows() == aRhs.rows() && aLhs.cols() == aRhs.cols()); } + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Index rows() const { // return the fixed size type if available to enable compile time optimizations if (internal::traits::type>::RowsAtCompileTime==Dynamic) @@ -138,6 +140,7 @@ class CwiseBinaryOp : internal::no_assignment_operator, else return m_lhs.rows(); } + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Index cols() const { // return the fixed size type if available to enable compile time optimizations if (internal::traits::type>::ColsAtCompileTime==Dynamic) @@ -147,10 +150,13 @@ class CwiseBinaryOp : internal::no_assignment_operator, } /** \returns the left hand side nested expression */ + EIGEN_DEVICE_FUNC const _LhsNested& lhs() const { return m_lhs; } /** \returns the right hand side nested expression */ + EIGEN_DEVICE_FUNC const _RhsNested& rhs() const { return m_rhs; } /** \returns the functor representing the binary operation */ + EIGEN_DEVICE_FUNC const BinaryOp& functor() const { return m_functor; } protected: @@ -169,6 +175,7 @@ class CwiseBinaryOpImpl typedef typename internal::dense_xpr_base >::type Base; EIGEN_DENSE_PUBLIC_INTERFACE( Derived ) + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE const Scalar coeff(Index rowId, Index colId) const { return derived().functor()(derived().lhs().coeff(rowId, colId), @@ -182,6 +189,7 @@ class CwiseBinaryOpImpl derived().rhs().template packet(rowId, colId)); } + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE const Scalar coeff(Index index) const { return derived().functor()(derived().lhs().coeff(index), diff --git a/Eigen/src/Core/CwiseUnaryOp.h b/Eigen/src/Core/CwiseUnaryOp.h index f2de749f9..aa7df197f 100644 --- a/Eigen/src/Core/CwiseUnaryOp.h +++ b/Eigen/src/Core/CwiseUnaryOp.h @@ -64,20 +64,26 @@ class CwiseUnaryOp : internal::no_assignment_operator, typedef typename CwiseUnaryOpImpl::StorageKind>::Base Base; EIGEN_GENERIC_PUBLIC_INTERFACE(CwiseUnaryOp) + EIGEN_DEVICE_FUNC inline CwiseUnaryOp(const XprType& xpr, const UnaryOp& func = UnaryOp()) : m_xpr(xpr), m_functor(func) {} + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Index rows() const { return m_xpr.rows(); } + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Index cols() const { return m_xpr.cols(); } /** \returns the functor representing the unary operation */ + EIGEN_DEVICE_FUNC const UnaryOp& functor() const { return m_functor; } /** \returns the nested expression */ + EIGEN_DEVICE_FUNC const typename internal::remove_all::type& nestedExpression() const { return m_xpr; } /** \returns the nested expression */ + EIGEN_DEVICE_FUNC typename internal::remove_all::type& nestedExpression() { return m_xpr.const_cast_derived(); } @@ -98,6 +104,7 @@ class CwiseUnaryOpImpl typedef typename internal::dense_xpr_base >::type Base; EIGEN_DENSE_PUBLIC_INTERFACE(Derived) + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE const Scalar coeff(Index rowId, Index colId) const { return derived().functor()(derived().nestedExpression().coeff(rowId, colId)); @@ -109,12 +116,14 @@ class CwiseUnaryOpImpl return derived().functor().packetOp(derived().nestedExpression().template packet(rowId, colId)); } + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE const Scalar coeff(Index index) const { return derived().functor()(derived().nestedExpression().coeff(index)); } template + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE PacketScalar packet(Index index) const { return derived().functor().packetOp(derived().nestedExpression().template packet(index)); diff --git a/Eigen/src/Core/DenseBase.h b/Eigen/src/Core/DenseBase.h index 62c73f1a9..12780354b 100644 --- a/Eigen/src/Core/DenseBase.h +++ b/Eigen/src/Core/DenseBase.h @@ -172,6 +172,7 @@ template class DenseBase /** \returns the number of nonzero coefficients which is in practice the number * of stored coefficients. */ + EIGEN_DEVICE_FUNC inline Index nonZeros() const { return size(); } /** \returns true if either the number of rows or the number of columns is equal to 1. * In other words, this function returns @@ -183,6 +184,7 @@ template class DenseBase * \note For a vector, this returns just 1. For a matrix (non-vector), this is the major dimension * with respect to the \ref TopicStorageOrders "storage order", i.e., the number of columns for a * column-major matrix, and the number of rows for a row-major matrix. */ + EIGEN_DEVICE_FUNC Index outerSize() const { return IsVectorAtCompileTime ? 1 @@ -194,6 +196,7 @@ template class DenseBase * \note For a vector, this is just the size. For a matrix (non-vector), this is the minor dimension * with respect to the \ref TopicStorageOrders "storage order", i.e., the number of rows for a * column-major matrix, and the number of columns for a row-major matrix. */ + EIGEN_DEVICE_FUNC Index innerSize() const { return IsVectorAtCompileTime ? this->size() @@ -204,6 +207,7 @@ template class DenseBase * Matrix::resize() and Array::resize(). The present method only asserts that the new size equals the old size, and does * nothing else. */ + EIGEN_DEVICE_FUNC void resize(Index newSize) { EIGEN_ONLY_USED_FOR_DEBUG(newSize); @@ -214,6 +218,7 @@ template class DenseBase * Matrix::resize() and Array::resize(). The present method only asserts that the new size equals the old size, and does * nothing else. */ + EIGEN_DEVICE_FUNC void resize(Index nbRows, Index nbCols) { EIGEN_ONLY_USED_FOR_DEBUG(nbRows); @@ -237,42 +242,54 @@ template class DenseBase /** Copies \a other into *this. \returns a reference to *this. */ template + EIGEN_DEVICE_FUNC Derived& operator=(const DenseBase& other); /** Special case of the template operator=, in order to prevent the compiler * from generating a default operator= (issue hit with g++ 4.1) */ + EIGEN_DEVICE_FUNC Derived& operator=(const DenseBase& other); template + EIGEN_DEVICE_FUNC Derived& operator=(const EigenBase &other); template + EIGEN_DEVICE_FUNC Derived& operator+=(const EigenBase &other); template + EIGEN_DEVICE_FUNC Derived& operator-=(const EigenBase &other); template + EIGEN_DEVICE_FUNC Derived& operator=(const ReturnByValue& func); #ifndef EIGEN_PARSED_BY_DOXYGEN /** Copies \a other into *this without evaluating other. \returns a reference to *this. */ template + EIGEN_DEVICE_FUNC Derived& lazyAssign(const DenseBase& other); #endif // not EIGEN_PARSED_BY_DOXYGEN + EIGEN_DEVICE_FUNC CommaInitializer operator<< (const Scalar& s); template const Flagged flagged() const; template + EIGEN_DEVICE_FUNC CommaInitializer operator<< (const DenseBase& other); + EIGEN_DEVICE_FUNC Eigen::Transpose transpose(); typedef const Transpose ConstTransposeReturnType; + EIGEN_DEVICE_FUNC ConstTransposeReturnType transpose() const; + EIGEN_DEVICE_FUNC void transposeInPlace(); #ifndef EIGEN_NO_DEBUG protected: @@ -346,6 +363,7 @@ template class DenseBase * Notice that in the case of a plain matrix or vector (not an expression) this function just returns * a const reference, in order to avoid a useless copy. */ + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE EvalReturnType eval() const { // Even though MSVC does not honor strong inlining when the return type @@ -380,14 +398,14 @@ template class DenseBase template inline const typename internal::conditional,Derived&>::type forceAlignedAccessIf() const; template inline typename internal::conditional,Derived&>::type forceAlignedAccessIf(); - Scalar sum() const; - Scalar mean() const; - Scalar trace() const; + EIGEN_DEVICE_FUNC Scalar sum() const; + EIGEN_DEVICE_FUNC Scalar mean() const; + EIGEN_DEVICE_FUNC Scalar trace() const; - Scalar prod() const; + EIGEN_DEVICE_FUNC Scalar prod() const; - typename internal::traits::Scalar minCoeff() const; - typename internal::traits::Scalar maxCoeff() const; + EIGEN_DEVICE_FUNC typename internal::traits::Scalar minCoeff() const; + EIGEN_DEVICE_FUNC typename internal::traits::Scalar maxCoeff() const; template typename internal::traits::Scalar minCoeff(IndexType* row, IndexType* col) const; @@ -399,15 +417,18 @@ template class DenseBase typename internal::traits::Scalar maxCoeff(IndexType* index) const; template + EIGEN_DEVICE_FUNC typename internal::result_of::Scalar)>::type redux(const BinaryOp& func) const; template + EIGEN_DEVICE_FUNC void visit(Visitor& func) const; inline const WithFormat format(const IOFormat& fmt) const; /** \returns the unique coefficient of a 1x1 expression */ + EIGEN_DEVICE_FUNC CoeffReturnType value() const { EIGEN_STATIC_ASSERT_SIZE_1x1(Derived) @@ -417,8 +438,8 @@ template class DenseBase /////////// Array module /////////// - bool all(void) const; - bool any(void) const; + bool all() const; + bool any() const; Index count() const; typedef VectorwiseOp RowwiseReturnType; @@ -480,14 +501,16 @@ template class DenseBase // disable the use of evalTo for dense objects with a nice compilation error - template inline void evalTo(Dest& ) const + template + EIGEN_DEVICE_FUNC + inline void evalTo(Dest& ) const { EIGEN_STATIC_ASSERT((internal::is_same::value),THE_EVAL_EVALTO_FUNCTION_SHOULD_NEVER_BE_CALLED_FOR_DENSE_OBJECTS); } protected: /** Default constructor. Do nothing. */ - DenseBase() + EIGEN_DEVICE_FUNC DenseBase() { /* Just checks for self-consistency of the flags. * Only do it when debugging Eigen, as this borders on paranoiac and could slow compilation down @@ -500,9 +523,9 @@ template class DenseBase } private: - explicit DenseBase(int); - DenseBase(int,int); - template explicit DenseBase(const DenseBase&); + EIGEN_DEVICE_FUNC explicit DenseBase(int); + EIGEN_DEVICE_FUNC DenseBase(int,int); + template EIGEN_DEVICE_FUNC explicit DenseBase(const DenseBase&); }; } // end namespace Eigen diff --git a/Eigen/src/Core/DenseCoeffsBase.h b/Eigen/src/Core/DenseCoeffsBase.h index 3c890f215..efabb5e67 100644 --- a/Eigen/src/Core/DenseCoeffsBase.h +++ b/Eigen/src/Core/DenseCoeffsBase.h @@ -61,6 +61,7 @@ class DenseCoeffsBase : public EigenBase using Base::size; using Base::derived; + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Index rowIndexByOuterInner(Index outer, Index inner) const { return int(Derived::RowsAtCompileTime) == 1 ? 0 @@ -69,6 +70,7 @@ class DenseCoeffsBase : public EigenBase : inner; } + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Index colIndexByOuterInner(Index outer, Index inner) const { return int(Derived::ColsAtCompileTime) == 1 ? 0 @@ -91,6 +93,7 @@ class DenseCoeffsBase : public EigenBase * * \sa operator()(Index,Index) const, coeffRef(Index,Index), coeff(Index) const */ + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE CoeffReturnType coeff(Index row, Index col) const { eigen_internal_assert(row >= 0 && row < rows() @@ -98,6 +101,7 @@ class DenseCoeffsBase : public EigenBase return derived().coeff(row, col); } + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE CoeffReturnType coeffByOuterInner(Index outer, Index inner) const { return coeff(rowIndexByOuterInner(outer, inner), @@ -108,6 +112,7 @@ class DenseCoeffsBase : public EigenBase * * \sa operator()(Index,Index), operator[](Index) */ + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE CoeffReturnType operator()(Index row, Index col) const { eigen_assert(row >= 0 && row < rows() @@ -130,6 +135,7 @@ class DenseCoeffsBase : public EigenBase * \sa operator[](Index) const, coeffRef(Index), coeff(Index,Index) const */ + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE CoeffReturnType coeff(Index index) const { @@ -146,6 +152,7 @@ class DenseCoeffsBase : public EigenBase * z() const, w() const */ + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE CoeffReturnType operator[](Index index) const { @@ -167,6 +174,7 @@ class DenseCoeffsBase : public EigenBase * z() const, w() const */ + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE CoeffReturnType operator()(Index index) const { @@ -176,21 +184,25 @@ class DenseCoeffsBase : public EigenBase /** equivalent to operator[](0). */ + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE CoeffReturnType x() const { return (*this)[0]; } /** equivalent to operator[](1). */ + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE CoeffReturnType y() const { return (*this)[1]; } /** equivalent to operator[](2). */ + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE CoeffReturnType z() const { return (*this)[2]; } /** equivalent to operator[](3). */ + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE CoeffReturnType w() const { return (*this)[3]; } @@ -311,6 +323,7 @@ class DenseCoeffsBase : public DenseCoeffsBase= 0 && row < rows() @@ -318,6 +331,7 @@ class DenseCoeffsBase : public DenseCoeffsBase : public DenseCoeffsBase : public DenseCoeffsBase : public DenseCoeffsBase : public DenseCoeffsBase : public DenseCoeffsBase : public DenseCoeffsBase + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void copyCoeff(Index row, Index col, const DenseBase& other) { eigen_internal_assert(row >= 0 && row < rows() @@ -489,6 +512,7 @@ class DenseCoeffsBase : public DenseCoeffsBase + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void copyCoeff(Index index, const DenseBase& other) { eigen_internal_assert(index >= 0 && index < size()); @@ -497,6 +521,7 @@ class DenseCoeffsBase : public DenseCoeffsBase + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void copyCoeffByOuterInner(Index outer, Index inner, const DenseBase& other) { const Index row = rowIndexByOuterInner(outer,inner); @@ -581,6 +606,7 @@ class DenseCoeffsBase : public DenseCoeffsBase : public DenseCoeffsBase : public DenseCoeffsBase : public DenseCoeffsBase * * \sa outerStride(), rowStride(), colStride() */ + EIGEN_DEVICE_FUNC inline Index innerStride() const { return derived().innerStride(); @@ -662,6 +692,7 @@ class DenseCoeffsBase * * \sa innerStride(), rowStride(), colStride() */ + EIGEN_DEVICE_FUNC inline Index outerStride() const { return derived().outerStride(); @@ -677,6 +708,7 @@ class DenseCoeffsBase * * \sa innerStride(), outerStride(), colStride() */ + EIGEN_DEVICE_FUNC inline Index rowStride() const { return Derived::IsRowMajor ? outerStride() : innerStride(); @@ -686,6 +718,7 @@ class DenseCoeffsBase * * \sa innerStride(), outerStride(), rowStride() */ + EIGEN_DEVICE_FUNC inline Index colStride() const { return Derived::IsRowMajor ? innerStride() : outerStride(); diff --git a/Eigen/src/Core/DenseStorage.h b/Eigen/src/Core/DenseStorage.h index 894dcf2c1..203944620 100644 --- a/Eigen/src/Core/DenseStorage.h +++ b/Eigen/src/Core/DenseStorage.h @@ -36,12 +36,14 @@ struct plain_array { T array[Size]; - plain_array() + EIGEN_DEVICE_FUNC + plain_array() { EIGEN_STATIC_ASSERT(Size * sizeof(T) <= 128 * 128 * 8, OBJECT_ALLOCATED_ON_STACK_IS_TOO_BIG); } - plain_array(constructor_without_unaligned_array_assert) + EIGEN_DEVICE_FUNC + plain_array(constructor_without_unaligned_array_assert) { EIGEN_STATIC_ASSERT(Size * sizeof(T) <= 128 * 128 * 8, OBJECT_ALLOCATED_ON_STACK_IS_TOO_BIG); } @@ -73,12 +75,14 @@ struct plain_array { EIGEN_USER_ALIGN16 T array[Size]; + EIGEN_DEVICE_FUNC plain_array() { EIGEN_MAKE_UNALIGNED_ARRAY_ASSERT(0xf); EIGEN_STATIC_ASSERT(Size * sizeof(T) <= 128 * 128 * 8, OBJECT_ALLOCATED_ON_STACK_IS_TOO_BIG); } + EIGEN_DEVICE_FUNC plain_array(constructor_without_unaligned_array_assert) { EIGEN_STATIC_ASSERT(Size * sizeof(T) <= 128 * 128 * 8, OBJECT_ALLOCATED_ON_STACK_IS_TOO_BIG); @@ -89,8 +93,8 @@ template struct plain_array { EIGEN_USER_ALIGN16 T array[1]; - plain_array() {} - plain_array(constructor_without_unaligned_array_assert) {} + EIGEN_DEVICE_FUNC plain_array() {} + EIGEN_DEVICE_FUNC plain_array(constructor_without_unaligned_array_assert) {} }; } // end namespace internal @@ -114,17 +118,17 @@ template class DenseSt { internal::plain_array m_data; public: - inline explicit DenseStorage() {} - inline DenseStorage(internal::constructor_without_unaligned_array_assert) + EIGEN_DEVICE_FUNC inline explicit DenseStorage() {} + EIGEN_DEVICE_FUNC inline DenseStorage(internal::constructor_without_unaligned_array_assert) : m_data(internal::constructor_without_unaligned_array_assert()) {} - inline DenseStorage(DenseIndex,DenseIndex,DenseIndex) {} - inline void swap(DenseStorage& other) { std::swap(m_data,other.m_data); } - static inline DenseIndex rows(void) {return _Rows;} - static inline DenseIndex cols(void) {return _Cols;} - inline void conservativeResize(DenseIndex,DenseIndex,DenseIndex) {} - inline void resize(DenseIndex,DenseIndex,DenseIndex) {} - inline const T *data() const { return m_data.array; } - inline T *data() { return m_data.array; } + EIGEN_DEVICE_FUNC inline DenseStorage(DenseIndex,DenseIndex,DenseIndex) {} + EIGEN_DEVICE_FUNC inline void swap(DenseStorage& other) { std::swap(m_data,other.m_data); } + EIGEN_DEVICE_FUNC static inline DenseIndex rows(void) {return _Rows;} + EIGEN_DEVICE_FUNC static inline DenseIndex cols(void) {return _Cols;} + EIGEN_DEVICE_FUNC inline void conservativeResize(DenseIndex,DenseIndex,DenseIndex) {} + EIGEN_DEVICE_FUNC inline void resize(DenseIndex,DenseIndex,DenseIndex) {} + EIGEN_DEVICE_FUNC inline const T *data() const { return m_data.array; } + EIGEN_DEVICE_FUNC inline T *data() { return m_data.array; } }; // null matrix diff --git a/Eigen/src/Core/EigenBase.h b/Eigen/src/Core/EigenBase.h index 2b8dd1b70..a25e823ab 100644 --- a/Eigen/src/Core/EigenBase.h +++ b/Eigen/src/Core/EigenBase.h @@ -31,29 +31,40 @@ template struct EigenBase typedef typename internal::traits::Index Index; /** \returns a reference to the derived object */ + EIGEN_DEVICE_FUNC Derived& derived() { return *static_cast(this); } /** \returns a const reference to the derived object */ + EIGEN_DEVICE_FUNC const Derived& derived() const { return *static_cast(this); } + EIGEN_DEVICE_FUNC inline Derived& const_cast_derived() const { return *static_cast(const_cast(this)); } + EIGEN_DEVICE_FUNC inline const Derived& const_derived() const { return *static_cast(this); } /** \returns the number of rows. \sa cols(), RowsAtCompileTime */ + EIGEN_DEVICE_FUNC inline Index rows() const { return derived().rows(); } /** \returns the number of columns. \sa rows(), ColsAtCompileTime*/ + EIGEN_DEVICE_FUNC inline Index cols() const { return derived().cols(); } /** \returns the number of coefficients, which is rows()*cols(). * \sa rows(), cols(), SizeAtCompileTime. */ + EIGEN_DEVICE_FUNC inline Index size() const { return rows() * cols(); } /** \internal Don't use it, but do the equivalent: \code dst = *this; \endcode */ - template inline void evalTo(Dest& dst) const + template + EIGEN_DEVICE_FUNC + inline void evalTo(Dest& dst) const { derived().evalTo(dst); } /** \internal Don't use it, but do the equivalent: \code dst += *this; \endcode */ - template inline void addTo(Dest& dst) const + template + EIGEN_DEVICE_FUNC + inline void addTo(Dest& dst) const { // This is the default implementation, // derived class can reimplement it in a more optimized way. @@ -63,7 +74,9 @@ template struct EigenBase } /** \internal Don't use it, but do the equivalent: \code dst -= *this; \endcode */ - template inline void subTo(Dest& dst) const + template + EIGEN_DEVICE_FUNC + inline void subTo(Dest& dst) const { // This is the default implementation, // derived class can reimplement it in a more optimized way. @@ -73,7 +86,8 @@ template struct EigenBase } /** \internal Don't use it, but do the equivalent: \code dst.applyOnTheRight(*this); \endcode */ - template inline void applyThisOnTheRight(Dest& dst) const + template + EIGEN_DEVICE_FUNC inline void applyThisOnTheRight(Dest& dst) const { // This is the default implementation, // derived class can reimplement it in a more optimized way. @@ -81,7 +95,8 @@ template struct EigenBase } /** \internal Don't use it, but do the equivalent: \code dst.applyOnTheLeft(*this); \endcode */ - template inline void applyThisOnTheLeft(Dest& dst) const + template + EIGEN_DEVICE_FUNC inline void applyThisOnTheLeft(Dest& dst) const { // This is the default implementation, // derived class can reimplement it in a more optimized way. diff --git a/Eigen/src/Core/Functors.h b/Eigen/src/Core/Functors.h index 2a6c3c003..6b6b36656 100644 --- a/Eigen/src/Core/Functors.h +++ b/Eigen/src/Core/Functors.h @@ -23,7 +23,7 @@ namespace internal { */ template struct scalar_sum_op { EIGEN_EMPTY_STRUCT_CTOR(scalar_sum_op) - EIGEN_STRONG_INLINE const Scalar operator() (const Scalar& a, const Scalar& b) const { return a + b; } + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE const Scalar operator() (const Scalar& a, const Scalar& b) const { return a + b; } template EIGEN_STRONG_INLINE const Packet packetOp(const Packet& a, const Packet& b) const { return internal::padd(a,b); } @@ -51,7 +51,7 @@ template struct scalar_product_op { }; typedef typename scalar_product_traits::ReturnType result_type; EIGEN_EMPTY_STRUCT_CTOR(scalar_product_op) - EIGEN_STRONG_INLINE const result_type operator() (const LhsScalar& a, const RhsScalar& b) const { return a * b; } + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE const result_type operator() (const LhsScalar& a, const RhsScalar& b) const { return a * b; } template EIGEN_STRONG_INLINE const Packet packetOp(const Packet& a, const Packet& b) const { return internal::pmul(a,b); } @@ -81,7 +81,7 @@ template struct scalar_conj_product_op { typedef typename scalar_product_traits::ReturnType result_type; EIGEN_EMPTY_STRUCT_CTOR(scalar_conj_product_op) - EIGEN_STRONG_INLINE const result_type operator() (const LhsScalar& a, const RhsScalar& b) const + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE const result_type operator() (const LhsScalar& a, const RhsScalar& b) const { return conj_helper().pmul(a,b); } template @@ -103,7 +103,7 @@ struct functor_traits > { */ template struct scalar_min_op { EIGEN_EMPTY_STRUCT_CTOR(scalar_min_op) - EIGEN_STRONG_INLINE const Scalar operator() (const Scalar& a, const Scalar& b) const { using std::min; return (min)(a, b); } + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE const Scalar operator() (const Scalar& a, const Scalar& b) const { using std::min; return (min)(a, b); } template EIGEN_STRONG_INLINE const Packet packetOp(const Packet& a, const Packet& b) const { return internal::pmin(a,b); } @@ -126,7 +126,7 @@ struct functor_traits > { */ template struct scalar_max_op { EIGEN_EMPTY_STRUCT_CTOR(scalar_max_op) - EIGEN_STRONG_INLINE const Scalar operator() (const Scalar& a, const Scalar& b) const { using std::max; return (max)(a, b); } + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE const Scalar operator() (const Scalar& a, const Scalar& b) const { using std::max; return (max)(a, b); } template EIGEN_STRONG_INLINE const Packet packetOp(const Packet& a, const Packet& b) const { return internal::pmax(a,b); } @@ -150,7 +150,7 @@ struct functor_traits > { template struct scalar_hypot_op { EIGEN_EMPTY_STRUCT_CTOR(scalar_hypot_op) // typedef typename NumTraits::Real result_type; - EIGEN_STRONG_INLINE const Scalar operator() (const Scalar& _x, const Scalar& _y) const + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE const Scalar operator() (const Scalar& _x, const Scalar& _y) const { using std::max; using std::min; @@ -170,7 +170,7 @@ struct functor_traits > { */ template struct scalar_binary_pow_op { EIGEN_EMPTY_STRUCT_CTOR(scalar_binary_pow_op) - inline Scalar operator() (const Scalar& a, const OtherScalar& b) const { return internal::pow(a, b); } + EIGEN_DEVICE_FUNC inline Scalar operator() (const Scalar& a, const OtherScalar& b) const { return internal::pow(a, b); } }; template struct functor_traits > { @@ -186,7 +186,7 @@ struct functor_traits > { */ template struct scalar_difference_op { EIGEN_EMPTY_STRUCT_CTOR(scalar_difference_op) - EIGEN_STRONG_INLINE const Scalar operator() (const Scalar& a, const Scalar& b) const { return a - b; } + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE const Scalar operator() (const Scalar& a, const Scalar& b) const { return a - b; } template EIGEN_STRONG_INLINE const Packet packetOp(const Packet& a, const Packet& b) const { return internal::psub(a,b); } @@ -211,7 +211,7 @@ template struct scalar_quotient_op { }; typedef typename scalar_product_traits::ReturnType result_type; EIGEN_EMPTY_STRUCT_CTOR(scalar_quotient_op) - EIGEN_STRONG_INLINE const result_type operator() (const LhsScalar& a, const RhsScalar& b) const { return a / b; } + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE const result_type operator() (const LhsScalar& a, const RhsScalar& b) const { return a / b; } template EIGEN_STRONG_INLINE const Packet packetOp(const Packet& a, const Packet& b) const { return internal::pdiv(a,b); } @@ -233,7 +233,7 @@ struct functor_traits > { */ struct scalar_boolean_and_op { EIGEN_EMPTY_STRUCT_CTOR(scalar_boolean_and_op) - EIGEN_STRONG_INLINE bool operator() (const bool& a, const bool& b) const { return a && b; } + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE bool operator() (const bool& a, const bool& b) const { return a && b; } }; template<> struct functor_traits { enum { @@ -249,7 +249,7 @@ template<> struct functor_traits { */ struct scalar_boolean_or_op { EIGEN_EMPTY_STRUCT_CTOR(scalar_boolean_or_op) - EIGEN_STRONG_INLINE bool operator() (const bool& a, const bool& b) const { return a || b; } + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE bool operator() (const bool& a, const bool& b) const { return a || b; } }; template<> struct functor_traits { enum { @@ -267,7 +267,7 @@ template<> struct functor_traits { */ template struct scalar_opposite_op { EIGEN_EMPTY_STRUCT_CTOR(scalar_opposite_op) - EIGEN_STRONG_INLINE const Scalar operator() (const Scalar& a) const { return -a; } + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE const Scalar operator() (const Scalar& a) const { return -a; } template EIGEN_STRONG_INLINE const Packet packetOp(const Packet& a) const { return internal::pnegate(a); } @@ -287,7 +287,7 @@ struct functor_traits > template struct scalar_abs_op { EIGEN_EMPTY_STRUCT_CTOR(scalar_abs_op) typedef typename NumTraits::Real result_type; - EIGEN_STRONG_INLINE const result_type operator() (const Scalar& a) const { using std::abs; return abs(a); } + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE const result_type operator() (const Scalar& a) const { using std::abs; return abs(a); } template EIGEN_STRONG_INLINE const Packet packetOp(const Packet& a) const { return internal::pabs(a); } @@ -309,7 +309,7 @@ struct functor_traits > template struct scalar_abs2_op { EIGEN_EMPTY_STRUCT_CTOR(scalar_abs2_op) typedef typename NumTraits::Real result_type; - EIGEN_STRONG_INLINE const result_type operator() (const Scalar& a) const { return internal::abs2(a); } + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE const result_type operator() (const Scalar& a) const { return internal::abs2(a); } template EIGEN_STRONG_INLINE const Packet packetOp(const Packet& a) const { return internal::pmul(a,a); } @@ -325,7 +325,7 @@ struct functor_traits > */ template struct scalar_conjugate_op { EIGEN_EMPTY_STRUCT_CTOR(scalar_conjugate_op) - EIGEN_STRONG_INLINE const Scalar operator() (const Scalar& a) const { using internal::conj; return conj(a); } + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE const Scalar operator() (const Scalar& a) const { using internal::conj; return conj(a); } template EIGEN_STRONG_INLINE const Packet packetOp(const Packet& a) const { return internal::pconj(a); } }; @@ -347,7 +347,7 @@ template struct scalar_cast_op { EIGEN_EMPTY_STRUCT_CTOR(scalar_cast_op) typedef NewType result_type; - EIGEN_STRONG_INLINE const NewType operator() (const Scalar& a) const { return cast(a); } + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE const NewType operator() (const Scalar& a) const { return cast(a); } }; template struct functor_traits > @@ -362,7 +362,7 @@ template struct scalar_real_op { EIGEN_EMPTY_STRUCT_CTOR(scalar_real_op) typedef typename NumTraits::Real result_type; - EIGEN_STRONG_INLINE result_type operator() (const Scalar& a) const { return internal::real(a); } + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE result_type operator() (const Scalar& a) const { return internal::real(a); } }; template struct functor_traits > @@ -377,7 +377,7 @@ template struct scalar_imag_op { EIGEN_EMPTY_STRUCT_CTOR(scalar_imag_op) typedef typename NumTraits::Real result_type; - EIGEN_STRONG_INLINE result_type operator() (const Scalar& a) const { return internal::imag(a); } + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE result_type operator() (const Scalar& a) const { return internal::imag(a); } }; template struct functor_traits > @@ -392,7 +392,7 @@ template struct scalar_real_ref_op { EIGEN_EMPTY_STRUCT_CTOR(scalar_real_ref_op) typedef typename NumTraits::Real result_type; - EIGEN_STRONG_INLINE result_type& operator() (const Scalar& a) const { return internal::real_ref(*const_cast(&a)); } + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE result_type& operator() (const Scalar& a) const { return internal::real_ref(*const_cast(&a)); } }; template struct functor_traits > @@ -407,7 +407,7 @@ template struct scalar_imag_ref_op { EIGEN_EMPTY_STRUCT_CTOR(scalar_imag_ref_op) typedef typename NumTraits::Real result_type; - EIGEN_STRONG_INLINE result_type& operator() (const Scalar& a) const { return internal::imag_ref(*const_cast(&a)); } + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE result_type& operator() (const Scalar& a) const { return internal::imag_ref(*const_cast(&a)); } }; template struct functor_traits > @@ -421,7 +421,7 @@ struct functor_traits > */ template struct scalar_exp_op { EIGEN_EMPTY_STRUCT_CTOR(scalar_exp_op) - inline const Scalar operator() (const Scalar& a) const { using std::exp; return exp(a); } + EIGEN_DEVICE_FUNC inline const Scalar operator() (const Scalar& a) const { using std::exp; return exp(a); } typedef typename packet_traits::type Packet; inline Packet packetOp(const Packet& a) const { return internal::pexp(a); } }; @@ -437,7 +437,7 @@ struct functor_traits > */ template struct scalar_log_op { EIGEN_EMPTY_STRUCT_CTOR(scalar_log_op) - inline const Scalar operator() (const Scalar& a) const { using std::log; return log(a); } + EIGEN_DEVICE_FUNC inline const Scalar operator() (const Scalar& a) const { using std::log; return log(a); } typedef typename packet_traits::type Packet; inline Packet packetOp(const Packet& a) const { return internal::plog(a); } }; @@ -462,8 +462,11 @@ template struct scalar_multiple_op { typedef typename packet_traits::type Packet; // FIXME default copy constructors seems bugged with std::complex<> + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE scalar_multiple_op(const scalar_multiple_op& other) : m_other(other.m_other) { } + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE scalar_multiple_op(const Scalar& other) : m_other(other) { } + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Scalar operator() (const Scalar& a) const { return a * m_other; } EIGEN_STRONG_INLINE const Packet packetOp(const Packet& a) const { return internal::pmul(a, pset1(m_other)); } @@ -478,7 +481,7 @@ struct scalar_multiple2_op { typedef typename scalar_product_traits::ReturnType result_type; EIGEN_STRONG_INLINE scalar_multiple2_op(const scalar_multiple2_op& other) : m_other(other.m_other) { } EIGEN_STRONG_INLINE scalar_multiple2_op(const Scalar2& other) : m_other(other) { } - EIGEN_STRONG_INLINE result_type operator() (const Scalar1& a) const { return a * m_other; } + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE result_type operator() (const Scalar1& a) const { return a * m_other; } typename add_const_on_value_type::Nested>::type m_other; }; template @@ -499,7 +502,7 @@ struct scalar_quotient1_op { // FIXME default copy constructors seems bugged with std::complex<> EIGEN_STRONG_INLINE scalar_quotient1_op(const scalar_quotient1_op& other) : m_other(other.m_other) { } EIGEN_STRONG_INLINE scalar_quotient1_op(const Scalar& other) : m_other(other) {} - EIGEN_STRONG_INLINE Scalar operator() (const Scalar& a) const { return a / m_other; } + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Scalar operator() (const Scalar& a) const { return a / m_other; } EIGEN_STRONG_INLINE const Packet packetOp(const Packet& a) const { return internal::pdiv(a, pset1(m_other)); } typename add_const_on_value_type::Nested>::type m_other; @@ -516,7 +519,7 @@ struct scalar_constant_op { EIGEN_STRONG_INLINE scalar_constant_op(const scalar_constant_op& other) : m_other(other.m_other) { } EIGEN_STRONG_INLINE scalar_constant_op(const Scalar& other) : m_other(other) { } template - EIGEN_STRONG_INLINE const Scalar operator() (Index, Index = 0) const { return m_other; } + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE const Scalar operator() (Index, Index = 0) const { return m_other; } template EIGEN_STRONG_INLINE const Packet packetOp(Index, Index = 0) const { return internal::pset1(m_other); } const Scalar m_other; @@ -529,7 +532,7 @@ struct functor_traits > template struct scalar_identity_op { EIGEN_EMPTY_STRUCT_CTOR(scalar_identity_op) template - EIGEN_STRONG_INLINE const Scalar operator() (Index row, Index col) const { return row==col ? Scalar(1) : Scalar(0); } + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE const Scalar operator() (Index row, Index col) const { return row==col ? Scalar(1) : Scalar(0); } }; template struct functor_traits > @@ -553,7 +556,7 @@ struct linspaced_op_impl m_base(padd(pset1(low),pmul(pset1(step),plset(-packet_traits::size)))) {} template - EIGEN_STRONG_INLINE const Scalar operator() (Index i) const { return m_low+i*m_step; } + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE const Scalar operator() (Index i) const { return m_low+i*m_step; } template EIGEN_STRONG_INLINE const Packet packetOp(Index) const { return m_base = padd(m_base,m_packetStep); } @@ -576,7 +579,7 @@ struct linspaced_op_impl m_lowPacket(pset1(m_low)), m_stepPacket(pset1(m_step)), m_interPacket(plset(0)) {} template - EIGEN_STRONG_INLINE const Scalar operator() (Index i) const { return m_low+i*m_step; } + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE const Scalar operator() (Index i) const { return m_low+i*m_step; } template EIGEN_STRONG_INLINE const Packet packetOp(Index i) const @@ -603,12 +606,12 @@ template struct linspaced_op linspaced_op(Scalar low, Scalar high, int num_steps) : impl((num_steps==1 ? high : low), (num_steps==1 ? Scalar() : (high-low)/(num_steps-1))) {} template - EIGEN_STRONG_INLINE const Scalar operator() (Index i) const { return impl(i); } + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE const Scalar operator() (Index i) const { return impl(i); } // We need this function when assigning e.g. a RowVectorXd to a MatrixXd since // there row==0 and col is used for the actual iteration. template - EIGEN_STRONG_INLINE const Scalar operator() (Index row, Index col) const + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE const Scalar operator() (Index row, Index col) const { eigen_assert(col==0 || row==0); return impl(col + row); @@ -657,9 +660,9 @@ template struct scalar_add_op { typedef typename packet_traits::type Packet; // FIXME default copy constructors seems bugged with std::complex<> - inline scalar_add_op(const scalar_add_op& other) : m_other(other.m_other) { } - inline scalar_add_op(const Scalar& other) : m_other(other) { } - inline Scalar operator() (const Scalar& a) const { return a + m_other; } + EIGEN_DEVICE_FUNC inline scalar_add_op(const scalar_add_op& other) : m_other(other.m_other) { } + EIGEN_DEVICE_FUNC inline scalar_add_op(const Scalar& other) : m_other(other) { } + EIGEN_DEVICE_FUNC inline Scalar operator() (const Scalar& a) const { return a + m_other; } inline const Packet packetOp(const Packet& a) const { return internal::padd(a, pset1(m_other)); } const Scalar m_other; @@ -674,7 +677,7 @@ struct functor_traits > */ template struct scalar_sqrt_op { EIGEN_EMPTY_STRUCT_CTOR(scalar_sqrt_op) - inline const Scalar operator() (const Scalar& a) const { using std::sqrt; return sqrt(a); } + EIGEN_DEVICE_FUNC inline const Scalar operator() (const Scalar& a) const { using std::sqrt; return sqrt(a); } typedef typename packet_traits::type Packet; inline Packet packetOp(const Packet& a) const { return internal::psqrt(a); } }; @@ -692,7 +695,7 @@ struct functor_traits > */ template struct scalar_cos_op { EIGEN_EMPTY_STRUCT_CTOR(scalar_cos_op) - inline Scalar operator() (const Scalar& a) const { using std::cos; return cos(a); } + EIGEN_DEVICE_FUNC inline Scalar operator() (const Scalar& a) const { using std::cos; return cos(a); } typedef typename packet_traits::type Packet; inline Packet packetOp(const Packet& a) const { return internal::pcos(a); } }; @@ -711,7 +714,7 @@ struct functor_traits > */ template struct scalar_sin_op { EIGEN_EMPTY_STRUCT_CTOR(scalar_sin_op) - inline const Scalar operator() (const Scalar& a) const { using std::sin; return sin(a); } + EIGEN_DEVICE_FUNC inline const Scalar operator() (const Scalar& a) const { using std::sin; return sin(a); } typedef typename packet_traits::type Packet; inline Packet packetOp(const Packet& a) const { return internal::psin(a); } }; @@ -731,7 +734,7 @@ struct functor_traits > */ template struct scalar_tan_op { EIGEN_EMPTY_STRUCT_CTOR(scalar_tan_op) - inline const Scalar operator() (const Scalar& a) const { using std::tan; return tan(a); } + EIGEN_DEVICE_FUNC inline const Scalar operator() (const Scalar& a) const { using std::tan; return tan(a); } typedef typename packet_traits::type Packet; inline Packet packetOp(const Packet& a) const { return internal::ptan(a); } }; @@ -750,7 +753,7 @@ struct functor_traits > */ template struct scalar_acos_op { EIGEN_EMPTY_STRUCT_CTOR(scalar_acos_op) - inline const Scalar operator() (const Scalar& a) const { using std::acos; return acos(a); } + EIGEN_DEVICE_FUNC inline const Scalar operator() (const Scalar& a) const { using std::acos; return acos(a); } typedef typename packet_traits::type Packet; inline Packet packetOp(const Packet& a) const { return internal::pacos(a); } }; @@ -769,7 +772,7 @@ struct functor_traits > */ template struct scalar_asin_op { EIGEN_EMPTY_STRUCT_CTOR(scalar_asin_op) - inline const Scalar operator() (const Scalar& a) const { using std::asin; return asin(a); } + EIGEN_DEVICE_FUNC inline const Scalar operator() (const Scalar& a) const { using std::asin; return asin(a); } typedef typename packet_traits::type Packet; inline Packet packetOp(const Packet& a) const { return internal::pasin(a); } }; @@ -791,7 +794,7 @@ struct scalar_pow_op { // FIXME default copy constructors seems bugged with std::complex<> inline scalar_pow_op(const scalar_pow_op& other) : m_exponent(other.m_exponent) { } inline scalar_pow_op(const Scalar& exponent) : m_exponent(exponent) {} - inline Scalar operator() (const Scalar& a) const { return internal::pow(a, m_exponent); } + EIGEN_DEVICE_FUNC inline Scalar operator() (const Scalar& a) const { return internal::pow(a, m_exponent); } const Scalar m_exponent; }; template @@ -805,7 +808,7 @@ struct functor_traits > template struct scalar_inverse_mult_op { scalar_inverse_mult_op(const Scalar& other) : m_other(other) {} - inline Scalar operator() (const Scalar& a) const { return m_other / a; } + EIGEN_DEVICE_FUNC inline Scalar operator() (const Scalar& a) const { return m_other / a; } template inline const Packet packetOp(const Packet& a) const { return internal::pdiv(pset1(m_other),a); } @@ -819,7 +822,7 @@ struct scalar_inverse_mult_op { template struct scalar_inverse_op { EIGEN_EMPTY_STRUCT_CTOR(scalar_inverse_op) - inline Scalar operator() (const Scalar& a) const { return Scalar(1)/a; } + EIGEN_DEVICE_FUNC inline Scalar operator() (const Scalar& a) const { return Scalar(1)/a; } template inline const Packet packetOp(const Packet& a) const { return internal::pdiv(pset1(Scalar(1)),a); } @@ -835,7 +838,7 @@ struct functor_traits > template struct scalar_square_op { EIGEN_EMPTY_STRUCT_CTOR(scalar_square_op) - inline Scalar operator() (const Scalar& a) const { return a*a; } + EIGEN_DEVICE_FUNC inline Scalar operator() (const Scalar& a) const { return a*a; } template inline const Packet packetOp(const Packet& a) const { return internal::pmul(a,a); } @@ -851,7 +854,7 @@ struct functor_traits > template struct scalar_cube_op { EIGEN_EMPTY_STRUCT_CTOR(scalar_cube_op) - inline Scalar operator() (const Scalar& a) const { return a*a*a; } + EIGEN_DEVICE_FUNC inline Scalar operator() (const Scalar& a) const { return a*a*a; } template inline const Packet packetOp(const Packet& a) const { return internal::pmul(a,pmul(a,a)); } diff --git a/Eigen/src/Core/GeneralProduct.h b/Eigen/src/Core/GeneralProduct.h index 9abc7b286..a070e618d 100644 --- a/Eigen/src/Core/GeneralProduct.h +++ b/Eigen/src/Core/GeneralProduct.h @@ -543,6 +543,7 @@ template<> struct gemv_selector * * \sa lazyProduct(), operator*=(const MatrixBase&), Cwise::operator*() */ +#ifndef __CUDACC__ template template inline const typename ProductReturnType::Type @@ -572,7 +573,7 @@ MatrixBase::operator*(const MatrixBase &other) const #endif return typename ProductReturnType::Type(derived(), other.derived()); } - +#endif /** \returns an expression of the matrix product of \c *this and \a other without implicit evaluation. * * The returned product will behave like any other expressions: the coefficients of the product will be diff --git a/Eigen/src/Core/Map.h b/Eigen/src/Core/Map.h index f804c89d6..e054f2202 100644 --- a/Eigen/src/Core/Map.h +++ b/Eigen/src/Core/Map.h @@ -118,11 +118,13 @@ template class Ma inline PointerType cast_to_pointer_type(PointerArgType ptr) { return ptr; } #endif + EIGEN_DEVICE_FUNC inline Index innerStride() const { return StrideType::InnerStrideAtCompileTime != 0 ? m_stride.inner() : 1; } + EIGEN_DEVICE_FUNC inline Index outerStride() const { return StrideType::OuterStrideAtCompileTime != 0 ? m_stride.outer() @@ -136,6 +138,7 @@ template class Ma * \param dataPtr pointer to the array to map * \param a_stride optional Stride object, passing the strides. */ + EIGEN_DEVICE_FUNC inline Map(PointerArgType dataPtr, const StrideType& a_stride = StrideType()) : Base(cast_to_pointer_type(dataPtr)), m_stride(a_stride) { @@ -148,6 +151,7 @@ template class Ma * \param a_size the size of the vector expression * \param a_stride optional Stride object, passing the strides. */ + EIGEN_DEVICE_FUNC inline Map(PointerArgType dataPtr, Index a_size, const StrideType& a_stride = StrideType()) : Base(cast_to_pointer_type(dataPtr), a_size), m_stride(a_stride) { @@ -161,6 +165,7 @@ template class Ma * \param nbCols the number of columns of the matrix expression * \param a_stride optional Stride object, passing the strides. */ + EIGEN_DEVICE_FUNC inline Map(PointerArgType dataPtr, Index nbRows, Index nbCols, const StrideType& a_stride = StrideType()) : Base(cast_to_pointer_type(dataPtr), nbRows, nbCols), m_stride(a_stride) { diff --git a/Eigen/src/Core/MapBase.h b/Eigen/src/Core/MapBase.h index 6876de588..8def7442d 100644 --- a/Eigen/src/Core/MapBase.h +++ b/Eigen/src/Core/MapBase.h @@ -76,8 +76,8 @@ template class MapBase typedef typename Base::CoeffReturnType CoeffReturnType; - inline Index rows() const { return m_rows.value(); } - inline Index cols() const { return m_cols.value(); } + EIGEN_DEVICE_FUNC inline Index rows() const { return m_rows.value(); } + EIGEN_DEVICE_FUNC inline Index cols() const { return m_cols.value(); } /** Returns a pointer to the first coefficient of the matrix or vector. * @@ -87,22 +87,26 @@ template class MapBase */ inline const Scalar* data() const { return m_data; } + EIGEN_DEVICE_FUNC inline const Scalar& coeff(Index rowId, Index colId) const { return m_data[colId * colStride() + rowId * rowStride()]; } + EIGEN_DEVICE_FUNC inline const Scalar& coeff(Index index) const { EIGEN_STATIC_ASSERT_INDEX_BASED_ACCESS(Derived) return m_data[index * innerStride()]; } + EIGEN_DEVICE_FUNC inline const Scalar& coeffRef(Index rowId, Index colId) const { return this->m_data[colId * colStride() + rowId * rowStride()]; } + EIGEN_DEVICE_FUNC inline const Scalar& coeffRef(Index index) const { EIGEN_STATIC_ASSERT_INDEX_BASED_ACCESS(Derived) @@ -123,12 +127,14 @@ template class MapBase return internal::ploadt(m_data + index * innerStride()); } + EIGEN_DEVICE_FUNC inline MapBase(PointerType dataPtr) : m_data(dataPtr), m_rows(RowsAtCompileTime), m_cols(ColsAtCompileTime) { EIGEN_STATIC_ASSERT_FIXED_SIZE(Derived) checkSanity(); } + EIGEN_DEVICE_FUNC inline MapBase(PointerType dataPtr, Index vecSize) : m_data(dataPtr), m_rows(RowsAtCompileTime == Dynamic ? vecSize : Index(RowsAtCompileTime)), @@ -140,6 +146,7 @@ template class MapBase checkSanity(); } + EIGEN_DEVICE_FUNC inline MapBase(PointerType dataPtr, Index nbRows, Index nbCols) : m_data(dataPtr), m_rows(nbRows), m_cols(nbCols) { @@ -151,6 +158,7 @@ template class MapBase protected: + EIGEN_DEVICE_FUNC void checkSanity() const { EIGEN_STATIC_ASSERT(EIGEN_IMPLIES(internal::traits::Flags&PacketAccessBit, @@ -198,11 +206,13 @@ template class MapBase inline const Scalar* data() const { return this->m_data; } inline ScalarWithConstIfNotLvalue* data() { return this->m_data; } // no const-cast here so non-const-correct code will give a compile error + EIGEN_DEVICE_FUNC inline ScalarWithConstIfNotLvalue& coeffRef(Index row, Index col) { return this->m_data[col * colStride() + row * rowStride()]; } + EIGEN_DEVICE_FUNC inline ScalarWithConstIfNotLvalue& coeffRef(Index index) { EIGEN_STATIC_ASSERT_INDEX_BASED_ACCESS(Derived) @@ -224,10 +234,11 @@ template class MapBase (this->m_data + index * innerStride(), val); } - explicit inline MapBase(PointerType dataPtr) : Base(dataPtr) {} - inline MapBase(PointerType dataPtr, Index vecSize) : Base(dataPtr, vecSize) {} - inline MapBase(PointerType dataPtr, Index nbRows, Index nbCols) : Base(dataPtr, nbRows, nbCols) {} + EIGEN_DEVICE_FUNC explicit inline MapBase(PointerType dataPtr) : Base(dataPtr) {} + EIGEN_DEVICE_FUNC inline MapBase(PointerType dataPtr, Index vecSize) : Base(dataPtr, vecSize) {} + EIGEN_DEVICE_FUNC inline MapBase(PointerType dataPtr, Index nbRows, Index nbCols) : Base(dataPtr, nbRows, nbCols) {} + EIGEN_DEVICE_FUNC Derived& operator=(const MapBase& other) { Base::Base::operator=(other); diff --git a/Eigen/src/Core/Matrix.h b/Eigen/src/Core/Matrix.h index 99160b591..61af9d9a3 100644 --- a/Eigen/src/Core/Matrix.h +++ b/Eigen/src/Core/Matrix.h @@ -151,6 +151,7 @@ class Matrix * * \callgraph */ + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Matrix& operator=(const Matrix& other) { return Base::_set(other); @@ -167,6 +168,7 @@ class Matrix * remain row-vectors and vectors remain vectors. */ template + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Matrix& operator=(const MatrixBase& other) { return Base::_set(other); @@ -179,12 +181,14 @@ class Matrix * \copydetails DenseBase::operator=(const EigenBase &other) */ template + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Matrix& operator=(const EigenBase &other) { return Base::operator=(other); } template + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Matrix& operator=(const ReturnByValue& func) { return Base::operator=(func); @@ -200,6 +204,7 @@ class Matrix * * \sa resize(Index,Index) */ + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE explicit Matrix() : Base() { Base::_check_template_params(); @@ -207,6 +212,7 @@ class Matrix } // FIXME is it still needed + EIGEN_DEVICE_FUNC Matrix(internal::constructor_without_unaligned_array_assert) : Base(internal::constructor_without_unaligned_array_assert()) { Base::_check_template_params(); EIGEN_INITIALIZE_BY_ZERO_IF_THAT_OPTION_IS_ENABLED } @@ -217,6 +223,7 @@ class Matrix * it is redundant to pass the dimension here, so it makes more sense to use the default * constructor Matrix() instead. */ + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE explicit Matrix(Index dim) : Base(dim, RowsAtCompileTime == 1 ? 1 : dim, ColsAtCompileTime == 1 ? 1 : dim) { @@ -229,6 +236,7 @@ class Matrix #ifndef EIGEN_PARSED_BY_DOXYGEN template + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Matrix(const T0& x, const T1& y) { Base::_check_template_params(); @@ -240,12 +248,14 @@ class Matrix * This is useful for dynamic-size matrices. For fixed-size matrices, * it is redundant to pass these parameters, so one should use the default constructor * Matrix() instead. */ + EIGEN_DEVICE_FUNC Matrix(Index rows, Index cols); /** \brief Constructs an initialized 2D vector with given coefficients */ Matrix(const Scalar& x, const Scalar& y); #endif /** \brief Constructs an initialized 3D vector with given coefficients */ + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Matrix(const Scalar& x, const Scalar& y, const Scalar& z) { Base::_check_template_params(); @@ -255,6 +265,7 @@ class Matrix m_storage.data()[2] = z; } /** \brief Constructs an initialized 4D vector with given coefficients */ + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Matrix(const Scalar& x, const Scalar& y, const Scalar& z, const Scalar& w) { Base::_check_template_params(); @@ -265,10 +276,12 @@ class Matrix m_storage.data()[3] = w; } + EIGEN_DEVICE_FUNC explicit Matrix(const Scalar *data); /** \brief Constructor copying the value of the expression \a other */ template + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Matrix(const MatrixBase& other) : Base(other.rows() * other.cols(), other.rows(), other.cols()) { @@ -281,6 +294,7 @@ class Matrix Base::_set_noalias(other); } /** \brief Copy constructor */ + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Matrix(const Matrix& other) : Base(other.rows() * other.cols(), other.rows(), other.cols()) { @@ -289,6 +303,7 @@ class Matrix } /** \brief Copy constructor with in-place evaluation */ template + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Matrix(const ReturnByValue& other) { Base::_check_template_params(); @@ -300,6 +315,7 @@ class Matrix * \sa MatrixBase::operator=(const EigenBase&) */ template + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Matrix(const EigenBase &other) : Base(other.derived().rows() * other.derived().cols(), other.derived().rows(), other.derived().cols()) { @@ -318,14 +334,16 @@ class Matrix void swap(MatrixBase const & other) { this->_swap(other.derived()); } - inline Index innerStride() const { return 1; } - inline Index outerStride() const { return this->innerSize(); } + EIGEN_DEVICE_FUNC inline Index innerStride() const { return 1; } + EIGEN_DEVICE_FUNC inline Index outerStride() const { return this->innerSize(); } /////////// Geometry module /////////// template + EIGEN_DEVICE_FUNC explicit Matrix(const RotationBase& r); template + EIGEN_DEVICE_FUNC Matrix& operator=(const RotationBase& r); #ifdef EIGEN2_SUPPORT diff --git a/Eigen/src/Core/MatrixBase.h b/Eigen/src/Core/MatrixBase.h index e12a6763e..c76192e46 100644 --- a/Eigen/src/Core/MatrixBase.h +++ b/Eigen/src/Core/MatrixBase.h @@ -145,22 +145,27 @@ template class MatrixBase /** Special case of the template operator=, in order to prevent the compiler * from generating a default operator= (issue hit with g++ 4.1) */ + EIGEN_DEVICE_FUNC Derived& operator=(const MatrixBase& other); // We cannot inherit here via Base::operator= since it is causing // trouble with MSVC. template + EIGEN_DEVICE_FUNC Derived& operator=(const DenseBase& other); template + EIGEN_DEVICE_FUNC Derived& operator=(const EigenBase& other); template + EIGEN_DEVICE_FUNC Derived& operator=(const ReturnByValue& other); #ifndef EIGEN_PARSED_BY_DOXYGEN template + EIGEN_DEVICE_FUNC Derived& lazyAssign(const ProductBase& other); template @@ -168,15 +173,26 @@ template class MatrixBase #endif // not EIGEN_PARSED_BY_DOXYGEN template + EIGEN_DEVICE_FUNC Derived& operator+=(const MatrixBase& other); template + EIGEN_DEVICE_FUNC Derived& operator-=(const MatrixBase& other); +#ifdef __CUDACC__ + template + EIGEN_DEVICE_FUNC + const typename LazyProductReturnType::Type + operator*(const MatrixBase &other) const + { return this->lazyProduct(other); } +#else template const typename ProductReturnType::Type operator*(const MatrixBase &other) const; +#endif template + EIGEN_DEVICE_FUNC const typename LazyProductReturnType::Type lazyProduct(const MatrixBase &other) const; @@ -194,6 +210,7 @@ template class MatrixBase operator*(const DiagonalBase &diagonal) const; template + EIGEN_DEVICE_FUNC typename internal::scalar_product_traits::Scalar,typename internal::traits::Scalar>::ReturnType dot(const MatrixBase& other) const; @@ -324,8 +341,8 @@ template class MatrixBase /////////// LU module /////////// - const FullPivLU fullPivLu() const; - const PartialPivLU partialPivLu() const; + EIGEN_DEVICE_FUNC const FullPivLU fullPivLu() const; + EIGEN_DEVICE_FUNC const PartialPivLU partialPivLu() const; #if EIGEN2_SUPPORT_STAGE < STAGE20_RESOLVE_API_CONFLICTS const LU lu() const; @@ -346,6 +363,7 @@ template class MatrixBase } #endif + EIGEN_DEVICE_FUNC const internal::inverse_impl inverse() const; template void computeInverseAndDetWithCheck( @@ -495,12 +513,12 @@ template class MatrixBase #endif protected: - MatrixBase() : Base() {} + EIGEN_DEVICE_FUNC MatrixBase() : Base() {} private: - explicit MatrixBase(int); - MatrixBase(int,int); - template explicit MatrixBase(const MatrixBase&); + EIGEN_DEVICE_FUNC explicit MatrixBase(int); + EIGEN_DEVICE_FUNC MatrixBase(int,int); + template EIGEN_DEVICE_FUNC explicit MatrixBase(const MatrixBase&); protected: // mixing arrays and matrices is not legal template Derived& operator+=(const ArrayBase& ) diff --git a/Eigen/src/Core/NoAlias.h b/Eigen/src/Core/NoAlias.h index 0112c865b..9e371538a 100644 --- a/Eigen/src/Core/NoAlias.h +++ b/Eigen/src/Core/NoAlias.h @@ -37,11 +37,13 @@ class NoAlias /** Behaves like MatrixBase::lazyAssign(other) * \sa MatrixBase::lazyAssign() */ template + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE ExpressionType& operator=(const StorageBase& other) { return internal::assign_selector::run(m_expression,other.derived()); } /** \sa MatrixBase::operator+= */ template + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE ExpressionType& operator+=(const StorageBase& other) { typedef SelfCwiseBinaryOp, ExpressionType, OtherDerived> SelfAdder; @@ -54,6 +56,7 @@ class NoAlias /** \sa MatrixBase::operator-= */ template + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE ExpressionType& operator-=(const StorageBase& other) { typedef SelfCwiseBinaryOp, ExpressionType, OtherDerived> SelfAdder; @@ -66,10 +69,12 @@ class NoAlias #ifndef EIGEN_PARSED_BY_DOXYGEN template + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE ExpressionType& operator+=(const ProductBase& other) { other.derived().addTo(m_expression); return m_expression; } template + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE ExpressionType& operator-=(const ProductBase& other) { other.derived().subTo(m_expression); return m_expression; } @@ -78,10 +83,12 @@ class NoAlias { return m_expression.derived() += CoeffBasedProduct(other.lhs(), other.rhs()); } template + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE ExpressionType& operator-=(const CoeffBasedProduct& other) { return m_expression.derived() -= CoeffBasedProduct(other.lhs(), other.rhs()); } #endif + EIGEN_DEVICE_FUNC ExpressionType& expression() const { return m_expression; diff --git a/Eigen/src/Core/PlainObjectBase.h b/Eigen/src/Core/PlainObjectBase.h index 5c94ef621..49a5518e3 100644 --- a/Eigen/src/Core/PlainObjectBase.h +++ b/Eigen/src/Core/PlainObjectBase.h @@ -23,6 +23,7 @@ namespace internal { template struct check_rows_cols_for_overflow { template + EIGEN_DEVICE_FUNC static EIGEN_ALWAYS_INLINE void run(Index, Index) { } @@ -30,6 +31,7 @@ template struct check_rows_cols_for_overflow { template<> struct check_rows_cols_for_overflow { template + EIGEN_DEVICE_FUNC static EIGEN_ALWAYS_INLINE void run(Index rows, Index cols) { // http://hg.mozilla.org/mozilla-central/file/6c8a909977d3/xpcom/ds/CheckedInt.h#l242 @@ -124,9 +126,12 @@ class PlainObjectBase : public internal::dense_xpr_base::type Base& base() { return *static_cast(this); } const Base& base() const { return *static_cast(this); } + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Index rows() const { return m_storage.rows(); } + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Index cols() const { return m_storage.cols(); } + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE const Scalar& coeff(Index rowId, Index colId) const { if(Flags & RowMajorBit) @@ -135,11 +140,13 @@ class PlainObjectBase : public internal::dense_xpr_base::type return m_storage.data()[rowId + colId * m_storage.rows()]; } + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE const Scalar& coeff(Index index) const { return m_storage.data()[index]; } + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Scalar& coeffRef(Index rowId, Index colId) { if(Flags & RowMajorBit) @@ -148,11 +155,13 @@ class PlainObjectBase : public internal::dense_xpr_base::type return m_storage.data()[rowId + colId * m_storage.rows()]; } + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Scalar& coeffRef(Index index) { return m_storage.data()[index]; } + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE const Scalar& coeffRef(Index rowId, Index colId) const { if(Flags & RowMajorBit) @@ -161,6 +170,7 @@ class PlainObjectBase : public internal::dense_xpr_base::type return m_storage.data()[rowId + colId * m_storage.rows()]; } + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE const Scalar& coeffRef(Index index) const { return m_storage.data()[index]; @@ -224,6 +234,7 @@ class PlainObjectBase : public internal::dense_xpr_base::type * * \sa resize(Index) for vectors, resize(NoChange_t, Index), resize(Index, NoChange_t) */ + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void resize(Index nbRows, Index nbCols) { eigen_assert( EIGEN_IMPLIES(RowsAtCompileTime!=Dynamic,nbRows==RowsAtCompileTime) @@ -254,6 +265,7 @@ class PlainObjectBase : public internal::dense_xpr_base::type * * \sa resize(Index,Index), resize(NoChange_t, Index), resize(Index, NoChange_t) */ + EIGEN_DEVICE_FUNC inline void resize(Index size) { EIGEN_STATIC_ASSERT_VECTOR_ONLY(PlainObjectBase) @@ -278,6 +290,7 @@ class PlainObjectBase : public internal::dense_xpr_base::type * * \sa resize(Index,Index) */ + EIGEN_DEVICE_FUNC inline void resize(NoChange_t, Index nbCols) { resize(rows(), nbCols); @@ -291,6 +304,7 @@ class PlainObjectBase : public internal::dense_xpr_base::type * * \sa resize(Index,Index) */ + EIGEN_DEVICE_FUNC inline void resize(Index nbRows, NoChange_t) { resize(nbRows, cols()); @@ -304,6 +318,7 @@ class PlainObjectBase : public internal::dense_xpr_base::type * remain row-vectors and vectors remain vectors. */ template + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void resizeLike(const EigenBase& _other) { const OtherDerived& other = _other.derived(); @@ -393,6 +408,7 @@ class PlainObjectBase : public internal::dense_xpr_base::type /** This is a special case of the templated operator=. Its purpose is to * prevent a default operator= from hiding the templated operator=. */ + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Derived& operator=(const PlainObjectBase& other) { return _set(other); @@ -400,6 +416,7 @@ class PlainObjectBase : public internal::dense_xpr_base::type /** \sa MatrixBase::lazyAssign() */ template + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Derived& lazyAssign(const DenseBase& other) { _resize_to_match(other); @@ -407,12 +424,14 @@ class PlainObjectBase : public internal::dense_xpr_base::type } template + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Derived& operator=(const ReturnByValue& func) { resize(func.rows(), func.cols()); return Base::operator=(func); } + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE explicit PlainObjectBase() : m_storage() { // _check_template_params(); @@ -422,6 +441,7 @@ class PlainObjectBase : public internal::dense_xpr_base::type #ifndef EIGEN_PARSED_BY_DOXYGEN // FIXME is it still needed ? /** \internal */ + EIGEN_DEVICE_FUNC PlainObjectBase(internal::constructor_without_unaligned_array_assert) : m_storage(internal::constructor_without_unaligned_array_assert()) { @@ -429,6 +449,7 @@ class PlainObjectBase : public internal::dense_xpr_base::type } #endif + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE PlainObjectBase(Index a_size, Index nbRows, Index nbCols) : m_storage(a_size, nbRows, nbCols) { @@ -439,6 +460,7 @@ class PlainObjectBase : public internal::dense_xpr_base::type /** \copydoc MatrixBase::operator=(const EigenBase&) */ template + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Derived& operator=(const EigenBase &other) { _resize_to_match(other); @@ -448,6 +470,7 @@ class PlainObjectBase : public internal::dense_xpr_base::type /** \sa MatrixBase::operator=(const EigenBase&) */ template + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE PlainObjectBase(const EigenBase &other) : m_storage(other.derived().rows() * other.derived().cols(), other.derived().rows(), other.derived().cols()) { @@ -558,6 +581,7 @@ class PlainObjectBase : public internal::dense_xpr_base::type * remain row-vectors and vectors remain vectors. */ template + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void _resize_to_match(const EigenBase& other) { #ifdef EIGEN_NO_AUTOMATIC_RESIZING @@ -585,6 +609,7 @@ class PlainObjectBase : public internal::dense_xpr_base::type * \internal */ template + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Derived& _set(const DenseBase& other) { _set_selector(other.derived(), typename internal::conditional(int(OtherDerived::Flags) & EvalBeforeAssigningBit), internal::true_type, internal::false_type>::type()); @@ -592,9 +617,11 @@ class PlainObjectBase : public internal::dense_xpr_base::type } template + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void _set_selector(const OtherDerived& other, const internal::true_type&) { _set_noalias(other.eval()); } template + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void _set_selector(const OtherDerived& other, const internal::false_type&) { _set_noalias(other); } /** \internal Like _set() but additionally makes the assumption that no aliasing effect can happen (which @@ -603,6 +630,7 @@ class PlainObjectBase : public internal::dense_xpr_base::type * \sa operator=(const MatrixBase&), _set() */ template + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Derived& _set_noalias(const DenseBase& other) { // I don't think we need this resize call since the lazyAssign will anyways resize @@ -622,6 +650,7 @@ class PlainObjectBase : public internal::dense_xpr_base::type resize(nbRows,nbCols); } template + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void _init2(const Scalar& val0, const Scalar& val1, typename internal::enable_if::type* = 0) { EIGEN_STATIC_ASSERT_VECTOR_SPECIFIC_SIZE(PlainObjectBase, 2) @@ -644,6 +673,7 @@ class PlainObjectBase : public internal::dense_xpr_base::type public: #ifndef EIGEN_PARSED_BY_DOXYGEN + EIGEN_DEVICE_FUNC static EIGEN_STRONG_INLINE void _check_template_params() { EIGEN_STATIC_ASSERT((EIGEN_IMPLIES(MaxRowsAtCompileTime==1 && MaxColsAtCompileTime!=1, (Options&RowMajor)==RowMajor) diff --git a/Eigen/src/Core/Redux.h b/Eigen/src/Core/Redux.h index b7ce7c658..12b3db584 100644 --- a/Eigen/src/Core/Redux.h +++ b/Eigen/src/Core/Redux.h @@ -82,6 +82,7 @@ struct redux_novec_unroller typedef typename Derived::Scalar Scalar; + EIGEN_DEVICE_FUNC static EIGEN_STRONG_INLINE Scalar run(const Derived &mat, const Func& func) { return func(redux_novec_unroller::run(mat,func), @@ -99,6 +100,7 @@ struct redux_novec_unroller typedef typename Derived::Scalar Scalar; + EIGEN_DEVICE_FUNC static EIGEN_STRONG_INLINE Scalar run(const Derived &mat, const Func&) { return mat.coeffByOuterInner(outer, inner); @@ -112,6 +114,7 @@ template struct redux_novec_unroller { typedef typename Derived::Scalar Scalar; + EIGEN_DEVICE_FUNC static EIGEN_STRONG_INLINE Scalar run(const Derived&, const Func&) { return Scalar(); } }; @@ -170,6 +173,7 @@ struct redux_impl { typedef typename Derived::Scalar Scalar; typedef typename Derived::Index Index; + EIGEN_DEVICE_FUNC static EIGEN_STRONG_INLINE Scalar run(const Derived& mat, const Func& func) { eigen_assert(mat.rows()>0 && mat.cols()>0 && "you are using an empty matrix"); diff --git a/Eigen/src/Core/ReturnByValue.h b/Eigen/src/Core/ReturnByValue.h index 613912ffa..87c6d9194 100644 --- a/Eigen/src/Core/ReturnByValue.h +++ b/Eigen/src/Core/ReturnByValue.h @@ -57,10 +57,11 @@ template class ReturnByValue EIGEN_DENSE_PUBLIC_INTERFACE(ReturnByValue) template + EIGEN_DEVICE_FUNC inline void evalTo(Dest& dst) const { static_cast(this)->evalTo(dst); } - inline Index rows() const { return static_cast(this)->rows(); } - inline Index cols() const { return static_cast(this)->cols(); } + EIGEN_DEVICE_FUNC inline Index rows() const { return static_cast(this)->rows(); } + EIGEN_DEVICE_FUNC inline Index cols() const { return static_cast(this)->cols(); } #ifndef EIGEN_PARSED_BY_DOXYGEN #define Unusable YOU_ARE_TRYING_TO_ACCESS_A_SINGLE_COEFFICIENT_IN_A_SPECIAL_EXPRESSION_WHERE_THAT_IS_NOT_ALLOWED_BECAUSE_THAT_WOULD_BE_INEFFICIENT diff --git a/Eigen/src/Core/Stride.h b/Eigen/src/Core/Stride.h index 1e3f5fe9f..d3d454e4e 100644 --- a/Eigen/src/Core/Stride.h +++ b/Eigen/src/Core/Stride.h @@ -51,6 +51,7 @@ class Stride }; /** Default constructor, for use when strides are fixed at compile time */ + EIGEN_DEVICE_FUNC Stride() : m_outer(OuterStrideAtCompileTime), m_inner(InnerStrideAtCompileTime) { @@ -58,6 +59,7 @@ class Stride } /** Constructor allowing to pass the strides at runtime */ + EIGEN_DEVICE_FUNC Stride(Index outerStride, Index innerStride) : m_outer(outerStride), m_inner(innerStride) { @@ -65,13 +67,16 @@ class Stride } /** Copy constructor */ + EIGEN_DEVICE_FUNC Stride(const Stride& other) : m_outer(other.outer()), m_inner(other.inner()) {} /** \returns the outer stride */ + EIGEN_DEVICE_FUNC inline Index outer() const { return m_outer.value(); } /** \returns the inner stride */ + EIGEN_DEVICE_FUNC inline Index inner() const { return m_inner.value(); } protected: @@ -87,8 +92,8 @@ class InnerStride : public Stride<0, Value> typedef Stride<0, Value> Base; public: typedef DenseIndex Index; - InnerStride() : Base() {} - InnerStride(Index v) : Base(0, v) {} + EIGEN_DEVICE_FUNC InnerStride() : Base() {} + EIGEN_DEVICE_FUNC InnerStride(Index v) : Base(0, v) {} }; /** \brief Convenience specialization of Stride to specify only an outer stride @@ -99,8 +104,8 @@ class OuterStride : public Stride typedef Stride Base; public: typedef DenseIndex Index; - OuterStride() : Base() {} - OuterStride(Index v) : Base(v,0) {} + EIGEN_DEVICE_FUNC OuterStride() : Base() {} + EIGEN_DEVICE_FUNC OuterStride(Index v) : Base(v,0) {} }; } // end namespace Eigen diff --git a/Eigen/src/Core/Transpose.h b/Eigen/src/Core/Transpose.h index 34944e055..b5e1468df 100644 --- a/Eigen/src/Core/Transpose.h +++ b/Eigen/src/Core/Transpose.h @@ -62,18 +62,21 @@ template class Transpose typedef typename TransposeImpl::StorageKind>::Base Base; EIGEN_GENERIC_PUBLIC_INTERFACE(Transpose) + EIGEN_DEVICE_FUNC inline Transpose(MatrixType& a_matrix) : m_matrix(a_matrix) {} EIGEN_INHERIT_ASSIGNMENT_OPERATORS(Transpose) - inline Index rows() const { return m_matrix.cols(); } - inline Index cols() const { return m_matrix.rows(); } + EIGEN_DEVICE_FUNC inline Index rows() const { return m_matrix.cols(); } + EIGEN_DEVICE_FUNC inline Index cols() const { return m_matrix.rows(); } /** \returns the nested expression */ + EIGEN_DEVICE_FUNC const typename internal::remove_all::type& nestedExpression() const { return m_matrix; } /** \returns the nested expression */ + EIGEN_DEVICE_FUNC typename internal::remove_all::type& nestedExpression() { return m_matrix.const_cast_derived(); } @@ -105,8 +108,8 @@ template class TransposeImpl typedef typename internal::TransposeImpl_base::type Base; EIGEN_DENSE_PUBLIC_INTERFACE(Transpose) - inline Index innerStride() const { return derived().nestedExpression().innerStride(); } - inline Index outerStride() const { return derived().nestedExpression().outerStride(); } + EIGEN_DEVICE_FUNC inline Index innerStride() const { return derived().nestedExpression().innerStride(); } + EIGEN_DEVICE_FUNC inline Index outerStride() const { return derived().nestedExpression().outerStride(); } typedef typename internal::conditional< internal::is_lvalue::value, @@ -117,33 +120,39 @@ template class TransposeImpl inline ScalarWithConstIfNotLvalue* data() { return derived().nestedExpression().data(); } inline const Scalar* data() const { return derived().nestedExpression().data(); } + EIGEN_DEVICE_FUNC inline ScalarWithConstIfNotLvalue& coeffRef(Index rowId, Index colId) { EIGEN_STATIC_ASSERT_LVALUE(MatrixType) return derived().nestedExpression().const_cast_derived().coeffRef(colId, rowId); } + EIGEN_DEVICE_FUNC inline ScalarWithConstIfNotLvalue& coeffRef(Index index) { EIGEN_STATIC_ASSERT_LVALUE(MatrixType) return derived().nestedExpression().const_cast_derived().coeffRef(index); } + EIGEN_DEVICE_FUNC inline const Scalar& coeffRef(Index rowId, Index colId) const { return derived().nestedExpression().coeffRef(colId, rowId); } + EIGEN_DEVICE_FUNC inline const Scalar& coeffRef(Index index) const { return derived().nestedExpression().coeffRef(index); } + EIGEN_DEVICE_FUNC inline CoeffReturnType coeff(Index rowId, Index colId) const { return derived().nestedExpression().coeff(colId, rowId); } + EIGEN_DEVICE_FUNC inline CoeffReturnType coeff(Index index) const { return derived().nestedExpression().coeff(index); diff --git a/Eigen/src/Core/VectorBlock.h b/Eigen/src/Core/VectorBlock.h index 1a7330f3c..216c568c4 100644 --- a/Eigen/src/Core/VectorBlock.h +++ b/Eigen/src/Core/VectorBlock.h @@ -72,6 +72,7 @@ template class VectorBlock /** Dynamic-size constructor */ + EIGEN_DEVICE_FUNC inline VectorBlock(VectorType& vector, Index start, Index size) : Base(vector, IsColVector ? start : 0, IsColVector ? 0 : start, @@ -82,6 +83,7 @@ template class VectorBlock /** Fixed-size constructor */ + EIGEN_DEVICE_FUNC inline VectorBlock(VectorType& vector, Index start) : Base(vector, IsColVector ? start : 0, IsColVector ? 0 : start) { diff --git a/Eigen/src/Core/products/CoeffBasedProduct.h b/Eigen/src/Core/products/CoeffBasedProduct.h index 403d25fa9..312a05c71 100644 --- a/Eigen/src/Core/products/CoeffBasedProduct.h +++ b/Eigen/src/Core/products/CoeffBasedProduct.h @@ -140,11 +140,13 @@ class CoeffBasedProduct public: + EIGEN_DEVICE_FUNC inline CoeffBasedProduct(const CoeffBasedProduct& other) : Base(), m_lhs(other.m_lhs), m_rhs(other.m_rhs) {} template + EIGEN_DEVICE_FUNC inline CoeffBasedProduct(const Lhs& lhs, const Rhs& rhs) : m_lhs(lhs), m_rhs(rhs) { @@ -157,9 +159,10 @@ class CoeffBasedProduct && "if you wanted a coeff-wise or a dot product use the respective explicit functions"); } - EIGEN_STRONG_INLINE Index rows() const { return m_lhs.rows(); } - EIGEN_STRONG_INLINE Index cols() const { return m_rhs.cols(); } + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Index rows() const { return m_lhs.rows(); } + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Index cols() const { return m_rhs.cols(); } + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE const Scalar coeff(Index row, Index col) const { Scalar res; @@ -170,6 +173,7 @@ class CoeffBasedProduct /* Allow index-based non-packet access. It is impossible though to allow index-based packed access, * which is why we don't set the LinearAccessBit. */ + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE const Scalar coeff(Index index) const { Scalar res; @@ -191,22 +195,26 @@ class CoeffBasedProduct } // Implicit conversion to the nested type (trigger the evaluation of the product) + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE operator const PlainObject& () const { m_result.lazyAssign(*this); return m_result; } - const _LhsNested& lhs() const { return m_lhs; } - const _RhsNested& rhs() const { return m_rhs; } + EIGEN_DEVICE_FUNC const _LhsNested& lhs() const { return m_lhs; } + EIGEN_DEVICE_FUNC const _RhsNested& rhs() const { return m_rhs; } + EIGEN_DEVICE_FUNC const Diagonal diagonal() const { return reinterpret_cast(*this); } template + EIGEN_DEVICE_FUNC const Diagonal diagonal() const { return reinterpret_cast(*this); } + EIGEN_DEVICE_FUNC const Diagonal diagonal(Index index) const { return reinterpret_cast(*this).diagonal(index); } @@ -239,6 +247,7 @@ template struct product_coeff_impl { typedef typename Lhs::Index Index; + EIGEN_DEVICE_FUNC static EIGEN_STRONG_INLINE void run(Index row, Index col, const Lhs& lhs, const Rhs& rhs, RetScalar &res) { product_coeff_impl::run(row, col, lhs, rhs, res); @@ -250,6 +259,7 @@ template struct product_coeff_impl { typedef typename Lhs::Index Index; + EIGEN_DEVICE_FUNC static EIGEN_STRONG_INLINE void run(Index row, Index col, const Lhs& lhs, const Rhs& rhs, RetScalar &res) { res = lhs.coeff(row, 0) * rhs.coeff(0, col); @@ -260,6 +270,7 @@ template struct product_coeff_impl { typedef typename Lhs::Index Index; + EIGEN_DEVICE_FUNC static EIGEN_STRONG_INLINE void run(Index row, Index col, const Lhs& lhs, const Rhs& rhs, RetScalar& res) { eigen_assert(lhs.cols()>0 && "you are using a non initialized matrix"); diff --git a/Eigen/src/Core/util/Macros.h b/Eigen/src/Core/util/Macros.h index 933a34c9d..bf6a9293c 100644 --- a/Eigen/src/Core/util/Macros.h +++ b/Eigen/src/Core/util/Macros.h @@ -396,7 +396,7 @@ #define EIGEN_MAKE_CWISE_BINARY_OP(METHOD,FUNCTOR) \ template \ - EIGEN_STRONG_INLINE const CwiseBinaryOp, const Derived, const OtherDerived> \ + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE const CwiseBinaryOp, const Derived, const OtherDerived> \ (METHOD)(const EIGEN_CURRENT_STORAGE_BASE_CLASS &other) const \ { \ return CwiseBinaryOp, const Derived, const OtherDerived>(derived(), other.derived()); \ diff --git a/Eigen/src/Core/util/XprHelper.h b/Eigen/src/Core/util/XprHelper.h index 3d1290cd2..f115d3779 100644 --- a/Eigen/src/Core/util/XprHelper.h +++ b/Eigen/src/Core/util/XprHelper.h @@ -16,8 +16,8 @@ // so currently we simply disable this optimization for gcc 4.3 #if (defined __GNUG__) && !((__GNUC__==4) && (__GNUC_MINOR__==3)) #define EIGEN_EMPTY_STRUCT_CTOR(X) \ - EIGEN_STRONG_INLINE X() {} \ - EIGEN_STRONG_INLINE X(const X& ) {} + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE X() {} \ + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE X(const X& ) {} #else #define EIGEN_EMPTY_STRUCT_CTOR(X) #endif @@ -50,19 +50,19 @@ template class variable_if_dynamic { public: EIGEN_EMPTY_STRUCT_CTOR(variable_if_dynamic) - explicit variable_if_dynamic(T v) { EIGEN_ONLY_USED_FOR_DEBUG(v); assert(v == T(Value)); } - static T value() { return T(Value); } - void setValue(T) {} + EIGEN_DEVICE_FUNC explicit variable_if_dynamic(T v) { EIGEN_ONLY_USED_FOR_DEBUG(v); eigen_assert(v == T(Value)); } + EIGEN_DEVICE_FUNC static T value() { return T(Value); } + EIGEN_DEVICE_FUNC void setValue(T) {} }; template class variable_if_dynamic { T m_value; - variable_if_dynamic() { assert(false); } + EIGEN_DEVICE_FUNC variable_if_dynamic() { eigen_assert(false); } public: - explicit variable_if_dynamic(T value) : m_value(value) {} - T value() const { return m_value; } - void setValue(T value) { m_value = value; } + EIGEN_DEVICE_FUNC explicit variable_if_dynamic(T value) : m_value(value) {} + EIGEN_DEVICE_FUNC T value() const { return m_value; } + EIGEN_DEVICE_FUNC void setValue(T value) { m_value = value; } }; /** \internal like variable_if_dynamic but for DynamicIndex @@ -71,19 +71,19 @@ template class variable_if_dynamicindex { public: EIGEN_EMPTY_STRUCT_CTOR(variable_if_dynamicindex) - explicit variable_if_dynamicindex(T v) { EIGEN_ONLY_USED_FOR_DEBUG(v); assert(v == T(Value)); } - static T value() { return T(Value); } - void setValue(T) {} + EIGEN_DEVICE_FUNC explicit variable_if_dynamicindex(T v) { EIGEN_ONLY_USED_FOR_DEBUG(v); eigen_assert(v == T(Value)); } + EIGEN_DEVICE_FUNC static T value() { return T(Value); } + EIGEN_DEVICE_FUNC void setValue(T) {} }; template class variable_if_dynamicindex { T m_value; - variable_if_dynamicindex() { assert(false); } + EIGEN_DEVICE_FUNC variable_if_dynamicindex() { eigen_assert(false); } public: - explicit variable_if_dynamicindex(T value) : m_value(value) {} - T value() const { return m_value; } - void setValue(T value) { m_value = value; } + EIGEN_DEVICE_FUNC explicit variable_if_dynamicindex(T value) : m_value(value) {} + EIGEN_DEVICE_FUNC T value() const { return m_value; } + EIGEN_DEVICE_FUNC void setValue(T value) { m_value = value; } }; template struct functor_traits @@ -340,6 +340,7 @@ template::type> str }; template +EIGEN_DEVICE_FUNC T* const_cast_ptr(const T* ptr) { return const_cast(ptr); diff --git a/Eigen/src/LU/FullPivLU.h b/Eigen/src/LU/FullPivLU.h index 14a9c402d..bcd30be00 100644 --- a/Eigen/src/LU/FullPivLU.h +++ b/Eigen/src/LU/FullPivLU.h @@ -727,12 +727,14 @@ struct solve_retval, Rhs> * * \sa class FullPivLU */ +#ifndef __CUDACC__ template inline const FullPivLU::PlainObject> MatrixBase::fullPivLu() const { return FullPivLU(eval()); } +#endif } // end namespace Eigen diff --git a/Eigen/src/LU/Inverse.h b/Eigen/src/LU/Inverse.h index a5ae83bf4..57f9f686c 100644 --- a/Eigen/src/LU/Inverse.h +++ b/Eigen/src/LU/Inverse.h @@ -21,6 +21,7 @@ namespace internal { template struct compute_inverse { + EIGEN_DEVICE_FUNC static inline void run(const MatrixType& matrix, ResultType& result) { result = matrix.partialPivLu().inverse(); @@ -37,6 +38,7 @@ struct compute_inverse_and_det_with_check { /* nothing! general case not support template struct compute_inverse { + EIGEN_DEVICE_FUNC static inline void run(const MatrixType& matrix, ResultType& result) { typedef typename MatrixType::Scalar Scalar; @@ -47,6 +49,7 @@ struct compute_inverse template struct compute_inverse_and_det_with_check { + EIGEN_DEVICE_FUNC static inline void run( const MatrixType& matrix, const typename MatrixType::RealScalar& absDeterminantThreshold, @@ -67,6 +70,7 @@ struct compute_inverse_and_det_with_check ****************************/ template +EIGEN_DEVICE_FUNC inline void compute_inverse_size2_helper( const MatrixType& matrix, const typename ResultType::Scalar& invdet, ResultType& result) @@ -80,6 +84,7 @@ inline void compute_inverse_size2_helper( template struct compute_inverse { + EIGEN_DEVICE_FUNC static inline void run(const MatrixType& matrix, ResultType& result) { typedef typename ResultType::Scalar Scalar; @@ -91,6 +96,7 @@ struct compute_inverse template struct compute_inverse_and_det_with_check { + EIGEN_DEVICE_FUNC static inline void run( const MatrixType& matrix, const typename MatrixType::RealScalar& absDeterminantThreshold, @@ -114,6 +120,7 @@ struct compute_inverse_and_det_with_check ****************************/ template +EIGEN_DEVICE_FUNC inline typename MatrixType::Scalar cofactor_3x3(const MatrixType& m) { enum { @@ -127,6 +134,7 @@ inline typename MatrixType::Scalar cofactor_3x3(const MatrixType& m) } template +EIGEN_DEVICE_FUNC inline void compute_inverse_size3_helper( const MatrixType& matrix, const typename ResultType::Scalar& invdet, @@ -145,6 +153,7 @@ inline void compute_inverse_size3_helper( template struct compute_inverse { + EIGEN_DEVICE_FUNC static inline void run(const MatrixType& matrix, ResultType& result) { typedef typename ResultType::Scalar Scalar; @@ -161,6 +170,7 @@ struct compute_inverse template struct compute_inverse_and_det_with_check { + EIGEN_DEVICE_FUNC static inline void run( const MatrixType& matrix, const typename MatrixType::RealScalar& absDeterminantThreshold, @@ -188,6 +198,7 @@ struct compute_inverse_and_det_with_check ****************************/ template +EIGEN_DEVICE_FUNC inline const typename Derived::Scalar general_det3_helper (const MatrixBase& matrix, int i1, int i2, int i3, int j1, int j2, int j3) { @@ -196,6 +207,7 @@ inline const typename Derived::Scalar general_det3_helper } template +EIGEN_DEVICE_FUNC inline typename MatrixType::Scalar cofactor_4x4(const MatrixType& matrix) { enum { @@ -214,6 +226,7 @@ inline typename MatrixType::Scalar cofactor_4x4(const MatrixType& matrix) template struct compute_inverse_size4 { + EIGEN_DEVICE_FUNC static void run(const MatrixType& matrix, ResultType& result) { result.coeffRef(0,0) = cofactor_4x4(matrix); @@ -246,6 +259,7 @@ struct compute_inverse template struct compute_inverse_and_det_with_check { + EIGEN_DEVICE_FUNC static inline void run( const MatrixType& matrix, const typename MatrixType::RealScalar& absDeterminantThreshold, @@ -279,14 +293,17 @@ struct inverse_impl : public ReturnByValue > typedef typename remove_all::type MatrixTypeNestedCleaned; MatrixTypeNested m_matrix; + EIGEN_DEVICE_FUNC inverse_impl(const MatrixType& matrix) : m_matrix(matrix) {} - inline Index rows() const { return m_matrix.rows(); } - inline Index cols() const { return m_matrix.cols(); } + EIGEN_DEVICE_FUNC inline Index rows() const { return m_matrix.rows(); } + EIGEN_DEVICE_FUNC inline Index cols() const { return m_matrix.cols(); } - template inline void evalTo(Dest& dst) const + template + EIGEN_DEVICE_FUNC + inline void evalTo(Dest& dst) const { const int Size = EIGEN_PLAIN_ENUM_MIN(MatrixType::ColsAtCompileTime,Dest::ColsAtCompileTime); EIGEN_ONLY_USED_FOR_DEBUG(Size); diff --git a/Eigen/src/LU/PartialPivLU.h b/Eigen/src/LU/PartialPivLU.h index c9ff9dd5a..9cf1d61d8 100644 --- a/Eigen/src/LU/PartialPivLU.h +++ b/Eigen/src/LU/PartialPivLU.h @@ -469,12 +469,14 @@ struct solve_retval, Rhs> * * \sa class PartialPivLU */ +#ifndef __CUDACC__ template inline const PartialPivLU::PlainObject> MatrixBase::partialPivLu() const { return PartialPivLU(eval()); } +#endif #if EIGEN2_SUPPORT_STAGE > STAGE20_RESOLVE_API_CONFLICTS /** \lu_module @@ -485,6 +487,7 @@ MatrixBase::partialPivLu() const * * \sa class PartialPivLU */ +#ifndef __CUDACC__ template inline const PartialPivLU::PlainObject> MatrixBase::lu() const @@ -493,6 +496,8 @@ MatrixBase::lu() const } #endif +#endif + } // end namespace Eigen #endif // EIGEN_PARTIALLU_H diff --git a/Eigen/src/plugins/ArrayCwiseUnaryOps.h b/Eigen/src/plugins/ArrayCwiseUnaryOps.h index a59636790..afb0fb6e3 100644 --- a/Eigen/src/plugins/ArrayCwiseUnaryOps.h +++ b/Eigen/src/plugins/ArrayCwiseUnaryOps.h @@ -7,6 +7,7 @@ * * \sa abs2() */ +EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE const CwiseUnaryOp, const Derived> abs() const { @@ -20,6 +21,7 @@ abs() const * * \sa abs(), square() */ +EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE const CwiseUnaryOp, const Derived> abs2() const { @@ -33,6 +35,7 @@ abs2() const * * \sa pow(), log(), sin(), cos() */ +EIGEN_DEVICE_FUNC inline const CwiseUnaryOp, const Derived> exp() const { @@ -46,6 +49,7 @@ exp() const * * \sa exp() */ +EIGEN_DEVICE_FUNC inline const CwiseUnaryOp, const Derived> log() const { @@ -59,6 +63,7 @@ log() const * * \sa pow(), square() */ +EIGEN_DEVICE_FUNC inline const CwiseUnaryOp, const Derived> sqrt() const { @@ -72,6 +77,7 @@ sqrt() const * * \sa sin(), acos() */ +EIGEN_DEVICE_FUNC inline const CwiseUnaryOp, const Derived> cos() const { @@ -86,6 +92,7 @@ cos() const * * \sa cos(), asin() */ +EIGEN_DEVICE_FUNC inline const CwiseUnaryOp, const Derived> sin() const { @@ -99,6 +106,7 @@ sin() const * * \sa cos(), asin() */ +EIGEN_DEVICE_FUNC inline const CwiseUnaryOp, const Derived> acos() const { @@ -112,6 +120,7 @@ acos() const * * \sa sin(), acos() */ +EIGEN_DEVICE_FUNC inline const CwiseUnaryOp, const Derived> asin() const { @@ -125,6 +134,7 @@ asin() const * * \sa cos(), sin() */ +EIGEN_DEVICE_FUNC inline const CwiseUnaryOp, Derived> tan() const { @@ -139,6 +149,7 @@ tan() const * * \sa exp(), log() */ +EIGEN_DEVICE_FUNC inline const CwiseUnaryOp, const Derived> pow(const Scalar& exponent) const { @@ -154,6 +165,7 @@ pow(const Scalar& exponent) const * * \sa operator/(), operator*() */ +EIGEN_DEVICE_FUNC inline const CwiseUnaryOp, const Derived> inverse() const { @@ -167,6 +179,7 @@ inverse() const * * \sa operator/(), operator*(), abs2() */ +EIGEN_DEVICE_FUNC inline const CwiseUnaryOp, const Derived> square() const { @@ -180,6 +193,7 @@ square() const * * \sa square(), pow() */ +EIGEN_DEVICE_FUNC inline const CwiseUnaryOp, const Derived> cube() const { diff --git a/Eigen/src/plugins/BlockMethods.h b/Eigen/src/plugins/BlockMethods.h index 19a491cf7..5ef373a81 100644 --- a/Eigen/src/plugins/BlockMethods.h +++ b/Eigen/src/plugins/BlockMethods.h @@ -53,12 +53,14 @@ template struct ConstFixedSegmentReturnType { typedef const VectorBloc * * \sa class Block, block(Index,Index) */ +EIGEN_DEVICE_FUNC inline Block block(Index startRow, Index startCol, Index blockRows, Index blockCols) { return Block(derived(), startRow, startCol, blockRows, blockCols); } /** This is the const version of block(Index,Index,Index,Index). */ +EIGEN_DEVICE_FUNC inline const Block block(Index startRow, Index startCol, Index blockRows, Index blockCols) const { return Block(derived(), startRow, startCol, blockRows, blockCols); @@ -77,12 +79,14 @@ inline const Block block(Index startRow, Index startCol, Index bl * * \sa class Block, block(Index,Index,Index,Index) */ +EIGEN_DEVICE_FUNC inline Block topRightCorner(Index cRows, Index cCols) { return Block(derived(), 0, cols() - cCols, cRows, cCols); } /** This is the const version of topRightCorner(Index, Index).*/ +EIGEN_DEVICE_FUNC inline const Block topRightCorner(Index cRows, Index cCols) const { return Block(derived(), 0, cols() - cCols, cRows, cCols); @@ -98,6 +102,7 @@ inline const Block topRightCorner(Index cRows, Index cCols) const * \sa class Block, block(Index,Index,Index,Index) */ template +EIGEN_DEVICE_FUNC inline Block topRightCorner() { return Block(derived(), 0, cols() - CCols); @@ -105,6 +110,7 @@ inline Block topRightCorner() /** This is the const version of topRightCorner().*/ template +EIGEN_DEVICE_FUNC inline const Block topRightCorner() const { return Block(derived(), 0, cols() - CCols); @@ -123,12 +129,14 @@ inline const Block topRightCorner() const * * \sa class Block, block(Index,Index,Index,Index) */ +EIGEN_DEVICE_FUNC inline Block topLeftCorner(Index cRows, Index cCols) { return Block(derived(), 0, 0, cRows, cCols); } /** This is the const version of topLeftCorner(Index, Index).*/ +EIGEN_DEVICE_FUNC inline const Block topLeftCorner(Index cRows, Index cCols) const { return Block(derived(), 0, 0, cRows, cCols); @@ -144,6 +152,7 @@ inline const Block topLeftCorner(Index cRows, Index cCols) const * \sa class Block, block(Index,Index,Index,Index) */ template +EIGEN_DEVICE_FUNC inline Block topLeftCorner() { return Block(derived(), 0, 0); @@ -151,6 +160,7 @@ inline Block topLeftCorner() /** This is the const version of topLeftCorner().*/ template +EIGEN_DEVICE_FUNC inline const Block topLeftCorner() const { return Block(derived(), 0, 0); @@ -168,12 +178,14 @@ inline const Block topLeftCorner() const * * \sa class Block, block(Index,Index,Index,Index) */ +EIGEN_DEVICE_FUNC inline Block bottomRightCorner(Index cRows, Index cCols) { return Block(derived(), rows() - cRows, cols() - cCols, cRows, cCols); } /** This is the const version of bottomRightCorner(Index, Index).*/ +EIGEN_DEVICE_FUNC inline const Block bottomRightCorner(Index cRows, Index cCols) const { return Block(derived(), rows() - cRows, cols() - cCols, cRows, cCols); @@ -189,6 +201,7 @@ inline const Block bottomRightCorner(Index cRows, Index cCols) co * \sa class Block, block(Index,Index,Index,Index) */ template +EIGEN_DEVICE_FUNC inline Block bottomRightCorner() { return Block(derived(), rows() - CRows, cols() - CCols); @@ -196,6 +209,7 @@ inline Block bottomRightCorner() /** This is the const version of bottomRightCorner().*/ template +EIGEN_DEVICE_FUNC inline const Block bottomRightCorner() const { return Block(derived(), rows() - CRows, cols() - CCols); @@ -213,12 +227,14 @@ inline const Block bottomRightCorner() const * * \sa class Block, block(Index,Index,Index,Index) */ +EIGEN_DEVICE_FUNC inline Block bottomLeftCorner(Index cRows, Index cCols) { return Block(derived(), rows() - cRows, 0, cRows, cCols); } /** This is the const version of bottomLeftCorner(Index, Index).*/ +EIGEN_DEVICE_FUNC inline const Block bottomLeftCorner(Index cRows, Index cCols) const { return Block(derived(), rows() - cRows, 0, cRows, cCols); @@ -234,6 +250,7 @@ inline const Block bottomLeftCorner(Index cRows, Index cCols) con * \sa class Block, block(Index,Index,Index,Index) */ template +EIGEN_DEVICE_FUNC inline Block bottomLeftCorner() { return Block(derived(), rows() - CRows, 0); @@ -241,6 +258,7 @@ inline Block bottomLeftCorner() /** This is the const version of bottomLeftCorner().*/ template +EIGEN_DEVICE_FUNC inline const Block bottomLeftCorner() const { return Block(derived(), rows() - CRows, 0); @@ -257,12 +275,14 @@ inline const Block bottomLeftCorner() const * * \sa class Block, block(Index,Index,Index,Index) */ +EIGEN_DEVICE_FUNC inline RowsBlockXpr topRows(Index n) { return RowsBlockXpr(derived(), 0, 0, n, cols()); } /** This is the const version of topRows(Index).*/ +EIGEN_DEVICE_FUNC inline ConstRowsBlockXpr topRows(Index n) const { return ConstRowsBlockXpr(derived(), 0, 0, n, cols()); @@ -278,6 +298,7 @@ inline ConstRowsBlockXpr topRows(Index n) const * \sa class Block, block(Index,Index,Index,Index) */ template +EIGEN_DEVICE_FUNC inline typename NRowsBlockXpr::Type topRows() { return typename NRowsBlockXpr::Type(derived(), 0, 0, N, cols()); @@ -285,6 +306,7 @@ inline typename NRowsBlockXpr::Type topRows() /** This is the const version of topRows().*/ template +EIGEN_DEVICE_FUNC inline typename ConstNRowsBlockXpr::Type topRows() const { return typename ConstNRowsBlockXpr::Type(derived(), 0, 0, N, cols()); @@ -301,12 +323,14 @@ inline typename ConstNRowsBlockXpr::Type topRows() const * * \sa class Block, block(Index,Index,Index,Index) */ +EIGEN_DEVICE_FUNC inline RowsBlockXpr bottomRows(Index n) { return RowsBlockXpr(derived(), rows() - n, 0, n, cols()); } /** This is the const version of bottomRows(Index).*/ +EIGEN_DEVICE_FUNC inline ConstRowsBlockXpr bottomRows(Index n) const { return ConstRowsBlockXpr(derived(), rows() - n, 0, n, cols()); @@ -322,6 +346,7 @@ inline ConstRowsBlockXpr bottomRows(Index n) const * \sa class Block, block(Index,Index,Index,Index) */ template +EIGEN_DEVICE_FUNC inline typename NRowsBlockXpr::Type bottomRows() { return typename NRowsBlockXpr::Type(derived(), rows() - N, 0, N, cols()); @@ -329,6 +354,7 @@ inline typename NRowsBlockXpr::Type bottomRows() /** This is the const version of bottomRows().*/ template +EIGEN_DEVICE_FUNC inline typename ConstNRowsBlockXpr::Type bottomRows() const { return typename ConstNRowsBlockXpr::Type(derived(), rows() - N, 0, N, cols()); @@ -346,12 +372,14 @@ inline typename ConstNRowsBlockXpr::Type bottomRows() const * * \sa class Block, block(Index,Index,Index,Index) */ +EIGEN_DEVICE_FUNC inline RowsBlockXpr middleRows(Index startRow, Index numRows) { return RowsBlockXpr(derived(), startRow, 0, numRows, cols()); } /** This is the const version of middleRows(Index,Index).*/ +EIGEN_DEVICE_FUNC inline ConstRowsBlockXpr middleRows(Index startRow, Index numRows) const { return ConstRowsBlockXpr(derived(), startRow, 0, numRows, cols()); @@ -368,6 +396,7 @@ inline ConstRowsBlockXpr middleRows(Index startRow, Index numRows) const * \sa class Block, block(Index,Index,Index,Index) */ template +EIGEN_DEVICE_FUNC inline typename NRowsBlockXpr::Type middleRows(Index startRow) { return typename NRowsBlockXpr::Type(derived(), startRow, 0, N, cols()); @@ -375,6 +404,7 @@ inline typename NRowsBlockXpr::Type middleRows(Index startRow) /** This is the const version of middleRows().*/ template +EIGEN_DEVICE_FUNC inline typename ConstNRowsBlockXpr::Type middleRows(Index startRow) const { return typename ConstNRowsBlockXpr::Type(derived(), startRow, 0, N, cols()); @@ -391,12 +421,14 @@ inline typename ConstNRowsBlockXpr::Type middleRows(Index startRow) const * * \sa class Block, block(Index,Index,Index,Index) */ +EIGEN_DEVICE_FUNC inline ColsBlockXpr leftCols(Index n) { return ColsBlockXpr(derived(), 0, 0, rows(), n); } /** This is the const version of leftCols(Index).*/ +EIGEN_DEVICE_FUNC inline ConstColsBlockXpr leftCols(Index n) const { return ConstColsBlockXpr(derived(), 0, 0, rows(), n); @@ -412,6 +444,7 @@ inline ConstColsBlockXpr leftCols(Index n) const * \sa class Block, block(Index,Index,Index,Index) */ template +EIGEN_DEVICE_FUNC inline typename NColsBlockXpr::Type leftCols() { return typename NColsBlockXpr::Type(derived(), 0, 0, rows(), N); @@ -419,6 +452,7 @@ inline typename NColsBlockXpr::Type leftCols() /** This is the const version of leftCols().*/ template +EIGEN_DEVICE_FUNC inline typename ConstNColsBlockXpr::Type leftCols() const { return typename ConstNColsBlockXpr::Type(derived(), 0, 0, rows(), N); @@ -435,12 +469,14 @@ inline typename ConstNColsBlockXpr::Type leftCols() const * * \sa class Block, block(Index,Index,Index,Index) */ +EIGEN_DEVICE_FUNC inline ColsBlockXpr rightCols(Index n) { return ColsBlockXpr(derived(), 0, cols() - n, rows(), n); } /** This is the const version of rightCols(Index).*/ +EIGEN_DEVICE_FUNC inline ConstColsBlockXpr rightCols(Index n) const { return ConstColsBlockXpr(derived(), 0, cols() - n, rows(), n); @@ -456,6 +492,7 @@ inline ConstColsBlockXpr rightCols(Index n) const * \sa class Block, block(Index,Index,Index,Index) */ template +EIGEN_DEVICE_FUNC inline typename NColsBlockXpr::Type rightCols() { return typename NColsBlockXpr::Type(derived(), 0, cols() - N, rows(), N); @@ -463,6 +500,7 @@ inline typename NColsBlockXpr::Type rightCols() /** This is the const version of rightCols().*/ template +EIGEN_DEVICE_FUNC inline typename ConstNColsBlockXpr::Type rightCols() const { return typename ConstNColsBlockXpr::Type(derived(), 0, cols() - N, rows(), N); @@ -480,12 +518,14 @@ inline typename ConstNColsBlockXpr::Type rightCols() const * * \sa class Block, block(Index,Index,Index,Index) */ +EIGEN_DEVICE_FUNC inline ColsBlockXpr middleCols(Index startCol, Index numCols) { return ColsBlockXpr(derived(), 0, startCol, rows(), numCols); } /** This is the const version of middleCols(Index,Index).*/ +EIGEN_DEVICE_FUNC inline ConstColsBlockXpr middleCols(Index startCol, Index numCols) const { return ConstColsBlockXpr(derived(), 0, startCol, rows(), numCols); @@ -502,6 +542,7 @@ inline ConstColsBlockXpr middleCols(Index startCol, Index numCols) const * \sa class Block, block(Index,Index,Index,Index) */ template +EIGEN_DEVICE_FUNC inline typename NColsBlockXpr::Type middleCols(Index startCol) { return typename NColsBlockXpr::Type(derived(), 0, startCol, rows(), N); @@ -509,6 +550,7 @@ inline typename NColsBlockXpr::Type middleCols(Index startCol) /** This is the const version of middleCols().*/ template +EIGEN_DEVICE_FUNC inline typename ConstNColsBlockXpr::Type middleCols(Index startCol) const { return typename ConstNColsBlockXpr::Type(derived(), 0, startCol, rows(), N); @@ -533,6 +575,7 @@ inline typename ConstNColsBlockXpr::Type middleCols(Index startCol) const * \sa class Block, block(Index,Index,Index,Index) */ template +EIGEN_DEVICE_FUNC inline Block block(Index startRow, Index startCol) { return Block(derived(), startRow, startCol); @@ -540,6 +583,7 @@ inline Block block(Index startRow, Index startCol /** This is the const version of block<>(Index, Index). */ template +EIGEN_DEVICE_FUNC inline const Block block(Index startRow, Index startCol) const { return Block(derived(), startRow, startCol); @@ -551,12 +595,14 @@ inline const Block block(Index startRow, In * Output: \verbinclude MatrixBase_col.out * * \sa row(), class Block */ +EIGEN_DEVICE_FUNC inline ColXpr col(Index i) { return ColXpr(derived(), i); } /** This is the const version of col(). */ +EIGEN_DEVICE_FUNC inline ConstColXpr col(Index i) const { return ConstColXpr(derived(), i); @@ -568,12 +614,14 @@ inline ConstColXpr col(Index i) const * Output: \verbinclude MatrixBase_row.out * * \sa col(), class Block */ +EIGEN_DEVICE_FUNC inline RowXpr row(Index i) { return RowXpr(derived(), i); } /** This is the const version of row(). */ +EIGEN_DEVICE_FUNC inline ConstRowXpr row(Index i) const { return ConstRowXpr(derived(), i); @@ -595,6 +643,7 @@ inline ConstRowXpr row(Index i) const * * \sa class Block, segment(Index) */ +EIGEN_DEVICE_FUNC inline SegmentReturnType segment(Index start, Index vecSize) { EIGEN_STATIC_ASSERT_VECTOR_ONLY(Derived) @@ -603,6 +652,7 @@ inline SegmentReturnType segment(Index start, Index vecSize) /** This is the const version of segment(Index,Index).*/ +EIGEN_DEVICE_FUNC inline ConstSegmentReturnType segment(Index start, Index vecSize) const { EIGEN_STATIC_ASSERT_VECTOR_ONLY(Derived) @@ -624,6 +674,7 @@ inline ConstSegmentReturnType segment(Index start, Index vecSize) const * * \sa class Block, block(Index,Index) */ +EIGEN_DEVICE_FUNC inline SegmentReturnType head(Index vecSize) { EIGEN_STATIC_ASSERT_VECTOR_ONLY(Derived) @@ -631,6 +682,7 @@ inline SegmentReturnType head(Index vecSize) } /** This is the const version of head(Index).*/ +EIGEN_DEVICE_FUNC inline ConstSegmentReturnType head(Index vecSize) const { @@ -653,6 +705,7 @@ inline ConstSegmentReturnType * * \sa class Block, block(Index,Index) */ +EIGEN_DEVICE_FUNC inline SegmentReturnType tail(Index vecSize) { EIGEN_STATIC_ASSERT_VECTOR_ONLY(Derived) @@ -660,6 +713,7 @@ inline SegmentReturnType tail(Index vecSize) } /** This is the const version of tail(Index).*/ +EIGEN_DEVICE_FUNC inline ConstSegmentReturnType tail(Index vecSize) const { EIGEN_STATIC_ASSERT_VECTOR_ONLY(Derived) @@ -680,6 +734,7 @@ inline ConstSegmentReturnType tail(Index vecSize) const * \sa class Block */ template +EIGEN_DEVICE_FUNC inline typename FixedSegmentReturnType::Type segment(Index start) { EIGEN_STATIC_ASSERT_VECTOR_ONLY(Derived) @@ -688,6 +743,7 @@ inline typename FixedSegmentReturnType::Type segment(Index start) /** This is the const version of segment(Index).*/ template +EIGEN_DEVICE_FUNC inline typename ConstFixedSegmentReturnType::Type segment(Index start) const { EIGEN_STATIC_ASSERT_VECTOR_ONLY(Derived) @@ -706,6 +762,7 @@ inline typename ConstFixedSegmentReturnType::Type segment(Index start) con * \sa class Block */ template +EIGEN_DEVICE_FUNC inline typename FixedSegmentReturnType::Type head() { EIGEN_STATIC_ASSERT_VECTOR_ONLY(Derived) @@ -714,6 +771,7 @@ inline typename FixedSegmentReturnType::Type head() /** This is the const version of head().*/ template +EIGEN_DEVICE_FUNC inline typename ConstFixedSegmentReturnType::Type head() const { EIGEN_STATIC_ASSERT_VECTOR_ONLY(Derived) @@ -732,6 +790,7 @@ inline typename ConstFixedSegmentReturnType::Type head() const * \sa class Block */ template +EIGEN_DEVICE_FUNC inline typename FixedSegmentReturnType::Type tail() { EIGEN_STATIC_ASSERT_VECTOR_ONLY(Derived) @@ -740,6 +799,7 @@ inline typename FixedSegmentReturnType::Type tail() /** This is the const version of tail.*/ template +EIGEN_DEVICE_FUNC inline typename ConstFixedSegmentReturnType::Type tail() const { EIGEN_STATIC_ASSERT_VECTOR_ONLY(Derived) diff --git a/Eigen/src/plugins/CommonCwiseBinaryOps.h b/Eigen/src/plugins/CommonCwiseBinaryOps.h index 688d22440..a8fa287c9 100644 --- a/Eigen/src/plugins/CommonCwiseBinaryOps.h +++ b/Eigen/src/plugins/CommonCwiseBinaryOps.h @@ -38,6 +38,7 @@ EIGEN_MAKE_CWISE_BINARY_OP(operator+,internal::scalar_sum_op) * \sa class CwiseBinaryOp, operator+(), operator-(), cwiseProduct() */ template +EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE const CwiseBinaryOp binaryExpr(const EIGEN_CURRENT_STORAGE_BASE_CLASS &other, const CustomBinaryOp& func = CustomBinaryOp()) const { diff --git a/Eigen/src/plugins/CommonCwiseUnaryOps.h b/Eigen/src/plugins/CommonCwiseUnaryOps.h index 08e931aad..a17153e64 100644 --- a/Eigen/src/plugins/CommonCwiseUnaryOps.h +++ b/Eigen/src/plugins/CommonCwiseUnaryOps.h @@ -40,11 +40,13 @@ typedef CwiseUnaryView, Derived> NonConstIm /** \returns an expression of the opposite of \c *this */ +EIGEN_DEVICE_FUNC inline const CwiseUnaryOp::Scalar>, const Derived> operator-() const { return derived(); } /** \returns an expression of \c *this scaled by the scalar factor \a scalar */ +EIGEN_DEVICE_FUNC inline const ScalarMultipleReturnType operator*(const Scalar& scalar) const { @@ -57,6 +59,7 @@ const ScalarMultipleReturnType operator*(const RealScalar& scalar) const; #endif /** \returns an expression of \c *this divided by the scalar value \a scalar */ +EIGEN_DEVICE_FUNC inline const CwiseUnaryOp::Scalar>, const Derived> operator/(const Scalar& scalar) const { @@ -65,6 +68,7 @@ operator/(const Scalar& scalar) const } /** Overloaded for efficient real matrix times complex scalar value */ +EIGEN_DEVICE_FUNC inline const CwiseUnaryOp >, const Derived> operator*(const std::complex& scalar) const { @@ -72,10 +76,12 @@ operator*(const std::complex& scalar) const (*static_cast(this), internal::scalar_multiple2_op >(scalar)); } +EIGEN_DEVICE_FUNC inline friend const ScalarMultipleReturnType operator*(const Scalar& scalar, const StorageBaseType& matrix) { return matrix*scalar; } +EIGEN_DEVICE_FUNC inline friend const CwiseUnaryOp >, const Derived> operator*(const std::complex& scalar, const StorageBaseType& matrix) { return matrix*scalar; } @@ -88,6 +94,7 @@ operator*(const std::complex& scalar, const StorageBaseType& matrix) * \sa class CwiseUnaryOp */ template +EIGEN_DEVICE_FUNC typename internal::cast_return_type::Scalar, NewType>, const Derived> >::type cast() const { @@ -97,6 +104,7 @@ cast() const /** \returns an expression of the complex conjugate of \c *this. * * \sa adjoint() */ +EIGEN_DEVICE_FUNC inline ConjugateReturnType conjugate() const { @@ -106,12 +114,14 @@ conjugate() const /** \returns a read-only expression of the real part of \c *this. * * \sa imag() */ +EIGEN_DEVICE_FUNC inline RealReturnType real() const { return derived(); } /** \returns an read-only expression of the imaginary part of \c *this. * * \sa real() */ +EIGEN_DEVICE_FUNC inline const ImagReturnType imag() const { return derived(); } @@ -135,6 +145,7 @@ imag() const { return derived(); } * \sa class CwiseUnaryOp, class CwiseBinaryOp */ template +EIGEN_DEVICE_FUNC inline const CwiseUnaryOp unaryExpr(const CustomUnaryOp& func = CustomUnaryOp()) const { @@ -153,6 +164,7 @@ unaryExpr(const CustomUnaryOp& func = CustomUnaryOp()) const * \sa class CwiseUnaryOp, class CwiseBinaryOp */ template +EIGEN_DEVICE_FUNC inline const CwiseUnaryView unaryViewExpr(const CustomViewOp& func = CustomViewOp()) const { @@ -162,11 +174,13 @@ unaryViewExpr(const CustomViewOp& func = CustomViewOp()) const /** \returns a non const expression of the real part of \c *this. * * \sa imag() */ +EIGEN_DEVICE_FUNC inline NonConstRealReturnType real() { return derived(); } /** \returns a non const expression of the imaginary part of \c *this. * * \sa real() */ +EIGEN_DEVICE_FUNC inline NonConstImagReturnType imag() { return derived(); } diff --git a/Eigen/src/plugins/MatrixCwiseBinaryOps.h b/Eigen/src/plugins/MatrixCwiseBinaryOps.h index 3a737df7b..52e75a174 100644 --- a/Eigen/src/plugins/MatrixCwiseBinaryOps.h +++ b/Eigen/src/plugins/MatrixCwiseBinaryOps.h @@ -18,6 +18,7 @@ * \sa class CwiseBinaryOp, cwiseAbs2 */ template +EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE const EIGEN_CWISE_PRODUCT_RETURN_TYPE(Derived,OtherDerived) cwiseProduct(const EIGEN_CURRENT_STORAGE_BASE_CLASS &other) const { @@ -37,6 +38,7 @@ cwiseProduct(const EIGEN_CURRENT_STORAGE_BASE_CLASS &other) const * \sa cwiseNotEqual(), isApprox(), isMuchSmallerThan() */ template +EIGEN_DEVICE_FUNC inline const CwiseBinaryOp, const Derived, const OtherDerived> cwiseEqual(const EIGEN_CURRENT_STORAGE_BASE_CLASS &other) const { @@ -56,6 +58,7 @@ cwiseEqual(const EIGEN_CURRENT_STORAGE_BASE_CLASS &other) const * \sa cwiseEqual(), isApprox(), isMuchSmallerThan() */ template +EIGEN_DEVICE_FUNC inline const CwiseBinaryOp, const Derived, const OtherDerived> cwiseNotEqual(const EIGEN_CURRENT_STORAGE_BASE_CLASS &other) const { @@ -70,6 +73,7 @@ cwiseNotEqual(const EIGEN_CURRENT_STORAGE_BASE_CLASS &other) const * \sa class CwiseBinaryOp, max() */ template +EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE const CwiseBinaryOp, const Derived, const OtherDerived> cwiseMin(const EIGEN_CURRENT_STORAGE_BASE_CLASS &other) const { @@ -80,6 +84,7 @@ cwiseMin(const EIGEN_CURRENT_STORAGE_BASE_CLASS &other) const * * \sa class CwiseBinaryOp, min() */ +EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE const CwiseBinaryOp, const Derived, const ConstantReturnType> cwiseMin(const Scalar &other) const { @@ -94,6 +99,7 @@ cwiseMin(const Scalar &other) const * \sa class CwiseBinaryOp, min() */ template +EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE const CwiseBinaryOp, const Derived, const OtherDerived> cwiseMax(const EIGEN_CURRENT_STORAGE_BASE_CLASS &other) const { @@ -104,6 +110,7 @@ cwiseMax(const EIGEN_CURRENT_STORAGE_BASE_CLASS &other) const * * \sa class CwiseBinaryOp, min() */ +EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE const CwiseBinaryOp, const Derived, const ConstantReturnType> cwiseMax(const Scalar &other) const { @@ -119,6 +126,7 @@ cwiseMax(const Scalar &other) const * \sa class CwiseBinaryOp, cwiseProduct(), cwiseInverse() */ template +EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE const CwiseBinaryOp, const Derived, const OtherDerived> cwiseQuotient(const EIGEN_CURRENT_STORAGE_BASE_CLASS &other) const { diff --git a/Eigen/src/plugins/MatrixCwiseUnaryOps.h b/Eigen/src/plugins/MatrixCwiseUnaryOps.h index 0cf0640ba..1bb15f862 100644 --- a/Eigen/src/plugins/MatrixCwiseUnaryOps.h +++ b/Eigen/src/plugins/MatrixCwiseUnaryOps.h @@ -17,6 +17,7 @@ * * \sa cwiseAbs2() */ +EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE const CwiseUnaryOp, const Derived> cwiseAbs() const { return derived(); } @@ -27,6 +28,7 @@ cwiseAbs() const { return derived(); } * * \sa cwiseAbs() */ +EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE const CwiseUnaryOp, const Derived> cwiseAbs2() const { return derived(); } @@ -37,6 +39,7 @@ cwiseAbs2() const { return derived(); } * * \sa cwisePow(), cwiseSquare() */ +EIGEN_DEVICE_FUNC inline const CwiseUnaryOp, const Derived> cwiseSqrt() const { return derived(); } @@ -47,6 +50,7 @@ cwiseSqrt() const { return derived(); } * * \sa cwiseProduct() */ +EIGEN_DEVICE_FUNC inline const CwiseUnaryOp, const Derived> cwiseInverse() const { return derived(); } @@ -59,6 +63,7 @@ cwiseInverse() const { return derived(); } * * \sa cwiseEqual(const MatrixBase &) const */ +EIGEN_DEVICE_FUNC inline const CwiseUnaryOp >, const Derived> cwiseEqual(const Scalar& s) const { -- cgit v1.2.3