ObjectOnGPU< SparseCompressed< Scalar_t, Options_ > > Class Template Reference

Inheritance diagram for ObjectOnGPU< SparseCompressed< Scalar_t, Options_ > >:

Collaboration diagram for ObjectOnGPU< SparseCompressed< Scalar_t, Options_ > >:

Public Types
using	Scalar = Scalar_t
typedef Eigen::NumTraits< Scalar_t >::Real	RealScalar

Public Member Functions
__host__	ObjectOnGPU (int rows=0, int cols=0)
template<typename Scalar2_t >
__host__	ObjectOnGPU (SparseCompressed< Scalar2_t, Options_ > const &hObject)
template<class Scalar_ >
__host__	ObjectOnGPU (Eigen::SparseMatrix< Scalar_ > const &hObject)
template<class Derived >
__host__	ObjectOnGPU (Eigen::SparseMatrixBase< Derived > const &hObject)
__host__ ObjectOnGPU &	resize (int rows, int cols, int reserved=1)
__host__ int	rows () const
__host__ int	cols () const
__host__ int	nonZeros () const
template<class Scalar2_t >
__host__ void	set (SparseCompressed< Scalar2_t, Options_ > hObject)
template<class Scalar2_t >
__host__ void	set (Eigen::SparseMatrix< Scalar2_t > hObject)
template<typename Derived >
__host__ void	set (Eigen::SparseMatrixBase< Derived > const &hObject_)
Public Member Functions inherited from ObjectOnGPU_Base< SparseCompressed< Scalar_t, Options_ > >
	ObjectOnGPU_Base (Args... args)
	Default constructor for ObjectOnGPU_Base object.
	ObjectOnGPU_Base (ObjectOnGPU_Base const &other)=delete
	Default constructor.
	ObjectOnGPU_Base (ObjectOnGPU_Base &&other)
	Move constructor for ObjectOnGPU_Base object.
	~ObjectOnGPU_Base ()
	Destructor.
ObjectOnGPU_Base &	operator= (ObjectOnGPU_Base const &other)=delete
	Copy assignment operator (yet to be implmented)
ObjectOnGPU_Base &	operator= (ObjectOnGPU_Base &&other)
	Move assignment operator.
	operator Object_t * () const
Object_t *	ptr () const
Object_t const *	const_ptr () const

Static Public Attributes
static constexpr bool	IsRowMajor = bool(Options_ == Eigen::RowMajor)
Static Public Attributes inherited from ObjectOnGPU_Base< SparseCompressed< Scalar_t, Options_ > >
static constexpr bool	on_GPU

Private Types
using	Object_t = SparseCompressed< Scalar_t, Options_ >

Private Member Functions
template<class SparseCompressed_ >
__host__ void	set_impl (SparseCompressed_ &hObject)

Private Attributes
bool	is_constructed = false

Friends
Operator overloads
__host__ friend std::ostream &	operator<< (std::ostream &os, ObjectOnGPU const &sparseMatOnGPU)

Additional Inherited Members
Protected Types inherited from ObjectOnGPU_Base< SparseCompressed< Scalar_t, Options_ > >
using	Object_t = SparseCompressed< Scalar_t, Options_ >
Protected Attributes inherited from ObjectOnGPU_Base< SparseCompressed< Scalar_t, Options_ > >
SparseCompressed< Scalar_t, Options_ > *	m_ptr
SparseCompressed< Scalar_t, Options_ > **	m_ptrptr

Member Typedef Documentation

Object_t

template<typename Scalar_t , Eigen::StorageOptions Options_>

using ObjectOnGPU< SparseCompressed< Scalar_t, Options_ > >::Object_t = SparseCompressed<Scalar_t, Options_>

private

RealScalar

template<typename Scalar_t , Eigen::StorageOptions Options_>

typedef Eigen::NumTraits<Scalar_t>::Real ObjectOnGPU< SparseCompressed< Scalar_t, Options_ > >::RealScalar

Scalar

template<typename Scalar_t , Eigen::StorageOptions Options_>

using ObjectOnGPU< SparseCompressed< Scalar_t, Options_ > >::Scalar = Scalar_t

Constructor & Destructor Documentation

