Container for CSC matrices. More...

#include <cu_csc_matrix.h>

Inheritance diagram for cuCSCMatrix< DataType >:

[legend]

Collaboration diagram for cuCSCMatrix< DataType >:

[legend]

Public Member Functions
	cuCSCMatrix ()
	Default constructor.

	cuCSCMatrix (const DataType A_data_, const LongIndexType A_indices_, const LongIndexType *A_index_pointer_, const LongIndexType num_rows_, const LongIndexType num_columns_, const FlagType A_is_symmetric_, const int num_gpu_devices_)
	Constructor.

virtual	~cuCSCMatrix ()
	Destructor.

virtual FlagType	is_identity_matrix () const
	Checks whether the matrix is identity.

LongIndexType	get_nnz () const
	Returns the number of non-zero elements of the sparse matrix.

virtual void	dot (const DataType device_vector, DataType device_product)
	Matrix vector product.

virtual void	dot_plus (const DataType device_vector, const DataType alpha, DataType device_product)
	Matrix vector product written in place.

virtual void	transpose_dot (const DataType device_vector, DataType device_product)
	Transposed-matrix vector product.

virtual void	transpose_dot_plus (const DataType device_vector, const DataType alpha, DataType device_product)
	Transposed-matrix vector product written in place.

Public Member Functions inherited from cuMatrix< DataType >
	cuMatrix ()
	Default constructor.

	cuMatrix (const FlagType A_is_symmetric_)
	Constructor.

virtual	~cuMatrix ()
	Destructor.

DataType	get_eigenvalue (const DataType known_parameters, const DataType known_eigenvalue, const DataType inquiry_parameters) const
	This virtual function is implemented from its pure virtual function of the base class. In this class, this functio has no use and was only implemented so that this class be able to be instantiated (due to the pure virtual function).

virtual void	set_symmetry (const FlagType symmetric)
	Specify whether the matrix is symmetic or non-symmetric.

Public Member Functions inherited from cuLinearOperator< DataType >
	cuLinearOperator ()
	Default constructor.

	cuLinearOperator (const int num_gpu_devices_)
	Constructor with setting `num_rows` and `num_columns`.

virtual	~cuLinearOperator ()
	Destructor.

cublasHandle_t	get_cublas_handle () const
	This function returns a reference to the `cublasHandle_t` object. The object will be created, if it is not created already.

void	set_parameters (DataType *parameters_)
	Sets the scalar parameter `this->parameters`. Parameter is initialized to `NULL`. However, before calling `dot` or `transpose_dot` functions, the parameters must be set.

Public Member Functions inherited from cLinearOperatorBase
	cLinearOperatorBase ()
	Default constructor.

	cLinearOperatorBase (const LongIndexType num_rows_, const LongIndexType num_columns_)
	Constructor with setting `num_rows` and `num_columns`.

virtual	~cLinearOperatorBase ()
	Destructor.

LongIndexType	get_num_rows () const
	Returns the number of rows of the matrix.

LongIndexType	get_num_columns () const
	Returns the number of columns of the matrix.

IndexType	get_num_parameters () const
	Returns the number of parameters of the linear operator.

FlagType	is_eigenvalue_relation_known () const
	Returns a flag that determines whether a relation between the parameters of the operator and its eigenvalue(s) is known.

Protected Member Functions
virtual void	copy_host_to_device ()
	Copies the member data from the host memory to the device memory.

void	allocate_buffer (const int device_id, cusparseOperation_t cusparse_operation, const DataType alpha, const DataType beta, cusparseDnVecDescr_t &cusparse_input_vector, cusparseDnVecDescr_t &cusparse_output_vector, cusparseSpMVAlg_t algorithm)
	Allocates an external buffer for matrix-vector multiplication using `cusparseSpMV` function.

Protected Member Functions inherited from cuLinearOperator< DataType >
int	query_gpu_devices () const
	Before any numerical computation, this method chechs if any gpu device is available on the machine, or notifies the user if nothing was found.

void	initialize_cublas_handle ()
	Creates a `cublasHandle_t` object, if not created already.

void	initialize_cusparse_handle ()
	Creates a `cusparseHandle_t` object, if not created already.

Protected Attributes
const DataType *	A_data

const LongIndexType *	A_indices

const LongIndexType *	A_index_pointer

DataType **	device_A_data

LongIndexType **	device_A_indices

LongIndexType **	device_A_index_pointer

void **	device_buffer

size_t *	device_buffer_num_bytes

