ObjectOnGPU< Eigen::Matrix< _Scalar, _Rows, _Cols, _Options, _MaxRows, _MaxCols > > Class Template Reference

Collaboration diagram for ObjectOnGPU< Eigen::Matrix< _Scalar, _Rows, _Cols, _Options, _MaxRows, _MaxCols > >:

Public Types
using	EigenMatrix_t = Eigen::Matrix< _Scalar, _Rows, _Cols, _Options, _MaxRows, _MaxCols >
using	Object_t = Eigen::Map< EigenMatrix_t, Eigen::Unaligned, Eigen::Stride< Eigen::Dynamic, 1 > >
using	Scalar = _Scalar
using	RealScalar = typename Eigen::NumTraits< Scalar >::Real

Public Member Functions
template<class Derived >
__host__	ObjectOnGPU (Eigen::PlainObjectBase< Derived > const &hObject)
__host__	ObjectOnGPU (int rows=_Rows, int cols=_Cols)
	Default constructor.
__host__	ObjectOnGPU (ObjectOnGPU const &other)=delete
	Copy constructor.
__host__	ObjectOnGPU (ObjectOnGPU &&other)
	Move constructor.
__host__	~ObjectOnGPU ()
	Destructor.
__host__ ObjectOnGPU &	operator= (ObjectOnGPU const &other)=delete
	Copy assignment operator.
__host__ ObjectOnGPU &	operator= (ObjectOnGPU &&other)
	Move assignment operator.
__host__	operator Object_t * ()
__host__ Object_t *	ptr () const
__host__ Object_t const *	const_ptr () const
__host__ void	resize (int rows, int cols=_Cols)
__host__ Scalar *	data () const
__host__ Scalar const *	const_data () const
__host__ int	rows () const
__host__ int	cols () const
__host__ int	LD () const
template<class Derived >
__host__ ObjectOnGPU &	set (Eigen::PlainObjectBase< Derived > const &inMat, magma_queue_t queue)
template<class Derived >
__host__ ObjectOnGPU &	set (Eigen::PlainObjectBase< Derived > const &inMat)
__host__ Eigen::MatrixX< Scalar >	get () const

Static Public Attributes
static constexpr bool	onGPU = true
static constexpr bool	IsRowMajor = EigenMatrix_t::IsRowMajor

Private Attributes
Object_t *	m_ptr = nullptr
Object_t **	m_ptrptr = nullptr
bool	is_constructed = false
_Scalar *	m_data = nullptr
int	m_rows = 0
int	m_cols = 0
int	m_LD = 0
int	m_reserved = 0

Static Private Attributes
static constexpr int	GPU_UNIT = 32

Friends
__host__ friend ObjectOnGPU	operator* (ObjectOnGPU const &mat1, ObjectOnGPU const &mat2)
__host__ friend std::ostream &	operator<< (std::ostream &os, const ObjectOnGPU &dMat)

Member Typedef Documentation

EigenMatrix_t

template<typename _Scalar , int _Rows, int _Cols, int _Options, int _MaxRows, int _MaxCols>

using ObjectOnGPU< Eigen::Matrix< _Scalar, _Rows, _Cols, _Options, _MaxRows, _MaxCols > >::EigenMatrix_t = Eigen::Matrix<_Scalar, _Rows, _Cols, _Options, _MaxRows, _MaxCols>

Object_t

template<typename _Scalar , int _Rows, int _Cols, int _Options, int _MaxRows, int _MaxCols>

using ObjectOnGPU< Eigen::Matrix< _Scalar, _Rows, _Cols, _Options, _MaxRows, _MaxCols > >::Object_t = Eigen::Map< EigenMatrix_t, Eigen::Unaligned, Eigen::Stride<Eigen::Dynamic, 1> >

RealScalar

template<typename _Scalar , int _Rows, int _Cols, int _Options, int _MaxRows, int _MaxCols>

using ObjectOnGPU< Eigen::Matrix< _Scalar, _Rows, _Cols, _Options, _MaxRows, _MaxCols > >::RealScalar = typename Eigen::NumTraits<Scalar>::Real

Scalar

template<typename _Scalar , int _Rows, int _Cols, int _Options, int _MaxRows, int _MaxCols>

using ObjectOnGPU< Eigen::Matrix< _Scalar, _Rows, _Cols, _Options, _MaxRows, _MaxCols > >::Scalar = _Scalar

Constructor & Destructor Documentation

ObjectOnGPU() [1/4]

template<typename _Scalar , int _Rows, int _Cols, int _Options, int _MaxRows, int _MaxCols>

template<class Derived >