ObjectOnGPU() [1/4]

template<typename Scalar_t , Eigen::StorageOptions Options_>

__host__ ObjectOnGPU< SparseCompressed< Scalar_t, Options_ > >::ObjectOnGPU ( int  rows = 0, int  cols = 0  )

inline

: ObjectOnGPU_Base<Object_t>(rows, cols) {}

ObjectOnGPU() [2/4]

template<typename Scalar_t , Eigen::StorageOptions Options_>

template<typename Scalar2_t >

__host__ ObjectOnGPU< SparseCompressed< Scalar_t, Options_ > >::ObjectOnGPU ( SparseCompressed< Scalar2_t, Options_ > const &  hObject )

inline

            : ObjectOnGPU(0, 0) {
            debug_constructor_printf(1);
            this->set(hObject);
        }

ObjectOnGPU() [3/4]

template<typename Scalar_t , Eigen::StorageOptions Options_>

template<class Scalar_ >

__host__ ObjectOnGPU< SparseCompressed< Scalar_t, Options_ > >::ObjectOnGPU ( Eigen::SparseMatrix< Scalar_ > const &  hObject )

inline

                                                                        : ObjectOnGPU(0, 0) {
            debug_constructor_printf(2);
            const_cast<Eigen::SparseMatrix<Scalar_>*>(&hObject)->makeCompressed();
            this->set(hObject);
        }

ObjectOnGPU() [4/4]

template<typename Scalar_t , Eigen::StorageOptions Options_>

template<class Derived >

__host__ ObjectOnGPU< SparseCompressed< Scalar_t, Options_ > >::ObjectOnGPU ( Eigen::SparseMatrixBase< Derived > const &  hObject )

inline

                                                                            : ObjectOnGPU(0, 0) {
            debug_constructor_printf(3);
            this->set(hObject);
        }

Member Function Documentation

cols()

template<typename Scalar_t , Eigen::StorageOptions Options_>

__host__ int ObjectOnGPU< SparseCompressed< Scalar_t, Options_ > >::cols ( ) const

inline

                                  {
            int* dRowsPtr = nullptr;
            cuCHECK(cudaMalloc(&dRowsPtr, sizeof(int)));
            SparseCompressed_cols_kernel<<<1, 1>>>(this->m_ptr, dRowsPtr);
            cuCHECK(cudaPeekAtLastError());
            cuCHECK(cudaDeviceSynchronize());
            int res;
            cuCHECK(cudaMemcpy(&res, dRowsPtr, sizeof(int), cudaMemcpyDeviceToHost));
            cuCHECK(cudaFree(dRowsPtr));
            return res;
        }

nonZeros()

template<typename Scalar_t , Eigen::StorageOptions Options_>

__host__ int ObjectOnGPU< SparseCompressed< Scalar_t, Options_ > >::nonZeros ( ) const

inline

                                      {
            int* dRowsPtr = nullptr;
            cuCHECK(cudaMalloc(&dRowsPtr, sizeof(int)));
            SparseCompressed_nonZeros_kernel<<<1, 1>>>(this->m_ptr, dRowsPtr);
            cuCHECK(cudaPeekAtLastError());
            cuCHECK(cudaDeviceSynchronize());
            int res;
            cuCHECK(cudaMemcpy(&res, dRowsPtr, sizeof(int), cudaMemcpyDeviceToHost));
            cuCHECK(cudaFree(dRowsPtr));
            return res;
        }

resize()

template<typename Scalar_t , Eigen::StorageOptions Options_>

__host__ ObjectOnGPU & ObjectOnGPU< SparseCompressed< Scalar_t, Options_ > >::resize ( int  rows, int  cols, int  reserved = 1  )

inline

                                                                           {
            SparseCompressed_resize_kernel<<<1, 1>>>(this->m_ptr, rows, cols, reserved);
            cuCHECK(cudaPeekAtLastError());
            cuCHECK(cudaDeviceSynchronize());
            return *this;
        }

rows()

template<typename Scalar_t , Eigen::StorageOptions Options_>

__host__ int ObjectOnGPU< SparseCompressed< Scalar_t, Options_ > >::rows ( ) const

