diff --git a/Eigen/Core b/Eigen/Core index dd8da38b5..87f87cd70 100644 --- a/Eigen/Core +++ b/Eigen/Core @@ -18,12 +18,21 @@ #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__ + #else #define EIGEN_DEVICE_FUNC + +#endif + +#if defined(__CUDA_ARCH__) + // Do not try asserts on CUDA! + #define EIGEN_NO_DEBUG + #define EIGEN_USING_STD_MATH(FUNC) +#else + #define EIGEN_USING_STD_MATH(FUNC) using std::FUNC; #endif // then include this file where all our macros are defined. It's really important to do it first because diff --git a/Eigen/src/Cholesky/LDLT.h b/Eigen/src/Cholesky/LDLT.h index 219a0e8b9..5dede8081 100644 --- a/Eigen/src/Cholesky/LDLT.h +++ b/Eigen/src/Cholesky/LDLT.h @@ -501,7 +501,7 @@ struct solve_retval, Rhs> // dst = D^-1 (L^-1 P b) // more precisely, use pseudo-inverse of D (see bug 241) using std::abs; - using std::max; + EIGEN_USING_STD_MATH(max); typedef typename LDLTType::MatrixType MatrixType; typedef typename LDLTType::Scalar Scalar; typedef typename LDLTType::RealScalar RealScalar; diff --git a/Eigen/src/Core/CwiseNullaryOp.h b/Eigen/src/Core/CwiseNullaryOp.h index a93bab2d0..14c78cb85 100644 --- a/Eigen/src/Core/CwiseNullaryOp.h +++ b/Eigen/src/Core/CwiseNullaryOp.h @@ -54,6 +54,7 @@ class CwiseNullaryOp : internal::no_assignment_operator, typedef typename internal::dense_xpr_base::type Base; EIGEN_DENSE_PUBLIC_INTERFACE(CwiseNullaryOp) + EIGEN_DEVICE_FUNC CwiseNullaryOp(Index nbRows, Index nbCols, const NullaryOp& func = NullaryOp()) : m_rows(nbRows), m_cols(nbCols), m_functor(func) { @@ -63,9 +64,12 @@ class CwiseNullaryOp : internal::no_assignment_operator, && (ColsAtCompileTime == Dynamic || ColsAtCompileTime == nbCols)); } + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Index rows() const { return m_rows.value(); } + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Index cols() const { return m_cols.value(); } + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE const Scalar coeff(Index rowId, Index colId) const { return m_functor(rowId, colId); @@ -77,6 +81,7 @@ class CwiseNullaryOp : internal::no_assignment_operator, return m_functor.packetOp(rowId, colId); } + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE const Scalar coeff(Index index) const { return m_functor(index); @@ -89,6 +94,7 @@ class CwiseNullaryOp : internal::no_assignment_operator, } /** \returns the functor representing the nullary operation */ + EIGEN_DEVICE_FUNC const NullaryOp& functor() const { return m_functor; } protected: diff --git a/Eigen/src/Core/Dot.h b/Eigen/src/Core/Dot.h index 9d513d699..6b058d3d9 100644 --- a/Eigen/src/Core/Dot.h +++ b/Eigen/src/Core/Dot.h @@ -59,6 +59,7 @@ struct dot_nocheck */ template template +EIGEN_DEVICE_FUNC typename internal::scalar_product_traits::Scalar,typename internal::traits::Scalar>::ReturnType MatrixBase::dot(const MatrixBase& other) const { @@ -151,6 +152,7 @@ MatrixBase::normalized() const * \sa norm(), normalized() */ template +EIGEN_DEVICE_FUNC inline void MatrixBase::normalize() { *this /= norm(); @@ -164,6 +166,7 @@ template struct lpNorm_selector { typedef typename NumTraits::Scalar>::Real RealScalar; + EIGEN_DEVICE_FUNC static inline RealScalar run(const MatrixBase& m) { return pow(m.cwiseAbs().array().pow(p).sum(), RealScalar(1)/p); @@ -173,6 +176,7 @@ struct lpNorm_selector template struct lpNorm_selector { + EIGEN_DEVICE_FUNC static inline typename NumTraits::Scalar>::Real run(const MatrixBase& m) { return m.cwiseAbs().sum(); @@ -182,6 +186,7 @@ struct lpNorm_selector template struct lpNorm_selector { + EIGEN_DEVICE_FUNC static inline typename NumTraits::Scalar>::Real run(const MatrixBase& m) { return m.norm(); @@ -191,6 +196,7 @@ struct lpNorm_selector template struct lpNorm_selector { + EIGEN_DEVICE_FUNC static inline typename NumTraits::Scalar>::Real run(const MatrixBase& m) { return m.cwiseAbs().maxCoeff(); diff --git a/Eigen/src/Core/Functors.h b/Eigen/src/Core/Functors.h index 6b6b36656..29cd739ce 100644 --- a/Eigen/src/Core/Functors.h +++ b/Eigen/src/Core/Functors.h @@ -103,7 +103,7 @@ struct functor_traits > { */ template struct scalar_min_op { EIGEN_EMPTY_STRUCT_CTOR(scalar_min_op) - EIGEN_DEVICE_FUNC 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 { EIGEN_USING_STD_MATH(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_DEVICE_FUNC 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 { EIGEN_USING_STD_MATH(max); return (max)(a, b); } template EIGEN_STRONG_INLINE const Packet packetOp(const Packet& a, const Packet& b) const { return internal::pmax(a,b); } @@ -152,8 +152,8 @@ template struct scalar_hypot_op { // typedef typename NumTraits::Real result_type; EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE const Scalar operator() (const Scalar& _x, const Scalar& _y) const { - using std::max; - using std::min; + EIGEN_USING_STD_MATH(max); + EIGEN_USING_STD_MATH(min); Scalar p = (max)(_x, _y); Scalar q = (min)(_x, _y); Scalar qp = q/p; @@ -479,8 +479,8 @@ struct functor_traits > template 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_DEVICE_FUNC EIGEN_STRONG_INLINE scalar_multiple2_op(const scalar_multiple2_op& other) : m_other(other.m_other) { } + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE scalar_multiple2_op(const Scalar2& other) : m_other(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; }; @@ -500,8 +500,8 @@ template struct scalar_quotient1_op { typedef typename packet_traits::type Packet; // 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_DEVICE_FUNC EIGEN_STRONG_INLINE scalar_quotient1_op(const scalar_quotient1_op& other) : m_other(other.m_other) { } + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE scalar_quotient1_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::pdiv(a, pset1(m_other)); } @@ -516,8 +516,8 @@ struct functor_traits > template struct scalar_constant_op { typedef typename packet_traits::type Packet; - 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) { } + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE scalar_constant_op(const scalar_constant_op& other) : m_other(other.m_other) { } + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE scalar_constant_op(const Scalar& other) : m_other(other) { } template EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE const Scalar operator() (Index, Index = 0) const { return m_other; } template diff --git a/Eigen/src/Core/Fuzzy.h b/Eigen/src/Core/Fuzzy.h index 8fb9a01dd..9581c81bd 100644 --- a/Eigen/src/Core/Fuzzy.h +++ b/Eigen/src/Core/Fuzzy.h @@ -21,7 +21,7 @@ struct isApprox_selector { static bool run(const Derived& x, const OtherDerived& y, const typename Derived::RealScalar& prec) { - using std::min; + EIGEN_USING_STD_MATH(min); typename internal::nested::type nested(x); typename internal::nested::type otherNested(y); return (nested - otherNested).cwiseAbs2().sum() <= prec * prec * (min)(nested.cwiseAbs2().sum(), otherNested.cwiseAbs2().sum()); diff --git a/Eigen/src/Core/GenericPacketMath.h b/Eigen/src/Core/GenericPacketMath.h index 51913fa5c..aee58f09b 100644 --- a/Eigen/src/Core/GenericPacketMath.h +++ b/Eigen/src/Core/GenericPacketMath.h @@ -121,12 +121,12 @@ pdiv(const Packet& a, /** \internal \returns the min of \a a and \a b (coeff-wise) */ template inline Packet pmin(const Packet& a, - const Packet& b) { using std::min; return (min)(a, b); } + const Packet& b) { EIGEN_USING_STD_MATH(min); return (min)(a, b); } /** \internal \returns the max of \a a and \a b (coeff-wise) */ template inline Packet pmax(const Packet& a, - const Packet& b) { using std::max; return (max)(a, b); } + const Packet& b) { EIGEN_USING_STD_MATH(max); return (max)(a, b); } /** \internal \returns the absolute value of \a a */ template inline Packet diff --git a/Eigen/src/Core/Map.h b/Eigen/src/Core/Map.h index e054f2202..8ea13cfb7 100644 --- a/Eigen/src/Core/Map.h +++ b/Eigen/src/Core/Map.h @@ -115,6 +115,7 @@ template class Ma inline PointerType cast_to_pointer_type(PointerArgType ptr) { return const_cast(ptr); } #else typedef PointerType PointerArgType; + EIGEN_DEVICE_FUNC inline PointerType cast_to_pointer_type(PointerArgType ptr) { return ptr; } #endif diff --git a/Eigen/src/Core/MapBase.h b/Eigen/src/Core/MapBase.h index 8def7442d..ffa1371c2 100644 --- a/Eigen/src/Core/MapBase.h +++ b/Eigen/src/Core/MapBase.h @@ -203,7 +203,9 @@ template class MapBase const Scalar >::type ScalarWithConstIfNotLvalue; + EIGEN_DEVICE_FUNC inline const Scalar* data() const { return this->m_data; } + EIGEN_DEVICE_FUNC inline ScalarWithConstIfNotLvalue* data() { return this->m_data; } // no const-cast here so non-const-correct code will give a compile error EIGEN_DEVICE_FUNC diff --git a/Eigen/src/Core/MathFunctions.h b/Eigen/src/Core/MathFunctions.h index 4ce616746..b8d32ecec 100644 --- a/Eigen/src/Core/MathFunctions.h +++ b/Eigen/src/Core/MathFunctions.h @@ -63,6 +63,7 @@ template struct real_impl { typedef typename NumTraits::Real RealScalar; + EIGEN_DEVICE_FUNC static inline RealScalar run(const Scalar& x) { return x; @@ -72,6 +73,7 @@ struct real_impl template struct real_impl > { + EIGEN_DEVICE_FUNC static inline RealScalar run(const std::complex& x) { using std::real; @@ -86,6 +88,7 @@ struct real_retval }; template +EIGEN_DEVICE_FUNC inline EIGEN_MATHFUNC_RETVAL(real, Scalar) real(const Scalar& x) { return EIGEN_MATHFUNC_IMPL(real, Scalar)::run(x); @@ -99,6 +102,7 @@ template struct imag_impl { typedef typename NumTraits::Real RealScalar; + EIGEN_DEVICE_FUNC static inline RealScalar run(const Scalar&) { return RealScalar(0); @@ -108,6 +112,7 @@ struct imag_impl template struct imag_impl > { + EIGEN_DEVICE_FUNC static inline RealScalar run(const std::complex& x) { using std::imag; @@ -122,6 +127,7 @@ struct imag_retval }; template +EIGEN_DEVICE_FUNC inline EIGEN_MATHFUNC_RETVAL(imag, Scalar) imag(const Scalar& x) { return EIGEN_MATHFUNC_IMPL(imag, Scalar)::run(x); @@ -135,10 +141,12 @@ template struct real_ref_impl { typedef typename NumTraits::Real RealScalar; + EIGEN_DEVICE_FUNC static inline RealScalar& run(Scalar& x) { return reinterpret_cast(&x)[0]; } + EIGEN_DEVICE_FUNC static inline const RealScalar& run(const Scalar& x) { return reinterpret_cast(&x)[0]; @@ -152,12 +160,14 @@ struct real_ref_retval }; template +EIGEN_DEVICE_FUNC inline typename add_const_on_value_type< EIGEN_MATHFUNC_RETVAL(real_ref, Scalar) >::type real_ref(const Scalar& x) { return real_ref_impl::run(x); } template +EIGEN_DEVICE_FUNC inline EIGEN_MATHFUNC_RETVAL(real_ref, Scalar) real_ref(Scalar& x) { return EIGEN_MATHFUNC_IMPL(real_ref, Scalar)::run(x); @@ -171,10 +181,12 @@ template struct imag_ref_default_impl { typedef typename NumTraits::Real RealScalar; + EIGEN_DEVICE_FUNC static inline RealScalar& run(Scalar& x) { return reinterpret_cast(&x)[1]; } + EIGEN_DEVICE_FUNC static inline const RealScalar& run(const Scalar& x) { return reinterpret_cast(&x)[1]; @@ -184,10 +196,12 @@ struct imag_ref_default_impl template struct imag_ref_default_impl { + EIGEN_DEVICE_FUNC static inline Scalar run(Scalar&) { return Scalar(0); } + EIGEN_DEVICE_FUNC static inline const Scalar run(const Scalar&) { return Scalar(0); @@ -204,12 +218,14 @@ struct imag_ref_retval }; template +EIGEN_DEVICE_FUNC inline typename add_const_on_value_type< EIGEN_MATHFUNC_RETVAL(imag_ref, Scalar) >::type imag_ref(const Scalar& x) { return imag_ref_impl::run(x); } template +EIGEN_DEVICE_FUNC inline EIGEN_MATHFUNC_RETVAL(imag_ref, Scalar) imag_ref(Scalar& x) { return EIGEN_MATHFUNC_IMPL(imag_ref, Scalar)::run(x); @@ -222,6 +238,7 @@ inline EIGEN_MATHFUNC_RETVAL(imag_ref, Scalar) imag_ref(Scalar& x) template struct conj_impl { + EIGEN_DEVICE_FUNC static inline Scalar run(const Scalar& x) { return x; @@ -231,6 +248,7 @@ struct conj_impl template struct conj_impl > { + EIGEN_DEVICE_FUNC static inline std::complex run(const std::complex& x) { using std::conj; @@ -245,6 +263,7 @@ struct conj_retval }; template +EIGEN_DEVICE_FUNC inline EIGEN_MATHFUNC_RETVAL(conj, Scalar) conj(const Scalar& x) { return EIGEN_MATHFUNC_IMPL(conj, Scalar)::run(x); @@ -258,6 +277,7 @@ template struct abs2_impl { typedef typename NumTraits::Real RealScalar; + EIGEN_DEVICE_FUNC static inline RealScalar run(const Scalar& x) { return x*x; @@ -267,6 +287,7 @@ struct abs2_impl template struct abs2_impl > { + EIGEN_DEVICE_FUNC static inline RealScalar run(const std::complex& x) { return real(x)*real(x) + imag(x)*imag(x); @@ -280,6 +301,7 @@ struct abs2_retval }; template +EIGEN_DEVICE_FUNC inline EIGEN_MATHFUNC_RETVAL(abs2, Scalar) abs2(const Scalar& x) { return EIGEN_MATHFUNC_IMPL(abs2, Scalar)::run(x); @@ -293,6 +315,7 @@ template struct norm1_default_impl { typedef typename NumTraits::Real RealScalar; + EIGEN_DEVICE_FUNC static inline RealScalar run(const Scalar& x) { using std::abs; @@ -303,6 +326,7 @@ struct norm1_default_impl template struct norm1_default_impl { + EIGEN_DEVICE_FUNC static inline Scalar run(const Scalar& x) { using std::abs; @@ -320,6 +344,7 @@ struct norm1_retval }; template +EIGEN_DEVICE_FUNC inline EIGEN_MATHFUNC_RETVAL(norm1, Scalar) norm1(const Scalar& x) { return EIGEN_MATHFUNC_IMPL(norm1, Scalar)::run(x); @@ -335,8 +360,8 @@ struct hypot_impl typedef typename NumTraits::Real RealScalar; static inline RealScalar run(const Scalar& x, const Scalar& y) { - using std::max; - using std::min; + EIGEN_USING_STD_MATH(max); + EIGEN_USING_STD_MATH(min); using std::abs; RealScalar _x = abs(x); RealScalar _y = abs(y); @@ -631,7 +656,7 @@ struct scalar_fuzzy_default_impl } static inline bool isApprox(const Scalar& x, const Scalar& y, const RealScalar& prec) { - using std::min; + EIGEN_USING_STD_MATH(min); using std::abs; return abs(x - y) <= (min)(abs(x), abs(y)) * prec; } @@ -671,7 +696,7 @@ struct scalar_fuzzy_default_impl } static inline bool isApprox(const Scalar& x, const Scalar& y, const RealScalar& prec) { - using std::min; + EIGEN_USING_STD_MATH(min); return abs2(x - y) <= (min)(abs2(x), abs2(y)) * prec * prec; } }; diff --git a/Eigen/src/Core/MatrixBase.h b/Eigen/src/Core/MatrixBase.h index c76192e46..7b9ea8c0a 100644 --- a/Eigen/src/Core/MatrixBase.h +++ b/Eigen/src/Core/MatrixBase.h @@ -219,16 +219,16 @@ template class MatrixBase Scalar eigen2_dot(const MatrixBase& other) const; #endif - RealScalar squaredNorm() const; - RealScalar norm() const; + EIGEN_DEVICE_FUNC RealScalar squaredNorm() const; + EIGEN_DEVICE_FUNC RealScalar norm() const; RealScalar stableNorm() const; RealScalar blueNorm() const; RealScalar hypotNorm() const; - const PlainObject normalized() const; - void normalize(); + EIGEN_DEVICE_FUNC const PlainObject normalized() const; + EIGEN_DEVICE_FUNC void normalize(); - const AdjointReturnType adjoint() const; - void adjointInPlace(); + EIGEN_DEVICE_FUNC const AdjointReturnType adjoint() const; + EIGEN_DEVICE_FUNC void adjointInPlace(); typedef Diagonal DiagonalReturnType; DiagonalReturnType diagonal(); @@ -329,15 +329,15 @@ template class MatrixBase /////////// Array module /////////// - template RealScalar lpNorm() const; + template EIGEN_DEVICE_FUNC RealScalar lpNorm() const; - MatrixBase& matrix() { return *this; } - const MatrixBase& matrix() const { return *this; } + EIGEN_DEVICE_FUNC MatrixBase& matrix() { return *this; } + EIGEN_DEVICE_FUNC const MatrixBase& matrix() const { return *this; } /** \returns an \link Eigen::ArrayBase Array \endlink expression of this matrix * \sa ArrayBase::matrix() */ - ArrayWrapper array() { return derived(); } - const ArrayWrapper array() const { return derived(); } + EIGEN_DEVICE_FUNC ArrayWrapper array() { return derived(); } + EIGEN_DEVICE_FUNC const ArrayWrapper array() const { return derived(); } /////////// LU module /////////// diff --git a/Eigen/src/Core/SelfCwiseBinaryOp.h b/Eigen/src/Core/SelfCwiseBinaryOp.h index 0caf2bab1..40f9eb618 100644 --- a/Eigen/src/Core/SelfCwiseBinaryOp.h +++ b/Eigen/src/Core/SelfCwiseBinaryOp.h @@ -52,21 +52,24 @@ template class SelfCwiseBinaryOp typedef typename internal::packet_traits::type Packet; + EIGEN_DEVICE_FUNC inline SelfCwiseBinaryOp(Lhs& xpr, const BinaryOp& func = BinaryOp()) : m_matrix(xpr), m_functor(func) {} - inline Index rows() const { return m_matrix.rows(); } - inline Index cols() const { return m_matrix.cols(); } - inline Index outerStride() const { return m_matrix.outerStride(); } - inline Index innerStride() const { return m_matrix.innerStride(); } - inline const Scalar* data() const { return m_matrix.data(); } + EIGEN_DEVICE_FUNC inline Index rows() const { return m_matrix.rows(); } + EIGEN_DEVICE_FUNC inline Index cols() const { return m_matrix.cols(); } + EIGEN_DEVICE_FUNC inline Index outerStride() const { return m_matrix.outerStride(); } + EIGEN_DEVICE_FUNC inline Index innerStride() const { return m_matrix.innerStride(); } + EIGEN_DEVICE_FUNC inline const Scalar* data() const { return m_matrix.data(); } // note that this function is needed by assign to correctly align loads/stores // TODO make Assign use .data() + EIGEN_DEVICE_FUNC inline Scalar& coeffRef(Index row, Index col) { EIGEN_STATIC_ASSERT_LVALUE(Lhs) return m_matrix.const_cast_derived().coeffRef(row, col); } + EIGEN_DEVICE_FUNC inline const Scalar& coeffRef(Index row, Index col) const { return m_matrix.coeffRef(row, col); @@ -74,17 +77,20 @@ template class SelfCwiseBinaryOp // note that this function is needed by assign to correctly align loads/stores // TODO make Assign use .data() + EIGEN_DEVICE_FUNC inline Scalar& coeffRef(Index index) { EIGEN_STATIC_ASSERT_LVALUE(Lhs) return m_matrix.const_cast_derived().coeffRef(index); } + EIGEN_DEVICE_FUNC inline const Scalar& coeffRef(Index index) const { return m_matrix.const_cast_derived().coeffRef(index); } template + EIGEN_DEVICE_FUNC void copyCoeff(Index row, Index col, const DenseBase& other) { OtherDerived& _other = other.const_cast_derived(); @@ -95,6 +101,7 @@ template class SelfCwiseBinaryOp } template + EIGEN_DEVICE_FUNC void copyCoeff(Index index, const DenseBase& other) { OtherDerived& _other = other.const_cast_derived(); @@ -125,6 +132,7 @@ template class SelfCwiseBinaryOp // reimplement lazyAssign to handle complex *= real // see CwiseBinaryOp ctor for details template + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE SelfCwiseBinaryOp& lazyAssign(const DenseBase& rhs) { EIGEN_STATIC_ASSERT_SAME_MATRIX_SIZE(Lhs,RhsDerived) @@ -144,17 +152,20 @@ template class SelfCwiseBinaryOp // overloaded to honor evaluation of special matrices // maybe another solution would be to not use SelfCwiseBinaryOp // at first... + EIGEN_DEVICE_FUNC SelfCwiseBinaryOp& operator=(const Rhs& _rhs) { typename internal::nested::type rhs(_rhs); return Base::operator=(rhs); } + EIGEN_DEVICE_FUNC Lhs& expression() const { return m_matrix; } + EIGEN_DEVICE_FUNC const BinaryOp& functor() const { return m_functor; diff --git a/Eigen/src/Core/StableNorm.h b/Eigen/src/Core/StableNorm.h index 35626be21..1aefd91a8 100644 --- a/Eigen/src/Core/StableNorm.h +++ b/Eigen/src/Core/StableNorm.h @@ -36,8 +36,8 @@ blueNorm_impl(const EigenBase& _vec) typedef typename Derived::RealScalar RealScalar; typedef typename Derived::Index Index; using std::pow; - using std::min; - using std::max; + EIGEN_USING_STD_MATH(min); + EIGEN_USING_STD_MATH(max); using std::sqrt; using std::abs; const Derived& vec(_vec.derived()); @@ -141,7 +141,7 @@ template inline typename NumTraits::Scalar>::Real MatrixBase::stableNorm() const { - using std::min; + EIGEN_USING_STD_MATH(min); using std::sqrt; const Index blockSize = 4096; RealScalar scale(0); diff --git a/Eigen/src/Eigenvalues/EigenSolver.h b/Eigen/src/Eigenvalues/EigenSolver.h index 43d0ffa76..201ea620d 100644 --- a/Eigen/src/Eigenvalues/EigenSolver.h +++ b/Eigen/src/Eigenvalues/EigenSolver.h @@ -568,7 +568,7 @@ void EigenSolver::doComputeEigenvectors() } // Overflow control - using std::max; + EIGEN_USING_STD_MATH(max); Scalar t = (max)(abs(m_matT.coeff(i,n-1)),abs(m_matT.coeff(i,n))); if ((eps * t) * t > Scalar(1)) m_matT.block(i, n-1, size-i, 2) /= t; diff --git a/Eigen/src/Geometry/AngleAxis.h b/Eigen/src/Geometry/AngleAxis.h index 553d38c74..f424e6d7d 100644 --- a/Eigen/src/Geometry/AngleAxis.h +++ b/Eigen/src/Geometry/AngleAxis.h @@ -159,8 +159,8 @@ template AngleAxis& AngleAxis::operator=(const QuaternionBase& q) { using std::acos; - using std::min; - using std::max; + EIGEN_USING_STD_MATH(min); + EIGEN_USING_STD_MATH(max); using std::sqrt; Scalar n2 = q.vec().squaredNorm(); if (n2 < NumTraits::dummy_precision()*NumTraits::dummy_precision()) diff --git a/Eigen/src/Geometry/Quaternion.h b/Eigen/src/Geometry/Quaternion.h index e7801abf9..45824fbeb 100644 --- a/Eigen/src/Geometry/Quaternion.h +++ b/Eigen/src/Geometry/Quaternion.h @@ -572,7 +572,7 @@ template template inline Derived& QuaternionBase::setFromTwoVectors(const MatrixBase& a, const MatrixBase& b) { - using std::max; + EIGEN_USING_STD_MATH(max); using std::sqrt; Vector3 v0 = a.normalized(); Vector3 v1 = b.normalized(); diff --git a/Eigen/src/SVD/JacobiSVD.h b/Eigen/src/SVD/JacobiSVD.h index 5741a8ba1..1616fe560 100644 --- a/Eigen/src/SVD/JacobiSVD.h +++ b/Eigen/src/SVD/JacobiSVD.h @@ -766,7 +766,7 @@ JacobiSVD::compute(const MatrixType& matrix, unsig // if this 2x2 sub-matrix is not diagonal already... // notice that this comparison will evaluate to false if any NaN is involved, ensuring that NaN's don't // keep us iterating forever. Similarly, small denormal numbers are considered zero. - using std::max; + EIGEN_USING_STD_MATH(max); RealScalar threshold = (max)(considerAsZero, precision * (max)(abs(m_workMatrix.coeff(p,p)), abs(m_workMatrix.coeff(q,q)))); if((max)(abs(m_workMatrix.coeff(p,q)),abs(m_workMatrix.coeff(q,p))) > threshold)