__host__ ObjectOnGPU< Eigen::Matrix< _Scalar, _Rows, _Cols, _Options, _MaxRows, _MaxCols > >::ObjectOnGPU ( Eigen::PlainObjectBase< Derived > const &  hObject )

inline

            : ObjectOnGPU(hObject.rows(), hObject.cols()) {
            debug_constructor_printf(1);
            this->set(hObject);
        }

ObjectOnGPU() [2/4]

template<typename _Scalar , int _Rows, int _Cols, int _Options, int _MaxRows, int _MaxCols>

__host__ ObjectOnGPU< Eigen::Matrix< _Scalar, _Rows, _Cols, _Options, _MaxRows, _MaxCols > >::ObjectOnGPU ( int  rows = _Rows, int  cols = _Cols  )

inline

Default constructor.

Parameters

rows
cols

                                                                 {
            debug_constructor_printf((Default));
            cuCHECK(cudaMalloc((void***)&m_ptrptr, sizeof(Object_t*)));
            this->resize(rows, cols);
        }

ObjectOnGPU() [3/4]

template<typename _Scalar , int _Rows, int _Cols, int _Options, int _MaxRows, int _MaxCols>

__host__ ObjectOnGPU< Eigen::Matrix< _Scalar, _Rows, _Cols, _Options, _MaxRows, _MaxCols > >::ObjectOnGPU ( ObjectOnGPU< Eigen::Matrix< _Scalar, _Rows, _Cols, _Options, _MaxRows, _MaxCols > > const &  other )

delete

Copy constructor.

Parameters

other

ObjectOnGPU() [4/4]

template<typename _Scalar , int _Rows, int _Cols, int _Options, int _MaxRows, int _MaxCols>

__host__ ObjectOnGPU< Eigen::Matrix< _Scalar, _Rows, _Cols, _Options, _MaxRows, _MaxCols > >::ObjectOnGPU ( ObjectOnGPU< Eigen::Matrix< _Scalar, _Rows, _Cols, _Options, _MaxRows, _MaxCols > > &&  other )

inline

Move constructor.

Parameters

other

            : m_rows{other.m_rows},
              m_cols{other.m_cols},
              m_LD{other.m_LD},
              m_reserved{other.m_reserved},
              m_ptr{other.m_ptr},
              m_ptrptr{other.m_ptrptr},
              m_data{other.m_data},
              is_constructed{other.is_constructed} {
            other.m_rows         = 0;
            other.m_cols         = 0;
            other.m_LD           = 0;
            other.m_reserved     = 0;
            other.m_ptr          = nullptr;
            other.m_ptrptr       = nullptr;
            other.m_data         = nullptr;
            other.is_constructed = false;
        }

~ObjectOnGPU()

template<typename _Scalar , int _Rows, int _Cols, int _Options, int _MaxRows, int _MaxCols>

__host__ ObjectOnGPU< Eigen::Matrix< _Scalar, _Rows, _Cols, _Options, _MaxRows, _MaxCols > >::~ObjectOnGPU ( )

inline

Destructor.

                                {
            debug_destructor_printf();
            debug_print("\tm_data = " << m_data << ", m_ptrptr=" << m_ptrptr
                                      << ", m_ptr=" << m_ptr);
            if(m_ptrptr != nullptr && is_constructed) {
                debug_print("\tCalling destructObject_kernel.");
                destructObject_kernel<<<1, 1>>>(m_ptrptr);
                cuCHECK(cudaPeekAtLastError());
                cuCHECK(cudaDeviceSynchronize());
                is_constructed = false;
            }
            cuCHECK(cudaFree(m_ptrptr));
            m_ptrptr = nullptr;
 
            CHECK(magma_free(m_data));
            m_data = nullptr;
        }

Member Function Documentation

cols()

template<typename _Scalar , int _Rows, int _Cols, int _Options, int _MaxRows, int _MaxCols>

__host__ int ObjectOnGPU< Eigen::Matrix< _Scalar, _Rows, _Cols, _Options, _MaxRows, _MaxCols > >::cols ( ) const

inline

{ return m_cols; }

const_data()

template<typename _Scalar , int _Rows, int _Cols, int _Options, int _MaxRows, int _MaxCols>

__host__ Scalar const * ObjectOnGPU< Eigen::Matrix< _Scalar, _Rows, _Cols, _Options, _MaxRows, _MaxCols > >::const_data ( ) const

inline

{ return m_data; }

const_ptr()

template<typename _Scalar , int _Rows, int _Cols, int _Options, int _MaxRows, int _MaxCols>