inline

                                  {
            int* dRowsPtr = nullptr;
            cuCHECK(cudaMalloc(&dRowsPtr, sizeof(int)));
            SparseCompressed_rows_kernel<<<1, 1>>>(this->m_ptr, dRowsPtr);
            cuCHECK(cudaPeekAtLastError());
            cuCHECK(cudaDeviceSynchronize());
            int res;
            cuCHECK(cudaMemcpy(&res, dRowsPtr, sizeof(int), cudaMemcpyDeviceToHost));
            cuCHECK(cudaFree(dRowsPtr));
            return res;
        }

set() [1/3]

template<typename Scalar_t , Eigen::StorageOptions Options_>

template<class Scalar2_t >

__host__ void ObjectOnGPU< SparseCompressed< Scalar_t, Options_ > >::set ( Eigen::SparseMatrix< Scalar2_t >  hObject )

inline

                                                                {
            debug_print(__PRETTY_FUNCTION__ << "\n");
            hObject.makeCompressed();
            this->set_impl(hObject);
        }

set() [2/3]

template<typename Scalar_t , Eigen::StorageOptions Options_>

template<typename Derived >

__host__ void ObjectOnGPU< SparseCompressed< Scalar_t, Options_ > >::set ( Eigen::SparseMatrixBase< Derived > const &  hObject_ )

inline

                                                                          {
            debug_print(__PRETTY_FUNCTION__ << "\n");
            Eigen::SparseMatrix<Scalar> hObject(hObject_.template cast<Scalar>());
            this->set(hObject);
        }

set() [3/3]

template<typename Scalar_t , Eigen::StorageOptions Options_>

template<class Scalar2_t >

__host__ void ObjectOnGPU< SparseCompressed< Scalar_t, Options_ > >::set ( SparseCompressed< Scalar2_t, Options_ >  hObject )

inline

                                                                         {
            debug_print(__PRETTY_FUNCTION__ << hObject.rows() << ", " << hObject.cols() << ", "
                                            << hObject.nonZeros() << "\n");
            this->set_impl(hObject);
        }

set_impl()

template<typename Scalar_t , Eigen::StorageOptions Options_>

template<class SparseCompressed_ >

__host__ void ObjectOnGPU< SparseCompressed< Scalar_t, Options_ > >::set_impl ( SparseCompressed_ &  hObject )

private

                                {
    using Scalar2_t = typename SparseCompressed_::Scalar;
    debug_print(__PRETTY_FUNCTION__ << hObject.rows() << ", " << hObject.cols() << ", "
                                    << hObject.nonZeros() << "\n");
    this->resize(hObject.rows(), hObject.cols(), hObject.nonZeros());
    SparseCompressed_setNonZeros_kernel<<<1, 1>>>(this->m_ptr, hObject.nonZeros());
    cuCHECK(cudaPeekAtLastError());
    cuCHECK(cudaDeviceSynchronize());
 
    int*    dSrcOuterIndex;
    int*    dSrcInnerIndex;
    Scalar* dSrcValueIndex;
    cuCHECK(cudaMalloc((void**)&dSrcOuterIndex, sizeof(int) * (hObject.outerSize() + 1)));
    cuCHECK(cudaMalloc((void**)&dSrcInnerIndex, sizeof(int) * hObject.nonZeros()));
    cuCHECK(cudaMalloc((void**)&dSrcValueIndex, sizeof(Scalar) * hObject.nonZeros()));
 
    cuCHECK(cudaMemcpy(dSrcOuterIndex, hObject.outerIndexPtr(),
                       sizeof(int) * (hObject.outerSize() + 1), cudaMemcpyHostToDevice));
    cuCHECK(cudaMemcpy(dSrcInnerIndex, hObject.innerIndexPtr(), sizeof(int) * hObject.nonZeros(),
                       cudaMemcpyHostToDevice));
    if constexpr(std::is_same_v<Scalar, Scalar2_t>) {
        cuCHECK(cudaMemcpy(dSrcValueIndex, hObject.valuePtr(), sizeof(Scalar) * hObject.nonZeros(),
                           cudaMemcpyHostToDevice));
    }
    else {
        std::vector<Scalar> hSrcValueIndex(hObject.nonZeros());
        for(int j = 0; j != hSrcValueIndex.size(); ++j) hSrcValueIndex[j] = hObject.valuePtr()[j];
        cuCHECK(cudaMemcpy(dSrcValueIndex, hSrcValueIndex.data(),
                           sizeof(Scalar) * hObject.nonZeros(), cudaMemcpyHostToDevice));
    }
 
    void (*kernel)(Object_t*, int*, int*, Scalar_t*) = &SparseCompressed_set_kernel;
    struct cudaFuncAttributes attr;
    cuCHECK(cudaFuncGetAttributes(&attr, kernel));
    int const numElem = (hObject.outerSize() + 1) + hObject.nonZeros();
    int const nThread = min(numElem, attr.maxThreadsPerBlock);
    int const nBlock  = (numElem % nThread == 0 ? numElem / nThread : numElem / nThread + 1);
 
    debug_print("\tnumElem = " << numElem << ", nThread = " << nThread << ", nBlock = " << nBlock);
    SparseCompressed_set_kernel<<<nBlock, nThread>>>(this->m_ptr, dSrcOuterIndex, dSrcInnerIndex,
                                                     dSrcValueIndex);
    cuCHECK(cudaPeekAtLastError());
 
    cuCHECK(cudaFree(dSrcOuterIndex));
    cuCHECK(cudaFree(dSrcInnerIndex));
    cuCHECK(cudaFree(dSrcValueIndex));
}

