Port SelfCwiseBinaryOp and Dot.h to nvcc, fix portability issue with std::min/max

This commit is contained in:
Gael Guennebaud 2013-04-05 16:35:49 +02:00
parent d93c1c113b
commit 12439e1249
17 changed files with 104 additions and 44 deletions

View File

@ -18,12 +18,21 @@
#ifdef __CUDACC__ #ifdef __CUDACC__
// Do not try to vectorize on CUDA! // Do not try to vectorize on CUDA!
#define EIGEN_DONT_VECTORIZE #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__ // All functions callable from CUDA code must be qualified with __device__
#define EIGEN_DEVICE_FUNC __host__ __device__ #define EIGEN_DEVICE_FUNC __host__ __device__
#else #else
#define EIGEN_DEVICE_FUNC #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 #endif
// then include this file where all our macros are defined. It's really important to do it first because // then include this file where all our macros are defined. It's really important to do it first because

View File

@ -501,7 +501,7 @@ struct solve_retval<LDLT<_MatrixType,_UpLo>, Rhs>
// dst = D^-1 (L^-1 P b) // dst = D^-1 (L^-1 P b)
// more precisely, use pseudo-inverse of D (see bug 241) // more precisely, use pseudo-inverse of D (see bug 241)
using std::abs; using std::abs;
using std::max; EIGEN_USING_STD_MATH(max);
typedef typename LDLTType::MatrixType MatrixType; typedef typename LDLTType::MatrixType MatrixType;
typedef typename LDLTType::Scalar Scalar; typedef typename LDLTType::Scalar Scalar;
typedef typename LDLTType::RealScalar RealScalar; typedef typename LDLTType::RealScalar RealScalar;

View File

@ -54,6 +54,7 @@ class CwiseNullaryOp : internal::no_assignment_operator,
typedef typename internal::dense_xpr_base<CwiseNullaryOp>::type Base; typedef typename internal::dense_xpr_base<CwiseNullaryOp>::type Base;
EIGEN_DENSE_PUBLIC_INTERFACE(CwiseNullaryOp) EIGEN_DENSE_PUBLIC_INTERFACE(CwiseNullaryOp)
EIGEN_DEVICE_FUNC
CwiseNullaryOp(Index nbRows, Index nbCols, const NullaryOp& func = NullaryOp()) CwiseNullaryOp(Index nbRows, Index nbCols, const NullaryOp& func = NullaryOp())
: m_rows(nbRows), m_cols(nbCols), m_functor(func) : m_rows(nbRows), m_cols(nbCols), m_functor(func)
{ {
@ -63,9 +64,12 @@ class CwiseNullaryOp : internal::no_assignment_operator,
&& (ColsAtCompileTime == Dynamic || ColsAtCompileTime == nbCols)); && (ColsAtCompileTime == Dynamic || ColsAtCompileTime == nbCols));
} }
EIGEN_DEVICE_FUNC
EIGEN_STRONG_INLINE Index rows() const { return m_rows.value(); } EIGEN_STRONG_INLINE Index rows() const { return m_rows.value(); }
EIGEN_DEVICE_FUNC
EIGEN_STRONG_INLINE Index cols() const { return m_cols.value(); } EIGEN_STRONG_INLINE Index cols() const { return m_cols.value(); }
EIGEN_DEVICE_FUNC
EIGEN_STRONG_INLINE const Scalar coeff(Index rowId, Index colId) const EIGEN_STRONG_INLINE const Scalar coeff(Index rowId, Index colId) const
{ {
return m_functor(rowId, colId); return m_functor(rowId, colId);
@ -77,6 +81,7 @@ class CwiseNullaryOp : internal::no_assignment_operator,
return m_functor.packetOp(rowId, colId); return m_functor.packetOp(rowId, colId);
} }
EIGEN_DEVICE_FUNC
EIGEN_STRONG_INLINE const Scalar coeff(Index index) const EIGEN_STRONG_INLINE const Scalar coeff(Index index) const
{ {
return m_functor(index); return m_functor(index);
@ -89,6 +94,7 @@ class CwiseNullaryOp : internal::no_assignment_operator,
} }
/** \returns the functor representing the nullary operation */ /** \returns the functor representing the nullary operation */
EIGEN_DEVICE_FUNC
const NullaryOp& functor() const { return m_functor; } const NullaryOp& functor() const { return m_functor; }
protected: protected:

View File