__host__ Object_t const * ObjectOnGPU< Eigen::Matrix< _Scalar, _Rows, _Cols, _Options, _MaxRows, _MaxCols > >::const_ptr ( ) const

inline

{ return m_ptr; }

data()

template<typename _Scalar , int _Rows, int _Cols, int _Options, int _MaxRows, int _MaxCols>

__host__ Scalar * ObjectOnGPU< Eigen::Matrix< _Scalar, _Rows, _Cols, _Options, _MaxRows, _MaxCols > >::data ( ) const

inline

{ return m_data; }

get()

template<typename _Scalar , int _Rows, int _Cols, int _Options, int _MaxRows, int _MaxCols>

__host__ Eigen::MatrixX< Scalar > ObjectOnGPU< Eigen::Matrix< _Scalar, _Rows, _Cols, _Options, _MaxRows, _MaxCols > >::get ( ) const

inline

                                                  {
            debug_print(__PRETTY_FUNCTION__);
            Eigen::MatrixX<Scalar> hMat(this->rows(), this->cols());
            magma_queue_t          queue = NULL;
            magma_int_t            dev   = 0;
            magma_getdevice(&dev);
            magma_queue_create(dev, &queue);
            magma_getmatrix(this->rows(), this->cols(), sizeof(Scalar), this->data(), this->LD(),
                            hMat.data(), hMat.rows(), queue);
            magma_queue_destroy(queue);
            return hMat;
        }

LD()

template<typename _Scalar , int _Rows, int _Cols, int _Options, int _MaxRows, int _MaxCols>

__host__ int ObjectOnGPU< Eigen::Matrix< _Scalar, _Rows, _Cols, _Options, _MaxRows, _MaxCols > >::LD ( ) const

inline

{ return m_LD; }

operator Object_t *()

template<typename _Scalar , int _Rows, int _Cols, int _Options, int _MaxRows, int _MaxCols>

__host__ ObjectOnGPU< Eigen::Matrix< _Scalar, _Rows, _Cols, _Options, _MaxRows, _MaxCols > >::operator Object_t * ( )

inline

{ return m_ptr; }

operator=() [1/2]

template<typename _Scalar , int _Rows, int _Cols, int _Options, int _MaxRows, int _MaxCols>

__host__ ObjectOnGPU & ObjectOnGPU< Eigen::Matrix< _Scalar, _Rows, _Cols, _Options, _MaxRows, _MaxCols > >::operator= ( ObjectOnGPU< Eigen::Matrix< _Scalar, _Rows, _Cols, _Options, _MaxRows, _MaxCols > > &&  other )

inline

Move assignment operator.

Parameters

other

                                                             {
            if(this == &other) {
                debug_printf("\tthis(%p) == &other(%p)\n", this, &other);
                return *this;
            }
            m_rows         = other.m_rows;
            m_cols         = other.m_cols;
            m_LD           = other.m_LD;
            m_reserved     = other.m_reserved;
            m_ptr          = other.m_ptr;
            m_ptrptr       = other.m_ptrptr;
            m_data         = other.m_data;
            is_constructed = other.is_constructed;
 
            other.m_rows         = 0;
            other.m_cols         = 0;
            other.m_LD           = 0;
            other.m_reserved     = 0;
            other.m_ptr          = nullptr;
            other.m_ptrptr       = nullptr;
            other.m_data         = nullptr;
            other.is_constructed = false;
 
            return *this;
        };

operator=() [2/2]

template<typename _Scalar , int _Rows, int _Cols, int _Options, int _MaxRows, int _MaxCols>

__host__ ObjectOnGPU & ObjectOnGPU< Eigen::Matrix< _Scalar, _Rows, _Cols, _Options, _MaxRows, _MaxCols > >::operator= ( ObjectOnGPU< Eigen::Matrix< _Scalar, _Rows, _Cols, _Options, _MaxRows, _MaxCols > > const &  other )

delete

Copy assignment operator.

Parameters

other

ptr()

template<typename _Scalar , int _Rows, int _Cols, int _Options, int _MaxRows, int _MaxCols>

__host__ Object_t * ObjectOnGPU< Eigen::Matrix< _Scalar, _Rows, _Cols, _Options, _MaxRows, _MaxCols > >::ptr ( ) const

inline

{ return m_ptr; }

resize()

template<typename _Scalar , int _Rows, int _Cols, int _Options, int _MaxRows, int _MaxCols>

__host__ void ObjectOnGPU< Eigen::Matrix< _Scalar, _Rows, _Cols, _Options, _MaxRows, _MaxCols > >::resize ( int  rows, int  cols = _Cols  )