cusparseSpMatDescr_t *	cusparse_matrix_A

Protected Attributes inherited from cuMatrix< DataType >
FlagType	A_is_symmetric

Protected Attributes inherited from cuLinearOperator< DataType >
int	num_gpu_devices

bool	copied_host_to_device

cublasHandle_t *	cublas_handle

cusparseHandle_t *	cusparse_handle

DataType *	parameters

Protected Attributes inherited from cLinearOperatorBase
const LongIndexType	num_rows

const LongIndexType	num_columns

FlagType	eigenvalue_relation_known

IndexType	num_parameters

Detailed Description

template<typename DataType>
class cuCSCMatrix< DataType >

Container for CSC matrices.

The cCSCMatrix holds a two-dimensional compressed sparse column matrix, and can perofrom matrix-vector product and transposed matrix-vector product.

See also: cuMatrix, cuDenseMatrix, cuCSRMatrix, cuCSCAffineMatrixFunction, cCSCMatrix

Definition at line 44 of file cu_csc_matrix.h.

Constructor & Destructor Documentation

◆ cuCSCMatrix() [1/2]

template<typename DataType >

cuCSCMatrix< DataType >::cuCSCMatrix ( )

Default constructor.

Definition at line 42 of file cu_csc_matrix.cu.

                                  :
    A_data(NULL),
    A_indices(NULL),
    A_index_pointer(NULL),
    device_A_data(NULL),
    device_A_indices(NULL),
    device_A_index_pointer(NULL),
    device_buffer(NULL),
    device_buffer_num_bytes(NULL),
    cusparse_matrix_A(NULL)
{
}

◆ cuCSCMatrix() [2/2]

template<typename DataType >

cuCSCMatrix< DataType >::cuCSCMatrix	(	const DataType *	A_data_,
		const LongIndexType *	A_indices_,
		const LongIndexType *	A_index_pointer_,
		const LongIndexType	num_rows_,
		const LongIndexType	num_columns_,
		const FlagType	A_is_symmetric_,
		const int	num_gpu_devices_
	)

Constructor.

Parameters

[in]	A_data_	1D array of the data content of sparse matrix. The size of the array is the nnz of the matrix.
[in]	A_indices_	1D array indicating the row of each element in `A_data_` . The size of this array is the nnz of the matrix.
[in]	A_index_pointer_	1D array pointing to the start of new rows in `A_indices_` . The size of this array is `num_rows+1` . The first element of this array is `0` and the last element of this array is the nnz of the matrix.
[in]	num_rows_	Number of rows of `A`
[in]	num_columns_	Number of columns of `A`
[in]	A_is_symmetric_	Boolean. If `A` is symmetric, set this value to `1`, otherwise `0`.
[in]	num_gpu_devices_	Number of GPU devices to be utilzied for parallel processing.

Definition at line 84 of file cu_csc_matrix.cu.

                                   :
 
    // Base class constructor
    cLinearOperatorBase(num_rows_, num_columns_),
    cuLinearOperator<DataType>(num_gpu_devices_),
    cuMatrix<DataType>(A_is_symmetric_),
 
    // Initializer list
    A_data(A_data_),
    A_indices(A_indices_),
    A_index_pointer(A_index_pointer_),
    device_A_data(NULL),
    device_A_indices(NULL),
    device_A_index_pointer(NULL),
    device_buffer(NULL),
    cusparse_matrix_A(NULL)
{
    this->initialize_cusparse_handle();
    this->copy_host_to_device();
 
    // Initialize device buffer
    this->device_buffer = new void*[this->num_gpu_devices];
    this->device_buffer_num_bytes = new size_t[this->num_gpu_devices];
    for (int device_id=0; device_id < this->num_gpu_devices; ++device_id)
    {
        this->device_buffer[device_id] = NULL;
        this->device_buffer_num_bytes[device_id] = 0;
    }
}

References cuCSCMatrix< DataType >::copy_host_to_device(), cuCSCMatrix< DataType >::device_buffer, cuCSCMatrix< DataType >::device_buffer_num_bytes, cuLinearOperator< DataType >::initialize_cusparse_handle(), and cuLinearOperator< DataType >::num_gpu_devices.

Here is the call graph for this function:

◆ ~cuCSCMatrix()

template<typename DataType >

cuCSCMatrix< DataType >::~cuCSCMatrix ( )

virtual

Destructor.

Definition at line 130 of file cu_csc_matrix.cu.

