mirror of
https://gitlab.com/libeigen/eigen.git
synced 2025-03-19 18:40:38 +08:00
Fixed some compiler bugs in NVCC, now compiles with CUDA.
(chtz: Manually joined sevaral commits to keep the history clean)
This commit is contained in:
Jonas Adler
parent
d259b719d1
commit
815fa0dbf6
@ -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__
|
||||
|
||||
@ -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"
|
||||
|
||||
@ -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)
|
||||
{
|
||||
|
||||
@ -114,27 +114,12 @@ 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
|
||||
*
|
||||
|
||||
@ -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.
|
||||
*
|
||||
|
||||
@ -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
|
||||
*
|
||||
|
||||
@ -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
|
||||
|
||||
@ -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); }
|
||||
};
|
||||
|
||||
|
||||
@ -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 {
|
||||
|
||||
|
||||
@ -1,6 +1,11 @@
|
||||
#include <Eigen/Eigen>
|
||||
#include <iostream>
|
||||
|
||||
#ifndef M_PI
|
||||
#define M_PI 3.1415926535897932384626433832795
|
||||
#endif
|
||||
|
||||
|
||||
using namespace Eigen;
|
||||
using namespace std;
|
||||
|
||||
|
||||
Loading…
x
Reference in New Issue
Block a user