@ -59,6 +59,7 @@ struct dot_nocheck<T, U, true>
*/ */
template<typename Derived> template<typename Derived>
template<typename OtherDerived> template<typename OtherDerived>
EIGEN_DEVICE_FUNC
typename internal::scalar_product_traits<typename internal::traits<Derived>::Scalar,typename internal::traits<OtherDerived>::Scalar>::ReturnType typename internal::scalar_product_traits<typename internal::traits<Derived>::Scalar,typename internal::traits<OtherDerived>::Scalar>::ReturnType
MatrixBase<Derived>::dot(const MatrixBase<OtherDerived>& other) const MatrixBase<Derived>::dot(const MatrixBase<OtherDerived>& other) const
{ {
@ -151,6 +152,7 @@ MatrixBase<Derived>::normalized() const
* \sa norm(), normalized() * \sa norm(), normalized()
*/ */
template<typename Derived> template<typename Derived>
EIGEN_DEVICE_FUNC
inline void MatrixBase<Derived>::normalize() inline void MatrixBase<Derived>::normalize()
{ {
*this /= norm(); *this /= norm();
@ -164,6 +166,7 @@ template<typename Derived, int p>
struct lpNorm_selector struct lpNorm_selector
{ {
typedef typename NumTraits<typename traits<Derived>::Scalar>::Real RealScalar; typedef typename NumTraits<typename traits<Derived>::Scalar>::Real RealScalar;
EIGEN_DEVICE_FUNC
static inline RealScalar run(const MatrixBase<Derived>& m) static inline RealScalar run(const MatrixBase<Derived>& m)
{ {
return pow(m.cwiseAbs().array().pow(p).sum(), RealScalar(1)/p); return pow(m.cwiseAbs().array().pow(p).sum(), RealScalar(1)/p);
@ -173,6 +176,7 @@ struct lpNorm_selector
template<typename Derived> template<typename Derived>
struct lpNorm_selector<Derived, 1> struct lpNorm_selector<Derived, 1>
{ {
EIGEN_DEVICE_FUNC
static inline typename NumTraits<typename traits<Derived>::Scalar>::Real run(const MatrixBase<Derived>& m) static inline typename NumTraits<typename traits<Derived>::Scalar>::Real run(const MatrixBase<Derived>& m)
{ {
return m.cwiseAbs().sum(); return m.cwiseAbs().sum();
@ -182,6 +186,7 @@ struct lpNorm_selector<Derived, 1>
template<typename Derived> template<typename Derived>
struct lpNorm_selector<Derived, 2> struct lpNorm_selector<Derived, 2>
{ {
EIGEN_DEVICE_FUNC
static inline typename NumTraits<typename traits<Derived>::Scalar>::Real run(const MatrixBase<Derived>& m) static inline typename NumTraits<typename traits<Derived>::Scalar>::Real run(const MatrixBase<Derived>& m)
{ {
return m.norm(); return m.norm();
@ -191,6 +196,7 @@ struct lpNorm_selector<Derived, 2>
template<typename Derived> template<typename Derived>
struct lpNorm_selector<Derived, Infinity> struct lpNorm_selector<Derived, Infinity>
{ {
EIGEN_DEVICE_FUNC
static inline typename NumTraits<typename traits<Derived>::Scalar>::Real run(const MatrixBase<Derived>& m) static inline typename NumTraits<typename traits<Derived>::Scalar>::Real run(const MatrixBase<Derived>& m)
{ {
return m.cwiseAbs().maxCoeff(); return m.cwiseAbs().maxCoeff();

View File

@ -103,7 +103,7 @@ struct functor_traits<scalar_conj_product_op<LhsScalar,RhsScalar> > {
*/ */
template<typename Scalar> struct scalar_min_op { template<typename Scalar> struct scalar_min_op {
EIGEN_EMPTY_STRUCT_CTOR(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<typename Packet> template<typename Packet>
EIGEN_STRONG_INLINE const Packet packetOp(const Packet& a, const Packet& b) const EIGEN_STRONG_INLINE const Packet packetOp(const Packet& a, const Packet& b) const
{ return internal::pmin(a,b); } { return internal::pmin(a,b); }
@ -126,7 +126,7 @@ struct functor_traits<scalar_min_op<Scalar> > {
*/ */
template<typename Scalar> struct scalar_max_op { template<typename Scalar> struct scalar_max_op {
EIGEN_EMPTY_STRUCT_CTOR(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<typename Packet> template<typename Packet>
EIGEN_STRONG_INLINE const Packet packetOp(const Packet& a, const Packet& b) const EIGEN_STRONG_INLINE const Packet packetOp(const Packet& a, const Packet& b) const
{ return internal::pmax(a,b); } { return internal::pmax(a,b); }
@ -152,8 +152,8 @@ template<typename Scalar> struct scalar_hypot_op {
// typedef typename NumTraits<Scalar>::Real result_type; // typedef typename NumTraits<Scalar>::Real result_type;
EIGEN_DEVICE_FUNC 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; EIGEN_USING_STD_MATH(max);
using std::min; EIGEN_USING_STD_MATH(min);
Scalar p = (max)(_x, _y); Scalar p = (max)(_x, _y);
Scalar q = (min)(_x, _y); Scalar q = (min)(_x, _y);
Scalar qp = q/p; Scalar qp = q/p;
@ -479,8 +479,8 @@ struct functor_traits<scalar_multiple_op<Scalar> >
template<typename Scalar1, typename Scalar2> template<typename Scalar1, typename Scalar2>
struct scalar_multiple2_op { struct scalar_multiple2_op {
typedef typename scalar_product_traits<Scalar1,Scalar2>::ReturnType result_type; typedef typename scalar_product_traits<Scalar1,Scalar2>::ReturnType result_type;
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 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 Scalar2& other) : m_other(other) { }
EIGEN_DEVICE_FUNC 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<typename NumTraits<Scalar2>::Nested>::type m_other; typename add_const_on_value_type<typename NumTraits<Scalar2>::Nested>::type m_other;
}; };
@ -500,8 +500,8 @@ template<typename Scalar>
struct scalar_quotient1_op { struct scalar_quotient1_op {
typedef typename packet_traits<Scalar>::type Packet; typedef typename packet_traits<Scalar>::type Packet;
// FIXME default copy constructors seems bugged with std::complex<> // 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_DEVICE_FUNC 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& other) : m_other(other) {}
EIGEN_DEVICE_FUNC 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 EIGEN_STRONG_INLINE const Packet packetOp(const Packet& a) const
{ return internal::pdiv(a, pset1<Packet>(m_other)); } { return internal::pdiv(a, pset1<Packet>(m_other)); }
@ -516,8 +516,8 @@ struct functor_traits<scalar_quotient1_op<Scalar> >
template<typename Scalar> template<typename Scalar>
struct scalar_constant_op { struct scalar_constant_op {
typedef typename packet_traits<Scalar>::type Packet; typedef typename packet_traits<Scalar>::type Packet;
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_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& other) : m_other(other) { }
template<typename Index> template<typename Index>
EIGEN_DEVICE_FUNC 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<typename Index> template<typename Index>

View File

@ -21,7 +21,7 @@ struct isApprox_selector
{ {
static bool run(const Derived& x, const OtherDerived& y, const typename Derived::RealScalar& prec) 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<Derived,2>::type nested(x); typename internal::nested<Derived,2>::type nested(x);
typename internal::nested<OtherDerived,2>::type otherNested(y); typename internal::nested<OtherDerived,2>::type otherNested(y);
return (nested - otherNested).cwiseAbs2().sum() <= prec * prec * (min)(nested.cwiseAbs2().sum(), otherNested.cwiseAbs2().sum()); return (nested - otherNested).cwiseAbs2().sum() <= prec * prec * (min)(nested.cwiseAbs2().sum(), otherNested.cwiseAbs2().sum());

View File

@ -121,12 +121,12 @@ pdiv(const Packet& a,
/** \internal \returns the min of \a a and \a b (coeff-wise) */ /** \internal \returns the min of \a a and \a b (coeff-wise) */
template<typename Packet> inline Packet template<typename Packet> inline Packet
pmin(const Packet& a, 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) */ /** \internal \returns the max of \a a and \a b (coeff-wise) */
template<typename Packet> inline Packet template<typename Packet> inline Packet
pmax(const Packet& a, 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 */ /** \internal \returns the absolute value of \a a */
template<typename Packet> inline Packet template<typename Packet> inline Packet

View File

@ -115,6 +115,7 @@ template<typename PlainObjectType, int MapOptions, typename StrideType> class Ma
inline PointerType cast_to_pointer_type(PointerArgType ptr) { return const_cast<PointerType>(ptr); } inline PointerType cast_to_pointer_type(PointerArgType ptr) { return const_cast<PointerType>(ptr); }
#else #else
typedef PointerType PointerArgType; typedef PointerType PointerArgType;
EIGEN_DEVICE_FUNC
inline PointerType cast_to_pointer_type(PointerArgType ptr) { return ptr; } inline PointerType cast_to_pointer_type(PointerArgType ptr) { return ptr; }
#endif #endif

View File

@ -203,7 +203,9 @@ template<typename Derived> class MapBase<Derived, WriteAccessors>
const Scalar const Scalar
>::type ScalarWithConstIfNotLvalue; >::type ScalarWithConstIfNotLvalue;
EIGEN_DEVICE_FUNC
inline const Scalar* data() const { return this->m_data; } 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 inline ScalarWithConstIfNotLvalue* data() { return this->m_data; } // no const-cast here so non-const-correct code will give a compile error
EIGEN_DEVICE_FUNC EIGEN_DEVICE_FUNC

View File

@ -63,6 +63,7 @@ template<typename Scalar>
struct real_impl struct real_impl
{ {
typedef typename NumTraits<Scalar>::Real RealScalar; typedef typename NumTraits<Scalar>::Real RealScalar;
EIGEN_DEVICE_FUNC
static inline RealScalar run(const Scalar& x) static inline RealScalar run(const Scalar& x)
{ {
return x; return x;
@ -72,6 +73,7 @@ struct real_impl
template<typename RealScalar> template<typename RealScalar>
struct real_impl<std::complex<RealScalar> > struct real_impl<std::complex<RealScalar> >
{ {
EIGEN_DEVICE_FUNC
static inline RealScalar run(const std::complex<RealScalar>& x) static inline RealScalar run(const std::complex<RealScalar>& x)
{ {
using std::real; using std::real;
@ -86,6 +88,7 @@ struct real_retval
}; };
template<typename Scalar> template<typename Scalar>
EIGEN_DEVICE_FUNC
inline EIGEN_MATHFUNC_RETVAL(real, Scalar) real(const Scalar& x) inline EIGEN_MATHFUNC_RETVAL(real, Scalar) real(const Scalar& x)
{ {
return EIGEN_MATHFUNC_IMPL(real, Scalar)::run(x); return EIGEN_MATHFUNC_IMPL(real, Scalar)::run(x);
@ -99,6 +102,7 @@ template<typename Scalar>
struct imag_impl struct imag_impl
{ {
typedef typename NumTraits<Scalar>::Real RealScalar; typedef typename NumTraits<Scalar>::Real RealScalar;
EIGEN_DEVICE_FUNC
static inline RealScalar run(const Scalar&) static inline RealScalar run(const Scalar&)
{ {
return RealScalar(0); return RealScalar(0);
@ -108,6 +112,7 @@ struct imag_impl
template<typename RealScalar> template<typename RealScalar>
struct imag_impl<std::complex<RealScalar> > struct imag_impl<std::complex<RealScalar> >
{ {
EIGEN_DEVICE_FUNC
static inline RealScalar run(const std::complex<RealScalar>& x) static inline RealScalar run(const std::complex<RealScalar>& x)
{ {
using std::imag; using std::imag;
@ -122,6 +127,7 @@ struct imag_retval
}; };
template<typename Scalar> template<typename Scalar>
EIGEN_DEVICE_FUNC
inline EIGEN_MATHFUNC_RETVAL(imag, Scalar) imag(const Scalar& x) inline EIGEN_MATHFUNC_RETVAL(imag, Scalar) imag(const Scalar& x)
{ {
return EIGEN_MATHFUNC_IMPL(imag, Scalar)::run(x); return EIGEN_MATHFUNC_IMPL(imag, Scalar)::run(x);
@ -135,10 +141,12 @@ template<typename Scalar>
struct real_ref_impl struct real_ref_impl
{ {
typedef typename NumTraits<Scalar>::Real RealScalar; typedef typename NumTraits<Scalar>::Real RealScalar;
EIGEN_DEVICE_FUNC
static inline RealScalar& run(Scalar& x) static inline RealScalar& run(Scalar& x)
{ {
return reinterpret_cast<RealScalar*>(&x)[0]; return reinterpret_cast<RealScalar*>(&x)[0];
} }
EIGEN_DEVICE_FUNC
static inline const RealScalar& run(const Scalar& x) static inline const RealScalar& run(const Scalar& x)
{ {
return reinterpret_cast<const RealScalar*>(&x)[0]; return reinterpret_cast<const RealScalar*>(&x)[0];
@ -152,12 +160,14 @@ struct real_ref_retval
}; };
template<typename Scalar> template<typename Scalar>
EIGEN_DEVICE_FUNC
inline typename add_const_on_value_type< EIGEN_MATHFUNC_RETVAL(real_ref, Scalar) >::type real_ref(const Scalar& x) inline typename add_const_on_value_type< EIGEN_MATHFUNC_RETVAL(real_ref, Scalar) >::type real_ref(const Scalar& x)
{ {
return real_ref_impl<Scalar>::run(x); return real_ref_impl<Scalar>::run(x);
} }
template<typename Scalar> template<typename Scalar>
EIGEN_DEVICE_FUNC
inline EIGEN_MATHFUNC_RETVAL(real_ref, Scalar) real_ref(Scalar& x) inline EIGEN_MATHFUNC_RETVAL(real_ref, Scalar) real_ref(Scalar& x)
{ {
return EIGEN_MATHFUNC_IMPL(real_ref, Scalar)::run(x); return EIGEN_MATHFUNC_IMPL(real_ref, Scalar)::run(x);
@ -171,10 +181,12 @@ template<typename Scalar, bool IsComplex>
struct imag_ref_default_impl struct imag_ref_default_impl
{ {
typedef typename NumTraits<Scalar>::Real RealScalar; typedef typename NumTraits<Scalar>::Real RealScalar;
EIGEN_DEVICE_FUNC
static inline RealScalar& run(Scalar& x) static inline RealScalar& run(Scalar& x)
{ {
return reinterpret_cast<RealScalar*>(&x)[1]; return reinterpret_cast<RealScalar*>(&x)[1];
} }
EIGEN_DEVICE_FUNC
static inline const RealScalar& run(const Scalar& x) static inline const RealScalar& run(const Scalar& x)
{ {
return reinterpret_cast<RealScalar*>(&x)[1]; return reinterpret_cast<RealScalar*>(&x)[1];
@ -184,10 +196,12 @@ struct imag_ref_default_impl
template<typename Scalar> template<typename Scalar>
struct imag_ref_default_impl<Scalar, false> struct imag_ref_default_impl<Scalar, false>
{ {
EIGEN_DEVICE_FUNC
static inline Scalar run(Scalar&) static inline Scalar run(Scalar&)
{ {
return Scalar(0); return Scalar(0);
} }
EIGEN_DEVICE_FUNC
static inline const Scalar run(const Scalar&) static inline const Scalar run(const Scalar&)
{ {
return Scalar(0); return Scalar(0);
@ -204,12 +218,14 @@ struct imag_ref_retval
}; };
template<typename Scalar> template<typename Scalar>
EIGEN_DEVICE_FUNC
inline typename add_const_on_value_type< EIGEN_MATHFUNC_RETVAL(imag_ref, Scalar) >::type imag_ref(const Scalar& x) inline typename add_const_on_value_type< EIGEN_MATHFUNC_RETVAL(imag_ref, Scalar) >::type imag_ref(const Scalar& x)
{ {
return imag_ref_impl<Scalar>::run(x); return imag_ref_impl<Scalar>::run(x);
} }
template<typename Scalar> template<typename Scalar>
EIGEN_DEVICE_FUNC
inline EIGEN_MATHFUNC_RETVAL(imag_ref, Scalar) imag_ref(Scalar& x) inline EIGEN_MATHFUNC_RETVAL(imag_ref, Scalar) imag_ref(Scalar& x)
{ {
return EIGEN_MATHFUNC_IMPL(imag_ref, Scalar)::run(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<typename Scalar> template<typename Scalar>
struct conj_impl struct conj_impl
{ {
EIGEN_DEVICE_FUNC
static inline Scalar run(const Scalar& x) static inline Scalar run(const Scalar& x)
{ {
return x; return x;
@ -231,6 +248,7 @@ struct conj_impl
template<typename RealScalar> template<typename RealScalar>
struct conj_impl<std::complex<RealScalar> > struct conj_impl<std::complex<RealScalar> >
{ {
EIGEN_DEVICE_FUNC
static inline std::complex<RealScalar> run(const std::complex<RealScalar>& x) static inline std::complex<RealScalar> run(const std::complex<RealScalar>& x)
{ {
using std::conj; using std::conj;
@ -245,6 +263,7 @@ struct conj_retval
}; };
template<typename Scalar> template<typename Scalar>
EIGEN_DEVICE_FUNC
inline EIGEN_MATHFUNC_RETVAL(conj, Scalar) conj(const Scalar& x) inline EIGEN_MATHFUNC_RETVAL(conj, Scalar) conj(const Scalar& x)
{ {
return EIGEN_MATHFUNC_IMPL(conj, Scalar)::run(x); return EIGEN_MATHFUNC_IMPL(conj, Scalar)::run(x);
@ -258,6 +277,7 @@ template<typename Scalar>
struct abs2_impl struct abs2_impl
{ {
typedef typename NumTraits<Scalar>::Real RealScalar; typedef typename NumTraits<Scalar>::Real RealScalar;
EIGEN_DEVICE_FUNC
static inline RealScalar run(const Scalar& x) static inline RealScalar run(const Scalar& x)
{ {
return x*x; return x*x;
@ -267,6 +287,7 @@ struct abs2_impl
template<typename RealScalar> template<typename RealScalar>
struct abs2_impl<std::complex<RealScalar> > struct abs2_impl<std::complex<RealScalar> >
{ {
EIGEN_DEVICE_FUNC
static inline RealScalar run(const std::complex<RealScalar>& x) static inline RealScalar run(const std::complex<RealScalar>& x)
{ {
return real(x)*real(x) + imag(x)*imag(x); return real(x)*real(x) + imag(x)*imag(x);
@ -280,6 +301,7 @@ struct abs2_retval
}; };
template<typename Scalar> template<typename Scalar>
EIGEN_DEVICE_FUNC
inline EIGEN_MATHFUNC_RETVAL(abs2, Scalar) abs2(const Scalar& x) inline EIGEN_MATHFUNC_RETVAL(abs2, Scalar) abs2(const Scalar& x)
{ {
return EIGEN_MATHFUNC_IMPL(abs2, Scalar)::run(x); return EIGEN_MATHFUNC_IMPL(abs2, Scalar)::run(x);
@ -293,6 +315,7 @@ template<typename Scalar, bool IsComplex>
struct norm1_default_impl struct norm1_default_impl
{ {
typedef typename NumTraits<Scalar>::Real RealScalar; typedef typename NumTraits<Scalar>::Real RealScalar;
EIGEN_DEVICE_FUNC
static inline RealScalar run(const Scalar& x) static inline RealScalar run(const Scalar& x)
{ {
using std::abs; using std::abs;
@ -303,6 +326,7 @@ struct norm1_default_impl
template<typename Scalar> template<typename Scalar>
struct norm1_default_impl<Scalar, false> struct norm1_default_impl<Scalar, false>
{ {
EIGEN_DEVICE_FUNC
static inline Scalar run(const Scalar& x) static inline Scalar run(const Scalar& x)
{ {
using std::abs; using std::abs;
@ -320,6 +344,7 @@ struct norm1_retval
}; };
template<typename Scalar> template<typename Scalar>
EIGEN_DEVICE_FUNC
inline EIGEN_MATHFUNC_RETVAL(norm1, Scalar) norm1(const Scalar& x) inline EIGEN_MATHFUNC_RETVAL(norm1, Scalar) norm1(const Scalar& x)
{ {
return EIGEN_MATHFUNC_IMPL(norm1, Scalar)::run(x); return EIGEN_MATHFUNC_IMPL(norm1, Scalar)::run(x);
@ -335,8 +360,8 @@ struct hypot_impl
typedef typename NumTraits<Scalar>::Real RealScalar; typedef typename NumTraits<Scalar>::Real RealScalar;
static inline RealScalar run(const Scalar& x, const Scalar& y) static inline RealScalar run(const Scalar& x, const Scalar& y)
{ {
using std::max; EIGEN_USING_STD_MATH(max);
using std::min; EIGEN_USING_STD_MATH(min);
using std::abs; using std::abs;
RealScalar _x = abs(x); RealScalar _x = abs(x);
RealScalar _y = abs(y); RealScalar _y = abs(y);
@ -631,7 +656,7 @@ struct scalar_fuzzy_default_impl<Scalar, false, false>
} }
static inline bool isApprox(const Scalar& x, const Scalar& y, const RealScalar& prec) static inline bool isApprox(const Scalar& x, const Scalar& y, const RealScalar& prec)
{ {
using std::min; EIGEN_USING_STD_MATH(min);
using std::abs; using std::abs;
return abs(x - y) <= (min)(abs(x), abs(y)) * prec; return abs(x - y) <= (min)(abs(x), abs(y)) * prec;
} }
@ -671,7 +696,7 @@ struct scalar_fuzzy_default_impl<Scalar, true, false>
} }
static inline bool isApprox(const Scalar& x, const Scalar& y, const RealScalar& prec) 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; return abs2(x - y) <= (min)(abs2(x), abs2(y)) * prec * prec;
} }
}; };

View File

@ -219,16 +219,16 @@ template<typename Derived> class MatrixBase
Scalar eigen2_dot(const MatrixBase<OtherDerived>& other) const; Scalar eigen2_dot(const MatrixBase<OtherDerived>& other) const;
#endif #endif
RealScalar squaredNorm() const; EIGEN_DEVICE_FUNC RealScalar squaredNorm() const;
RealScalar norm() const; EIGEN_DEVICE_FUNC RealScalar norm() const;
RealScalar stableNorm() const; RealScalar stableNorm() const;
RealScalar blueNorm() const; RealScalar blueNorm() const;
RealScalar hypotNorm() const; RealScalar hypotNorm() const;
const PlainObject normalized() const; EIGEN_DEVICE_FUNC const PlainObject normalized() const;
void normalize(); EIGEN_DEVICE_FUNC void normalize();
const AdjointReturnType adjoint() const; EIGEN_DEVICE_FUNC const AdjointReturnType adjoint() const;
void adjointInPlace(); EIGEN_DEVICE_FUNC void adjointInPlace();
typedef Diagonal<Derived> DiagonalReturnType; typedef Diagonal<Derived> DiagonalReturnType;
DiagonalReturnType diagonal(); DiagonalReturnType diagonal();
@ -329,15 +329,15 @@ template<typename Derived> class MatrixBase
/////////// Array module /////////// /////////// Array module ///////////
template<int p> RealScalar lpNorm() const; template<int p> EIGEN_DEVICE_FUNC RealScalar lpNorm() const;
MatrixBase<Derived>& matrix() { return *this; } EIGEN_DEVICE_FUNC MatrixBase<Derived>& matrix() { return *this; }
const MatrixBase<Derived>& matrix() const { return *this; } EIGEN_DEVICE_FUNC const MatrixBase<Derived>& matrix() const { return *this; }
/** \returns an \link Eigen::ArrayBase Array \endlink expression of this matrix /** \returns an \link Eigen::ArrayBase Array \endlink expression of this matrix
* \sa ArrayBase::matrix() */ * \sa ArrayBase::matrix() */
ArrayWrapper<Derived> array() { return derived(); } EIGEN_DEVICE_FUNC ArrayWrapper<Derived> array() { return derived(); }
const ArrayWrapper<const Derived> array() const { return derived(); } EIGEN_DEVICE_FUNC const ArrayWrapper<const Derived> array() const { return derived(); }
/////////// LU module /////////// /////////// LU module ///////////

View File

@ -52,21 +52,24 @@ template<typename BinaryOp, typename Lhs, typename Rhs> class SelfCwiseBinaryOp
typedef typename internal::packet_traits<Scalar>::type Packet; typedef typename internal::packet_traits<Scalar>::type Packet;
EIGEN_DEVICE_FUNC
inline SelfCwiseBinaryOp(Lhs& xpr, const BinaryOp& func = BinaryOp()) : m_matrix(xpr), m_functor(func) {} inline SelfCwiseBinaryOp(Lhs& xpr, const BinaryOp& func = BinaryOp()) : m_matrix(xpr), m_functor(func) {}
inline Index rows() const { return m_matrix.rows(); } EIGEN_DEVICE_FUNC inline Index rows() const { return m_matrix.rows(); }
inline Index cols() const { return m_matrix.cols(); } EIGEN_DEVICE_FUNC inline Index cols() const { return m_matrix.cols(); }
inline Index outerStride() const { return m_matrix.outerStride(); } EIGEN_DEVICE_FUNC inline Index outerStride() const { return m_matrix.outerStride(); }
inline Index innerStride() const { return m_matrix.innerStride(); } EIGEN_DEVICE_FUNC inline Index innerStride() const { return m_matrix.innerStride(); }
inline const Scalar* data() const { return m_matrix.data(); } 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 // note that this function is needed by assign to correctly align loads/stores
// TODO make Assign use .data() // TODO make Assign use .data()
EIGEN_DEVICE_FUNC
inline Scalar& coeffRef(Index row, Index col) inline Scalar& coeffRef(Index row, Index col)
{ {
EIGEN_STATIC_ASSERT_LVALUE(Lhs) EIGEN_STATIC_ASSERT_LVALUE(Lhs)
return m_matrix.const_cast_derived().coeffRef(row, col); return m_matrix.const_cast_derived().coeffRef(row, col);
} }
EIGEN_DEVICE_FUNC
inline const Scalar& coeffRef(Index row, Index col) const inline const Scalar& coeffRef(Index row, Index col) const
{ {
return m_matrix.coeffRef(row, col); return m_matrix.coeffRef(row, col);
@ -74,17 +77,20 @@ template<typename BinaryOp, typename Lhs, typename Rhs> class SelfCwiseBinaryOp
// note that this function is needed by assign to correctly align loads/stores // note that this function is needed by assign to correctly align loads/stores
// TODO make Assign use .data() // TODO make Assign use .data()
EIGEN_DEVICE_FUNC
inline Scalar& coeffRef(Index index) inline Scalar& coeffRef(Index index)
{ {
EIGEN_STATIC_ASSERT_LVALUE(Lhs) EIGEN_STATIC_ASSERT_LVALUE(Lhs)
return m_matrix.const_cast_derived().coeffRef(index); return m_matrix.const_cast_derived().coeffRef(index);
} }
EIGEN_DEVICE_FUNC
inline const Scalar& coeffRef(Index index) const inline const Scalar& coeffRef(Index index) const
{ {
return m_matrix.const_cast_derived().coeffRef(index); return m_matrix.const_cast_derived().coeffRef(index);
} }
template<typename OtherDerived> template<typename OtherDerived>
EIGEN_DEVICE_FUNC
void copyCoeff(Index row, Index col, const DenseBase<OtherDerived>& other) void copyCoeff(Index row, Index col, const DenseBase<OtherDerived>& other)
{ {
OtherDerived& _other = other.const_cast_derived(); OtherDerived& _other = other.const_cast_derived();
@ -95,6 +101,7 @@ template<typename BinaryOp, typename Lhs, typename Rhs> class SelfCwiseBinaryOp
} }
template<typename OtherDerived> template<typename OtherDerived>
EIGEN_DEVICE_FUNC
void copyCoeff(Index index, const DenseBase<OtherDerived>& other) void copyCoeff(Index index, const DenseBase<OtherDerived>& other)
{ {
OtherDerived& _other = other.const_cast_derived(); OtherDerived& _other = other.const_cast_derived();
@ -125,6 +132,7 @@ template<typename BinaryOp, typename Lhs, typename Rhs> class SelfCwiseBinaryOp
// reimplement lazyAssign to handle complex *= real // reimplement lazyAssign to handle complex *= real
// see CwiseBinaryOp ctor for details // see CwiseBinaryOp ctor for details
template<typename RhsDerived> template<typename RhsDerived>
EIGEN_DEVICE_FUNC
EIGEN_STRONG_INLINE SelfCwiseBinaryOp& lazyAssign(const DenseBase<RhsDerived>& rhs) EIGEN_STRONG_INLINE SelfCwiseBinaryOp& lazyAssign(const DenseBase<RhsDerived>& rhs)
{ {
EIGEN_STATIC_ASSERT_SAME_MATRIX_SIZE(Lhs,RhsDerived) EIGEN_STATIC_ASSERT_SAME_MATRIX_SIZE(Lhs,RhsDerived)
@ -144,17 +152,20 @@ template<typename BinaryOp, typename Lhs, typename Rhs> class SelfCwiseBinaryOp
// overloaded to honor evaluation of special matrices // overloaded to honor evaluation of special matrices
// maybe another solution would be to not use SelfCwiseBinaryOp // maybe another solution would be to not use SelfCwiseBinaryOp
// at first... // at first...
EIGEN_DEVICE_FUNC
SelfCwiseBinaryOp& operator=(const Rhs& _rhs) SelfCwiseBinaryOp& operator=(const Rhs& _rhs)
{ {
typename internal::nested<Rhs>::type rhs(_rhs); typename internal::nested<Rhs>::type rhs(_rhs);
return Base::operator=(rhs); return Base::operator=(rhs);
} }
EIGEN_DEVICE_FUNC
Lhs& expression() const Lhs& expression() const
{ {
return m_matrix; return m_matrix;
} }
EIGEN_DEVICE_FUNC
const BinaryOp& functor() const const BinaryOp& functor() const
{ {
return m_functor; return m_functor;

View File

@ -36,8 +36,8 @@ blueNorm_impl(const EigenBase<Derived>& _vec)
typedef typename Derived::RealScalar RealScalar; typedef typename Derived::RealScalar RealScalar;
typedef typename Derived::Index Index; typedef typename Derived::Index Index;
using std::pow; using std::pow;
using std::min; EIGEN_USING_STD_MATH(min);
using std::max; EIGEN_USING_STD_MATH(max);
using std::sqrt; using std::sqrt;
using std::abs; using std::abs;
const Derived& vec(_vec.derived()); const Derived& vec(_vec.derived());
@ -141,7 +141,7 @@ template<typename Derived>
inline typename NumTraits<typename internal::traits<Derived>::Scalar>::Real inline typename NumTraits<typename internal::traits<Derived>::Scalar>::Real
MatrixBase<Derived>::stableNorm() const MatrixBase<Derived>::stableNorm() const
{ {
using std::min; EIGEN_USING_STD_MATH(min);
using std::sqrt; using std::sqrt;
const Index blockSize = 4096; const Index blockSize = 4096;
RealScalar scale(0); RealScalar scale(0);

View File

@ -568,7 +568,7 @@ void EigenSolver<MatrixType>::doComputeEigenvectors()
} }
// Overflow control // 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))); Scalar t = (max)(abs(m_matT.coeff(i,n-1)),abs(m_matT.coeff(i,n)));
if ((eps * t) * t > Scalar(1)) if ((eps * t) * t > Scalar(1))
m_matT.block(i, n-1, size-i, 2) /= t; m_matT.block(i, n-1, size-i, 2) /= t;

View File

@ -159,8 +159,8 @@ template<typename QuatDerived>
AngleAxis<Scalar>& AngleAxis<Scalar>::operator=(const QuaternionBase<QuatDerived>& q) AngleAxis<Scalar>& AngleAxis<Scalar>::operator=(const QuaternionBase<QuatDerived>& q)
{ {
using std::acos; using std::acos;
using std::min; EIGEN_USING_STD_MATH(min);
using std::max; EIGEN_USING_STD_MATH(max);
using std::sqrt; using std::sqrt;
Scalar n2 = q.vec().squaredNorm(); Scalar n2 = q.vec().squaredNorm();
if (n2 < NumTraits<Scalar>::dummy_precision()*NumTraits<Scalar>::dummy_precision()) if (n2 < NumTraits<Scalar>::dummy_precision()*NumTraits<Scalar>::dummy_precision())

View File

@ -572,7 +572,7 @@ template<class Derived>
template<typename Derived1, typename Derived2> template<typename Derived1, typename Derived2>
inline Derived& QuaternionBase<Derived>::setFromTwoVectors(const MatrixBase<Derived1>& a, const MatrixBase<Derived2>& b) inline Derived& QuaternionBase<Derived>::setFromTwoVectors(const MatrixBase<Derived1>& a, const MatrixBase<Derived2>& b)
{ {
using std::max; EIGEN_USING_STD_MATH(max);
using std::sqrt; using std::sqrt;
Vector3 v0 = a.normalized(); Vector3 v0 = a.normalized();
Vector3 v1 = b.normalized(); Vector3 v1 = b.normalized();

View File

@ -766,7 +766,7 @@ JacobiSVD<MatrixType, QRPreconditioner>::compute(const MatrixType& matrix, unsig
// if this 2x2 sub-matrix is not diagonal already... // 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 // 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. // 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)), RealScalar threshold = (max)(considerAsZero, precision * (max)(abs(m_workMatrix.coeff(p,p)),
abs(m_workMatrix.coeff(q,q)))); abs(m_workMatrix.coeff(q,q))));
if((max)(abs(m_workMatrix.coeff(p,q)),abs(m_workMatrix.coeff(q,p))) > threshold) if((max)(abs(m_workMatrix.coeff(p,q)),abs(m_workMatrix.coeff(q,p))) > threshold)