{
    // Member objects exist if the second constructor was called.
    if (this->copied_host_to_device)
    {
        // Deallocate arrays of data on gpu
        for (int device_id=0; device_id < this->num_gpu_devices; ++device_id)
        {
            // Switch to a device
            CudaAPI<DataType>::set_device(device_id);
 
            // Deallocate
            CudaAPI<DataType>::del(this->device_A_data[device_id]);
            CudaAPI<LongIndexType>::del(
                    this->device_A_indices[device_id]);
            CudaAPI<LongIndexType>::del(
                    this->device_A_index_pointer[device_id]);
            CudaAPI<LongIndexType>::del(this->device_buffer[device_id]);
            cusparse_api::destroy_cusparse_matrix(
                    this->cusparse_matrix_A[device_id]);
        }
    }
 
    // Deallocate arrays of pointers on cpu
    if (this->device_A_data != NULL)
    {
        delete[] this->device_A_data;
        this->device_A_data = NULL;
    }
 
    if (this->device_A_indices != NULL)
    {
        delete[] this->device_A_indices;
        this->device_A_indices = NULL;
    }
 
    if (this->device_A_index_pointer != NULL)
    {
        delete[] this->device_A_index_pointer;
        this->device_A_index_pointer = NULL;
    }
 
    if (this->device_buffer != NULL)
    {
        delete[] this->device_buffer;
        this->device_buffer = NULL;
    }
 
    if (this->device_buffer_num_bytes != NULL)
    {
        delete[] this->device_buffer_num_bytes;
        this->device_buffer_num_bytes = NULL;
    }
 
    if (this->cusparse_matrix_A != NULL)
    {
        delete[] this->cusparse_matrix_A;
        this->cusparse_matrix_A = NULL;
    }
}

References CudaAPI< ArrayType >::del(), cusparse_api::destroy_cusparse_matrix(), and CudaAPI< ArrayType >::set_device().

Here is the call graph for this function:

Member Function Documentation

◆ allocate_buffer()

template<typename DataType >

void cuCSCMatrix< DataType >::allocate_buffer	(	const int	device_id,
		cusparseOperation_t	cusparse_operation,
		const DataType	alpha,
		const DataType	beta,
		cusparseDnVecDescr_t &	cusparse_input_vector,
		cusparseDnVecDescr_t &	cusparse_output_vector,
		cusparseSpMVAlg_t	algorithm
	)

protected

Allocates an external buffer for matrix-vector multiplication using cusparseSpMV function.

If buffer size if not the same as required buffer size, allocate (or reallocate) memory. The allocation is always performed in the first call of this function since buffer size is initialized to zero in constructor. But for the next calls it might not be reallocated if the buffer size is the same.

Parameters

[in]	device_id	The ID of the GPU device, from `0` to `num_gpu_devices-1`.
[in]	cusparse_operation	The CuSparfse operation, which can be `CUSPARSE_OPERATION_NON_TRANSPOSE` or `CUSPARSE_OPERATION_TRANSPOSE`.
[in]	alpha	Scalar. The parameter \( \alpha \) in matrix-vector multiplication.
[in]	beta	Scalar. The parameter \( \beta \) in matrix-vector multiplication.
[in]	cusparse_input_vector	Input vector in the matrix-vector multiplication.
[in]	cusparse_output_vector	Output vector in the matrix-vector multiplication.
[in]	algorithm	CuSparse algorithm for sparse matrix-vector product. Possible values can be `CUSPARSE_SPMV_ALG_DEFAULT`, `CUSPARSE_SPMV_CSR_ALG1`, `CUSPARSE_SPMV_CSR_ALG2`, etc.

Definition at line 335 of file cu_csc_matrix.cu.

{
    // Find the buffer size needed for matrix-vector multiplication
    size_t required_buffer_size;
    cusparse_api::cusparse_matrix_buffer_size(
            this->cusparse_handle[device_id], cusparse_operation, alpha,
            this->cusparse_matrix_A[device_id], cusparse_input_vector, beta,
            cusparse_output_vector, algorithm, &required_buffer_size);
 
    if (this->device_buffer_num_bytes[device_id] != required_buffer_size)
    {
        // Update the buffer size
        this->device_buffer_num_bytes[device_id] = required_buffer_size;
 
        // Delete buffer if it was allocated previously
        CudaAPI<DataType>::del(this->device_buffer[device_id]);
 
        // Allocate (or reallocate) buffer on device.
        CudaAPI<DataType>::alloc_bytes(
                this->device_buffer[device_id],
                this->device_buffer_num_bytes[device_id]);
    }
}