inline

                                                         {
            debug_print(__PRETTY_FUNCTION__);
            m_rows = rows;
            m_cols = cols;
            int newReserved;
            if constexpr(IsRowMajor) {
                m_LD        = magma_roundup(m_cols, GPU_UNIT);
                newReserved = m_LD * m_rows;
            }
            else {
                m_LD        = magma_roundup(m_rows, GPU_UNIT);
                newReserved = m_LD * m_cols;
            }
 
            if(m_reserved < newReserved) {
                if(is_constructed) {
                    destructObject_kernel<<<1, 1>>>(m_ptrptr);
                    cuCHECK(cudaPeekAtLastError());
                    cuCHECK(cudaDeviceSynchronize());
                }
 
                is_constructed = false;
                CHECK(magma_free(m_data));
                m_data = nullptr;
 
                m_reserved = newReserved;
                CHECK(magma_malloc((void**)&m_data, m_reserved * sizeof(Scalar)));
                // Only pass by value to the constructor is supported
                constructObject_kernel<<<1, 1>>>(m_ptrptr, m_data, m_rows, m_cols,
                                                 Eigen::Stride<Eigen::Dynamic, 1>(m_LD, 1));
                cuCHECK(cudaPeekAtLastError());
                cuCHECK(cudaMemcpy(&m_ptr, m_ptrptr, sizeof(Object_t*), cudaMemcpyDeviceToHost));
                is_constructed = true;
            }
            return;
        }

rows()

template<typename _Scalar , int _Rows, int _Cols, int _Options, int _MaxRows, int _MaxCols>

__host__ int ObjectOnGPU< Eigen::Matrix< _Scalar, _Rows, _Cols, _Options, _MaxRows, _MaxCols > >::rows ( ) const

inline

{ return m_rows; }

set() [1/2]

template<typename _Scalar , int _Rows, int _Cols, int _Options, int _MaxRows, int _MaxCols>

template<class Derived >

__host__ ObjectOnGPU & ObjectOnGPU< Eigen::Matrix< _Scalar, _Rows, _Cols, _Options, _MaxRows, _MaxCols > >::set ( Eigen::PlainObjectBase< Derived > const &  inMat )

inline

                                                                              {
            debug_print(__PRETTY_FUNCTION__);
            magma_queue_t queue = NULL;
            magma_int_t   dev   = 0;
            magma_getdevice(&dev);
            magma_queue_create(dev, &queue);
            this->set(inMat, queue);
            magma_queue_destroy(queue);
            return *this;
        }

set() [2/2]

template<typename _Scalar , int _Rows, int _Cols, int _Options, int _MaxRows, int _MaxCols>

template<class Derived >

__host__ ObjectOnGPU & ObjectOnGPU< Eigen::Matrix< _Scalar, _Rows, _Cols, _Options, _MaxRows, _MaxCols > >::set ( Eigen::PlainObjectBase< Derived > const &  inMat, magma_queue_t  queue  )

inline

                                                                                {
            debug_print(__PRETTY_FUNCTION__);
            this->resize(inMat.rows(), inMat.cols());
            if constexpr(Derived::IsRowMajor == IsRowMajor
                         && std::is_same_v<Scalar, typename Derived::Scalar>) {
                debug_print("\tcase1\n");
                magma_setmatrix(inMat.innerSize(), inMat.outerSize(), sizeof(Scalar), inMat.data(),
                                inMat.innerSize(), m_data, m_LD, queue);
            }
            else {
                // static_assert(std::is_same_v<Scalar, Complex_t<RealScalar> >);
                debug_print("\tcase2\n");
                EigenMatrix_t inMatCopy(inMat.template cast<_Scalar>());
                magma_setmatrix(inMat.innerSize(), inMat.outerSize(), sizeof(_Scalar),
                                inMatCopy.data(), inMat.innerSize(), m_data, m_LD, queue);
            }
            return *this;
        }

Friends And Related Symbol Documentation

operator*

template<typename _Scalar , int _Rows, int _Cols, int _Options, int _MaxRows, int _MaxCols>

__host__ friend ObjectOnGPU operator* ( ObjectOnGPU< Eigen::Matrix< _Scalar, _Rows, _Cols, _Options, _MaxRows, _MaxCols > > const &  mat1, ObjectOnGPU< Eigen::Matrix< _Scalar, _Rows, _Cols, _Options, _MaxRows, _MaxCols > > const &  mat2  )

