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