References CudaAPI< ArrayType >::alloc_bytes(), cusparse_api::cusparse_matrix_buffer_size(), and CudaAPI< ArrayType >::del().

Here is the call graph for this function:

◆ copy_host_to_device()

template<typename DataType >

void cuCSCMatrix< DataType >::copy_host_to_device ( )

protectedvirtual

Copies the member data from the host memory to the device memory.

Note: CUDA below version 12 does not have a cusparse api for CSC matrices. As such, for CUDA<12, we treat a CSC matrix as a CSR matrix, but using trapnsposed operations. In addition, we swap the number of columns and rows from the input matrix to the cusparse matrix.

Implements cuMatrix< DataType >.

Definition at line 205 of file cu_csc_matrix.cu.

{
    if (!this->copied_host_to_device)
    {
        // Set the number of threads
        #if defined(USE_OPENMP) && (USE_OPENMP == 1)
            omp_set_num_threads(this->num_gpu_devices);
        #endif
 
        // Array sizes
        LongIndexType A_data_size = this->get_nnz();
        LongIndexType A_indices_size = A_data_size;
        LongIndexType A_index_pointer_size = this->num_rows + 1;
        LongIndexType A_nnz = this->get_nnz();
 
        // CuSparse API in CUDA below 12 does not support CSC matrix
        #ifndef CUDA_VERSION
            #error CUDA_VERSION Undefined!
        #elif CUDA_VERSION < 12000
            // Swapping the number of rows and columns to treat the input CSC
            // matrix as a CSR matrix.
            LongIndexType csc_num_rows = this->num_columns;
            LongIndexType csc_num_columns = this->num_rows;
        #endif
 
        // Create array of pointers for data on each gpu device
        this->device_A_data = new DataType*[this->num_gpu_devices];
        this->device_A_indices = new LongIndexType*[this->num_gpu_devices];
        this->device_A_index_pointer = \
            new LongIndexType*[this->num_gpu_devices];
        this->cusparse_matrix_A = \
            new cusparseSpMatDescr_t[this->num_gpu_devices];
 
        #if defined(USE_OPENMP) && (USE_OPENMP == 1)
        #pragma omp parallel
        #endif
        {
            // Switch to a device with the same device id as the cpu thread id
            unsigned int thread_id;
            #if defined(USE_OPENMP) && (USE_OPENMP == 1)
                thread_id = omp_get_thread_num();
            #else
                thread_id = 0;
            #endif
 
            CudaAPI<DataType>::set_device(thread_id);
 
            // A_data
            CudaAPI<DataType>::alloc(this->device_A_data[thread_id],
                                           A_data_size);
            CudaAPI<DataType>::copy_to_device(
                    this->A_data, A_data_size, this->device_A_data[thread_id]);
 
            // A_indices
            CudaAPI<LongIndexType>::alloc(
                    this->device_A_indices[thread_id], A_indices_size);
            CudaAPI<LongIndexType>::copy_to_device(
                    this->A_indices, A_indices_size,
                    this->device_A_indices[thread_id]);
 
            // A_index_pointer
            CudaAPI<LongIndexType>::alloc(
                    this->device_A_index_pointer[thread_id],
                    A_index_pointer_size);
            CudaAPI<LongIndexType>::copy_to_device(
                    this->A_index_pointer, A_index_pointer_size,
                    this->device_A_index_pointer[thread_id]);
 
            // Create cusparse matrix
            #ifndef CUDA_VERSION
                #error CUDA_VERSION Undefined!
            #elif CUDA_VERSION < 12000
                // Treat CSC as CSR matrix with swapped columns and rows
                cusparse_api::create_cusparse_csr_matrix(
                        this->cusparse_matrix_A[thread_id], csc_num_rows,
                        csc_num_columns, A_nnz, this->device_A_data[thread_id],
                        this->device_A_indices[thread_id],
                        this->device_A_index_pointer[thread_id]);
            #else
                // Use CSC api in CUDA >= 12
                cusparse_api::create_cusparse_csc_matrix(
                        this->cusparse_matrix_A[thread_id], this->num_rows,
                        this->num_columns, A_nnz,
                        this->device_A_data[thread_id],
                        this->device_A_indices[thread_id],
                        this->device_A_index_pointer[thread_id]);
            #endif
        }
 
        // Flag to prevent reinitialization
        this->copied_host_to_device = true;
    }
}