friend

                                                                                                {
            constexpr Scalar alpha = {1, 0};
            constexpr Scalar beta  = {0, 0};
            ObjectOnGPU      res(mat1.rows(), mat2.cols());
 
            magma_queue_t queue = NULL;
            magma_int_t   dev   = 0;
            magma_getdevice(&dev);
            magma_queue_create(dev, &queue);
            magma_hemm_wrapper(MagmaLeft, MagmaLower, res.rows(), res.cols(), alpha, mat1.data(),
                               mat1.LD(), mat2.data(), mat2.LD(), beta, res.data(), res.LD(),
                               queue);
            // magma_queue_sync(queue);
            magma_queue_destroy(queue);
 
            return res;
        }

operator<<

template<typename _Scalar , int _Rows, int _Cols, int _Options, int _MaxRows, int _MaxCols>

__host__ friend std::ostream & operator<< ( std::ostream &  os, const ObjectOnGPU< Eigen::Matrix< _Scalar, _Rows, _Cols, _Options, _MaxRows, _MaxCols > > &  dMat  )

friend

                                                                                        {
            debug_print(__PRETTY_FUNCTION__);
            std::cout << dMat.get() << std::endl;
            return os;
        };

Member Data Documentation

GPU_UNIT

template<typename _Scalar , int _Rows, int _Cols, int _Options, int _MaxRows, int _MaxCols>

constexpr int ObjectOnGPU< Eigen::Matrix< _Scalar, _Rows, _Cols, _Options, _MaxRows, _MaxCols > >::GPU_UNIT = 32

staticconstexprprivate

is_constructed

template<typename _Scalar , int _Rows, int _Cols, int _Options, int _MaxRows, int _MaxCols>

bool ObjectOnGPU< Eigen::Matrix< _Scalar, _Rows, _Cols, _Options, _MaxRows, _MaxCols > >::is_constructed = false

private

IsRowMajor

template<typename _Scalar , int _Rows, int _Cols, int _Options, int _MaxRows, int _MaxCols>

constexpr bool ObjectOnGPU< Eigen::Matrix< _Scalar, _Rows, _Cols, _Options, _MaxRows, _MaxCols > >::IsRowMajor = EigenMatrix_t::IsRowMajor

staticconstexpr

m_cols

template<typename _Scalar , int _Rows, int _Cols, int _Options, int _MaxRows, int _MaxCols>

int ObjectOnGPU< Eigen::Matrix< _Scalar, _Rows, _Cols, _Options, _MaxRows, _MaxCols > >::m_cols = 0

private

m_data

template<typename _Scalar , int _Rows, int _Cols, int _Options, int _MaxRows, int _MaxCols>

_Scalar* ObjectOnGPU< Eigen::Matrix< _Scalar, _Rows, _Cols, _Options, _MaxRows, _MaxCols > >::m_data = nullptr

private

m_LD

template<typename _Scalar , int _Rows, int _Cols, int _Options, int _MaxRows, int _MaxCols>

int ObjectOnGPU< Eigen::Matrix< _Scalar, _Rows, _Cols, _Options, _MaxRows, _MaxCols > >::m_LD = 0

private

m_ptr

template<typename _Scalar , int _Rows, int _Cols, int _Options, int _MaxRows, int _MaxCols>

Object_t* ObjectOnGPU< Eigen::Matrix< _Scalar, _Rows, _Cols, _Options, _MaxRows, _MaxCols > >::m_ptr = nullptr

private

m_ptrptr

template<typename _Scalar , int _Rows, int _Cols, int _Options, int _MaxRows, int _MaxCols>

Object_t** ObjectOnGPU< Eigen::Matrix< _Scalar, _Rows, _Cols, _Options, _MaxRows, _MaxCols > >::m_ptrptr = nullptr

private

m_reserved

template<typename _Scalar , int _Rows, int _Cols, int _Options, int _MaxRows, int _MaxCols>

int ObjectOnGPU< Eigen::Matrix< _Scalar, _Rows, _Cols, _Options, _MaxRows, _MaxCols > >::m_reserved = 0

private

m_rows

template<typename _Scalar , int _Rows, int _Cols, int _Options, int _MaxRows, int _MaxCols>

int ObjectOnGPU< Eigen::Matrix< _Scalar, _Rows, _Cols, _Options, _MaxRows, _MaxCols > >::m_rows = 0

private

onGPU

template<typename _Scalar , int _Rows, int _Cols, int _Options, int _MaxRows, int _MaxCols>

constexpr bool ObjectOnGPU< Eigen::Matrix< _Scalar, _Rows, _Cols, _Options, _MaxRows, _MaxCols > >::onGPU = true

staticconstexpr

---

The documentation for this class was generated from the following file:

• /Users/shoki/GitHub/Locality/Headers/MatrixUtils/MatrixUtils.cuh