Friends And Related Symbol Documentation

operator<<

template<typename Scalar_t , Eigen::StorageOptions Options_>

__host__ friend std::ostream & operator<< ( std::ostream &  os, ObjectOnGPU< SparseCompressed< Scalar_t, Options_ > > const &  sparseMatOnGPU  )

friend

                                                                                    {
            SparseCompressed sparseMat;
            sparseMat.resize(sparseMatOnGPU.rows(), sparseMatOnGPU.cols(),
                             sparseMatOnGPU.nonZeros());
            sparseMat.setNonZeros(sparseMatOnGPU.nonZeros());
            int*    srcOuterIndexPtr;
            int*    srcInnerIndexPtr;
            Scalar* srcValuePtr;
            cuCHECK(cudaMalloc(&srcOuterIndexPtr, (sparseMat.outerSize() + 1) * sizeof(int)));
            cuCHECK(cudaMalloc(&srcInnerIndexPtr, sparseMat.nonZeros() * sizeof(int)));
            cuCHECK(cudaMalloc(&srcValuePtr, sparseMat.nonZeros() * sizeof(Scalar_t)));
 
            SparseCompressed_copyData_kernel<<<1, 1>>>(sparseMatOnGPU.m_ptr, srcOuterIndexPtr,
                                                       srcInnerIndexPtr, srcValuePtr);
            cuCHECK(cudaPeekAtLastError());
 
            cuCHECK(cudaMemcpy(sparseMat.outerIndexPtr(), srcOuterIndexPtr,
                               (sparseMat.outerSize() + 1) * sizeof(int), cudaMemcpyDeviceToHost));
            cuCHECK(cudaMemcpy(sparseMat.innerIndexPtr(), srcInnerIndexPtr,
                               sparseMat.nonZeros() * sizeof(int), cudaMemcpyDeviceToHost));
            cuCHECK(cudaMemcpy(sparseMat.valuePtr(), srcValuePtr,
                               sparseMat.nonZeros() * sizeof(Scalar), cudaMemcpyDeviceToHost));
 
            cuCHECK(cudaFree(srcOuterIndexPtr));
            cuCHECK(cudaFree(srcInnerIndexPtr));
            cuCHECK(cudaFree(srcValuePtr));
 
            os << sparseMat;
            return os;
        }

Member Data Documentation

is_constructed

template<typename Scalar_t , Eigen::StorageOptions Options_>

bool ObjectOnGPU< SparseCompressed< Scalar_t, Options_ > >::is_constructed = false

private

IsRowMajor

template<typename Scalar_t , Eigen::StorageOptions Options_>

constexpr bool ObjectOnGPU< SparseCompressed< Scalar_t, Options_ > >::IsRowMajor = bool(Options_ == Eigen::RowMajor)

staticconstexpr

---

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

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