References CudaAPI< ArrayType >::alloc(), CudaAPI< ArrayType >::copy_to_device(), cusparse_api::create_cusparse_csc_matrix(), cusparse_api::create_cusparse_csr_matrix(), omp_get_thread_num(), omp_set_num_threads(), and CudaAPI< ArrayType >::set_device().

Referenced by cuCSCMatrix< DataType >::cuCSCMatrix().

Here is the call graph for this function:

Here is the caller graph for this function:

◆ dot()

template<typename DataType >

void cuCSCMatrix< DataType >::dot	(	const DataType *	device_vector,
		DataType *	device_product
	)

virtual

Matrix vector product.

Performs the matrix vector product \( \boldsymbol{y} = \mathbf{A} \boldsymbol{x} \).

Parameters

[in]	device_vector	A one-dimensional input vector \( \boldsymbol{x} \) with size the of the number of columns of the matrix \( \mathbf{A} \). This array should be on GPU device.
[out]	device_product	A one-dimensional output vector \( \boldsymbol{y} \) with the size of the number of rows of \( \mathbf{A} \). This vector will be overwritten. This array should be on GPU device.

See also: cuCSCMatrix::dot_plus, cuCSCMatrix::transposed_dot cuCSCMatrix::transposed_dot_plus

Implements cuLinearOperator< DataType >.

Definition at line 478 of file cu_csc_matrix.cu.

{
    assert(this->copied_host_to_device);
 
    // Create cusparse vector for the input vector
    cusparseDnVecDescr_t cusparse_input_vector;
    cusparse_api::create_cusparse_vector(
            cusparse_input_vector, this->num_columns,
            const_cast<DataType*>(device_vector));
 
    // Create cusparse vector for the output vector
    cusparseDnVecDescr_t cusparse_output_vector;
    cusparse_api::create_cusparse_vector(
            cusparse_output_vector, this->num_rows, device_product);
 
    // Matrix vector settings
    DataType alpha = cu_arithmetics::cast<float, DataType>(1.0f);
    DataType beta = cu_arithmetics::cast<float, DataType>(0.0f);
 
    #ifndef CUDA_VERSION
        #error CUDA_VERSION Undefined!
    #elif CUDA_VERSION < 12000
        // Using transpose operation since we treat CSC matrix as CSR
        cusparseOperation_t cusparse_operation = CUSPARSE_OPERATION_TRANSPOSE;
    #else
        cusparseOperation_t cusparse_operation = \
            CUSPARSE_OPERATION_NON_TRANSPOSE;
    #endif
 
    cusparseSpMVAlg_t algorithm = CUSPARSE_SPMV_ALG_DEFAULT;
 
    // Get device id
    int device_id = CudaAPI<DataType>::get_device();
 
    // Allocate device buffer (or reallocation if needed)
    this->allocate_buffer(device_id, cusparse_operation, alpha, beta,
                          cusparse_input_vector, cusparse_output_vector,
                          algorithm);
 
    // Matrix vector multiplication
    cusparse_api::cusparse_matvec(
            this->cusparse_handle[device_id], cusparse_operation, alpha,
            this->cusparse_matrix_A[device_id], cusparse_input_vector, beta,
            cusparse_output_vector, algorithm, this->device_buffer[device_id]);
 
    // Destroy cusparse vectors
    cusparse_api::destroy_cusparse_vector(cusparse_input_vector);
    cusparse_api::destroy_cusparse_vector(cusparse_output_vector);
}

References cu_arithmetics::abs(), cusparse_api::create_cusparse_vector(), cusparse_api::cusparse_matvec(), CUSPARSE_SPMV_ALG_DEFAULT, cusparse_api::destroy_cusparse_vector(), and CudaAPI< ArrayType >::get_device().

Here is the call graph for this function:

◆ dot_plus()

template<typename DataType >

void cuCSCMatrix< DataType >::dot_plus	(	const DataType *	device_vector,
		const DataType	alpha,
		DataType *	device_product
	)

virtual

Matrix vector product written in place.

Performs the matrix vector product \( \boldsymbol{y} = \boldsymbol{y} + \alpha \mathbf{A} \boldsymbol{x} \).

Parameters

[in]	device_vector	A one-dimensional input vector \( \boldsymbol{x} \) with size the of the number of columns of the matrix \( \mathbf{A} \). This array should be on GPU device.
[in]	alpha	A scalar.
[out]	device_product	A one-dimensional output vector \( \boldsymbol{y} \) with the size of the number of rows of \( \mathbf{A} \). This array should be on GPU device.

See also: cuCSCMatrix::dot, cuCSCMatrix::transposed_dot cuCSCMatrix::transposed_dot_plus

Implements cuMatrix< DataType >.

Definition at line 556 of file cu_csc_matrix.cu.

{
    assert(this->copied_host_to_device);
 
    // Create cusparse vector for the input vector
    cusparseDnVecDescr_t cusparse_input_vector;
    cusparse_api::create_cusparse_vector(
            cusparse_input_vector, this->num_columns,
            const_cast<DataType*>(device_vector));
 
    // Create cusparse vector for the output vector
    cusparseDnVecDescr_t cusparse_output_vector;
    cusparse_api::create_cusparse_vector(
            cusparse_output_vector, this->num_rows, device_product);
 
    // Matrix vector settings
    DataType beta = cu_arithmetics::cast<float, DataType>(1.0f);
 
    #ifndef CUDA_VERSION
        #error CUDA_VERSION Undefined!
    #elif CUDA_VERSION < 12000
        // Using transpose operation since we treat CSC matrix as CSR
        cusparseOperation_t cusparse_operation = CUSPARSE_OPERATION_TRANSPOSE;
    #else
        cusparseOperation_t cusparse_operation = \
            CUSPARSE_OPERATION_NON_TRANSPOSE;
    #endif
 
    cusparseSpMVAlg_t algorithm = CUSPARSE_SPMV_ALG_DEFAULT;
 
    // Get device id
    int device_id = CudaAPI<DataType>::get_device();
 
    // Allocate device buffer (or reallocation if needed)
    this->allocate_buffer(device_id, cusparse_operation, alpha, beta,
                          cusparse_input_vector, cusparse_output_vector,
                          algorithm);
 
    // Matrix vector multiplication
    cusparse_api::cusparse_matvec(
            this->cusparse_handle[device_id], cusparse_operation, alpha,
            this->cusparse_matrix_A[device_id], cusparse_input_vector, beta,
            cusparse_output_vector, algorithm, this->device_buffer[device_id]);
 
    // Destroy cusparse vectors
    cusparse_api::destroy_cusparse_vector(cusparse_input_vector);
    cusparse_api::destroy_cusparse_vector(cusparse_output_vector);
}

References cu_arithmetics::abs(), cusparse_api::create_cusparse_vector(), cusparse_api::cusparse_matvec(), CUSPARSE_SPMV_ALG_DEFAULT, cusparse_api::destroy_cusparse_vector(), and CudaAPI< ArrayType >::get_device().

Here is the call graph for this function:

◆ get_nnz()

template<typename DataType >

LongIndexType cuCSCMatrix< DataType >::get_nnz ( ) const

Returns the number of non-zero elements of the sparse matrix.

The nnz of a CSC matrix can be obtained from the last element of A_index_pointer. The size of array A_index_pointer is one plus the number of columns of the matrix.

Returns: The nnz of the matrix.

Definition at line 449 of file cu_csc_matrix.cu.

{
    return this->A_index_pointer[this->num_columns];
}

◆ is_identity_matrix()

template<typename DataType >

FlagType cuCSCMatrix< DataType >::is_identity_matrix ( ) const

virtual

Checks whether the matrix is identity.

The identity check is primarily performed in the cAffineMatrixFunction class.

Returns: Returns 1 if the input matrix is identity, and 0 otherwise.

See also: cAffineMatrixFunction

Implements cuMatrix< DataType >.

Definition at line 381 of file cu_csc_matrix.cu.

{
    FlagType matrix_is_identity = 1;
    LongIndexType index_pointer;
    LongIndexType row;
    DataType matrix_element;
    const DataType diagonal = 1.0;
    const DataType off_diagonal = 0.0;
 
    // Check matrix element-wise
    #if defined(USE_OPENMP) && (USE_OPENMP == 1)
    #pragma omp parallel for \
        schedule(static) \
        if (!omp_in_parallel()) \
        default(none) \
        shared(matrix_is_identity, diagonal, off_diagonal) \
        private(index_pointer, row, matrix_element)
    #endif
    for (LongIndexType column=0; column < this->num_columns; ++column)
    {
        if (matrix_is_identity)
        {
            for (index_pointer=this->A_index_pointer[column];
                 index_pointer < this->A_index_pointer[column+1];
                 ++index_pointer)
            {
                row = this->A_indices[index_pointer];
 
                if (!((this->A_is_symmetric) && (column >= row)))
                {
                    matrix_element = this->A_data[index_pointer];
 
                    if (((row == column) && \
                         (!cu_arithmetics::is_equal(matrix_element,
                                                    diagonal))) || \
                        ((row != column) && \
                         (!cu_arithmetics::is_equal(matrix_element,
                                                    off_diagonal))))
                    {
                        #if defined(USE_OPENMP) && (USE_OPENMP == 1)
                        #pragma omp atomic write
                        #endif
                        matrix_is_identity = 0;
 
                        break;
                    }
                }
            }
        }
    }
 
    return matrix_is_identity;
}

References cu_arithmetics::is_equal().

Here is the call graph for this function:

◆ transpose_dot()

template<typename DataType >

void cuCSCMatrix< DataType >::transpose_dot	(	const DataType *	device_vector,
		DataType *	device_product
	)

virtual

Transposed-matrix vector product.

Performs the matrix vector product \( \boldsymbol{y} = \mathbf{A}^{\intercal} \boldsymbol{x} \).

Parameters

[in]	device_vector	A one-dimensional input vector \( \boldsymbol{x} \) with size the of the number of columns of the matrix \( \mathbf{A} \). This array should be on GPU device.
[out]	device_product	A one-dimensional output vector \( \boldsymbol{y} \) with the size of the number of rows of \( \mathbf{A} \). This vector will be overwritten. This array should be on GPU device.

See also: cuCSCMatrix::dot_plus, cuCSCMatrix::dot cuCSCMatrix::transposed_dot_plus

Implements cuLinearOperator< DataType >.

Definition at line 632 of file cu_csc_matrix.cu.

{
    assert(this->copied_host_to_device);
 
    // Create cusparse vector for the input vector
    cusparseDnVecDescr_t cusparse_input_vector;
    cusparse_api::create_cusparse_vector(
            cusparse_input_vector, this->num_columns,
            const_cast<DataType*>(device_vector));
 
    // Create cusparse vector for the output vector
    cusparseDnVecDescr_t cusparse_output_vector;
    cusparse_api::create_cusparse_vector(
            cusparse_output_vector, this->num_rows, device_product);
 
    // Matrix vector settings
    DataType alpha = cu_arithmetics::cast<float, DataType>(1.0f);
    DataType beta = cu_arithmetics::cast<float, DataType>(0.0f);
 
    #ifndef CUDA_VERSION
        #error CUDA_VERSION Undefined!
    #elif CUDA_VERSION < 12000
        // Using non-transpose operation since we treat CSC matrix as CSR
        cusparseOperation_t cusparse_operation = \
            CUSPARSE_OPERATION_NON_TRANSPOSE;
    #else
        cusparseOperation_t cusparse_operation = CUSPARSE_OPERATION_TRANSPOSE;
    #endif
 
    cusparseSpMVAlg_t algorithm = CUSPARSE_SPMV_ALG_DEFAULT;
 
    // Get device id
    int device_id = CudaAPI<DataType>::get_device();
 
    // Allocate device buffer (or reallocation if needed)
    this->allocate_buffer(device_id, cusparse_operation, alpha, beta,
                          cusparse_input_vector, cusparse_output_vector,
                          algorithm);
 
    // Matrix vector multiplication
    cusparse_api::cusparse_matvec(
            this->cusparse_handle[device_id], cusparse_operation, alpha,
            this->cusparse_matrix_A[device_id], cusparse_input_vector, beta,
            cusparse_output_vector, algorithm, this->device_buffer[device_id]);
 
    // Destroy cusparse vectors
    cusparse_api::destroy_cusparse_vector(cusparse_input_vector);
    cusparse_api::destroy_cusparse_vector(cusparse_output_vector);
}

References cu_arithmetics::abs(), cusparse_api::create_cusparse_vector(), cusparse_api::cusparse_matvec(), CUSPARSE_SPMV_ALG_DEFAULT, cusparse_api::destroy_cusparse_vector(), and CudaAPI< ArrayType >::get_device().

Here is the call graph for this function:

◆ transpose_dot_plus()

template<typename DataType >

void cuCSCMatrix< DataType >::transpose_dot_plus	(	const DataType *	device_vector,
		const DataType	alpha,
		DataType *	device_product
	)

virtual

Transposed-matrix vector product written in place.

Performs the matrix vector product \( \boldsymbol{y} = \boldsymbol{y} + \alpha \mathbf{A}^{\intercal} \boldsymbol{x} \).

Parameters

[in]	device_vector	A one-dimensional input vector \( \boldsymbol{x} \) with size the of the number of columns of the matrix \( \mathbf{A} \). This array should be on GPU device.
[in]	alpha	A scalar.
[out]	device_product	A one-dimensional output vector \( \boldsymbol{y} \) with the size of the number of rows of \( \mathbf{A} \). This array should be on GPU device.

See also: cuCSCMatrix::dot_plus, cuCSCMatrix::transposed_dot cuCSCMatrix::dot

Implements cuMatrix< DataType >.

Definition at line 711 of file cu_csc_matrix.cu.

{
    assert(this->copied_host_to_device);
 
    // Create cusparse vector for the input vector
    cusparseDnVecDescr_t cusparse_input_vector;
    cusparse_api::create_cusparse_vector(
            cusparse_input_vector, this->num_columns,
            const_cast<DataType*>(device_vector));
 
    // Create cusparse vector for the output vector
    cusparseDnVecDescr_t cusparse_output_vector;
    cusparse_api::create_cusparse_vector(
            cusparse_output_vector, this->num_rows, device_product);
 
    // Matrix vector settings
    DataType beta = cu_arithmetics::cast<float, DataType>(1.0f);
 
    #ifndef CUDA_VERSION
        #error CUDA_VERSION Undefined!
    #elif CUDA_VERSION < 12000
        // Using non-transpose operation since we treat CSC matrix as CSR
        cusparseOperation_t cusparse_operation = \
            CUSPARSE_OPERATION_NON_TRANSPOSE;
    #else
        cusparseOperation_t cusparse_operation = CUSPARSE_OPERATION_TRANSPOSE;
    #endif
 
    cusparseSpMVAlg_t algorithm = CUSPARSE_SPMV_ALG_DEFAULT;
 
    // Get device id
    int device_id = CudaAPI<DataType>::get_device();
 
    // Allocate device buffer (or reallocation if needed)
    this->allocate_buffer(device_id, cusparse_operation, alpha, beta,
                          cusparse_input_vector, cusparse_output_vector,
                          algorithm);
 
    // Matrix vector multiplication
    cusparse_api::cusparse_matvec(
            this->cusparse_handle[device_id], cusparse_operation, alpha,
            this->cusparse_matrix_A[device_id], cusparse_input_vector, beta,
            cusparse_output_vector, algorithm, this->device_buffer[device_id]);
 
    // Destroy cusparse vectors
    cusparse_api::destroy_cusparse_vector(cusparse_input_vector);
    cusparse_api::destroy_cusparse_vector(cusparse_output_vector);
}

References cu_arithmetics::abs(), cusparse_api::create_cusparse_vector(), cusparse_api::cusparse_matvec(), CUSPARSE_SPMV_ALG_DEFAULT, cusparse_api::destroy_cusparse_vector(), and CudaAPI< ArrayType >::get_device().

Here is the call graph for this function:

Definition at line 105 of file cu_csc_matrix.h.

Referenced by cuCSCMatrix< DataType >::cuCSCMatrix().

◆ device_buffer_num_bytes

template<typename DataType >

size_t* cuCSCMatrix< DataType >::device_buffer_num_bytes

protected

Definition at line 106 of file cu_csc_matrix.h.

Referenced by cuCSCMatrix< DataType >::cuCSCMatrix().

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

/home/runner/work/imate/imate/imate/_cu_linear_operator/cu_csc_matrix.h
/home/runner/work/imate/imate/imate/_cu_linear_operator/cu_csc_matrix.cu

Public Member Functions

Protected Member Functions

Protected Attributes

Detailed Description

Constructor & Destructor Documentation

◆ cuCSCMatrix() [1/2]

◆ cuCSCMatrix() [2/2]

◆ ~cuCSCMatrix()

Member Function Documentation

◆ allocate_buffer()

◆ copy_host_to_device()

◆ dot()

◆ dot_plus()

◆ get_nnz()

◆ is_identity_matrix()

◆ transpose_dot()

◆ transpose_dot_plus()

Member Data Documentation

◆ A_data

◆ A_index_pointer

◆ A_indices

◆ cusparse_matrix_A

◆ device_A_data

◆ device_A_index_pointer

◆ device_A_indices

◆ device_buffer

◆ device_buffer_num_bytes