Sparse level 3 functions

Sparse level 3 functions#

This module contains all sparse level 3 routines.

The sparse level 3 routines describe operations between a matrix in sparse format and multiple vectors in dense format that can also be seen as a dense matrix.

hipsparseXbsrmm()#

hipsparseStatus_t hipsparseSbsrmm(hipsparseHandle_t handle, hipsparseDirection_t dirA, hipsparseOperation_t transA, hipsparseOperation_t transB, int mb, int n, int kb, int nnzb, const float *alpha, const hipsparseMatDescr_t descrA, const float *bsrValA, const int *bsrRowPtrA, const int *bsrColIndA, int blockDim, const float *B, int ldb, const float *beta, float *C, int ldc)#

hipsparseStatus_t hipsparseDbsrmm(hipsparseHandle_t handle, hipsparseDirection_t dirA, hipsparseOperation_t transA, hipsparseOperation_t transB, int mb, int n, int kb, int nnzb, const double *alpha, const hipsparseMatDescr_t descrA, const double *bsrValA, const int *bsrRowPtrA, const int *bsrColIndA, int blockDim, const double *B, int ldb, const double *beta, double *C, int ldc)#

hipsparseStatus_t hipsparseCbsrmm(hipsparseHandle_t handle, hipsparseDirection_t dirA, hipsparseOperation_t transA, hipsparseOperation_t transB, int mb, int n, int kb, int nnzb, const hipComplex *alpha, const hipsparseMatDescr_t descrA, const hipComplex *bsrValA, const int *bsrRowPtrA, const int *bsrColIndA, int blockDim, const hipComplex *B, int ldb, const hipComplex *beta, hipComplex *C, int ldc)#

hipsparseStatus_t hipsparseZbsrmm(hipsparseHandle_t handle, hipsparseDirection_t dirA, hipsparseOperation_t transA, hipsparseOperation_t transB, int mb, int n, int kb, int nnzb, const hipDoubleComplex *alpha, const hipsparseMatDescr_t descrA, const hipDoubleComplex *bsrValA, const int *bsrRowPtrA, const int *bsrColIndA, int blockDim, const hipDoubleComplex *B, int ldb, const hipDoubleComplex *beta, hipDoubleComplex *C, int ldc)#

Sparse matrix dense matrix multiplication using BSR storage format.

hipsparseXbsrmm multiplies the scalar \(\alpha\) with a sparse \(m \times k\) matrix \(A\), defined in BSR storage format, and the column-oriented dense \(k \times n\) matrix \(B\) and adds the result to the column-oriented dense \(m \times n\) matrix \(C\) that is multiplied by the scalar \(\beta\), such that

\[ C := \alpha \cdot op(A) \cdot op(B) + \beta \cdot C, \]

with

\[\begin{split} op(A) = \left\{ \begin{array}{ll} A, & \text{if transA == HIPSPARSE_OPERATION_NON_TRANSPOSE} \\ \end{array} \right. \end{split}\]

and

\[\begin{split} op(B) = \left\{ \begin{array}{ll} B, & \text{if transB == HIPSPARSE_OPERATION_NON_TRANSPOSE} \\ B^T, & \text{if transB == HIPSPARSE_OPERATION_TRANSPOSE} \\ \end{array} \right. \end{split}\]

and where \(k = blockDim \times kb\) and \(m = blockDim \times mb\).

Note

This function is non blocking and executed asynchronously with respect to the host. It may return before the actual computation has finished.

Note

Currently, only transA == HIPSPARSE_OPERATION_NON_TRANSPOSE is supported.

Parameters:

handle – [in] handle to the hipsparse library context queue.
dirA – [in] the storage format of the blocks. Can be HIPSPARSE_DIRECTION_ROW or HIPSPARSE_DIRECTION_COLUMN.
transA – [in] matrix \(A\) operation type. Currently, only HIPSPARSE_OPERATION_NON_TRANSPOSE is supported.
transB – [in] matrix \(B\) operation type. Currently, only HIPSPARSE_OPERATION_NON_TRANSPOSE and HIPSPARSE_OPERATION_TRANSPOSE are supported.
mb – [in] number of block rows of the sparse BSR matrix \(A\). Must be non-negative.
n – [in] number of columns of the dense matrix \(op(B)\) and \(C\). Must be non-negative.
kb – [in] number of block columns of the sparse BSR matrix \(A\). Must be non-negative.
nnzb – [in] number of non-zero blocks of the sparse BSR matrix \(A\). Must be non-negative.
alpha – [in] scalar \(\alpha\).
descrA – [in] descriptor of the sparse BSR matrix \(A\). Currently, only HIPSPARSE_MATRIX_TYPE_GENERAL is supported.
bsrValA – [in] array of nnzb*blockDim*blockDim elements of the sparse BSR matrix \(A\).
bsrRowPtrA – [in] array of mb+1 elements that point to the start of every block row of the sparse BSR matrix \(A\).
bsrColIndA – [in] array of nnzb elements containing the block column indices of the sparse BSR matrix \(A\).
blockDim – [in] size of the blocks in the sparse BSR matrix. Must be positive.
B – [in] array of dimension ldb*n ( \(op(B) == B\)), ldb*k otherwise.
ldb – [in] leading dimension of \(B\), must be at least \(\max{(1, k)}\) ( \( op(B) == B\)) where k=blockDim*kb, \(\max{(1, n)}\) otherwise.
beta – [in] scalar \(\beta\).
C – [inout] array of dimension ldc*n.
ldc – [in] leading dimension of \(C\), must be at least \(\max{(1, m)}\) ( \( op(A) == A\)) where m=blockDim*mb, \(\max{(1, k)}\) where k=blockDim*kb otherwise.

Return values:

HIPSPARSE_STATUS_SUCCESS – the operation completed successfully.
HIPSPARSE_STATUS_NOT_INITIALIZED – handle is not initialized.
HIPSPARSE_STATUS_INVALID_VALUE – handle, descrA, alpha or beta is nullptr, mb, n, kb or nnzb is negative, ldb or ldc is invalid, blockDim is less than or equal to zero, or bsrValA, bsrRowPtrA, bsrColIndA, B or C is nullptr.
HIPSPARSE_STATUS_ARCH_MISMATCH – the device is not supported.
HIPSPARSE_STATUS_NOT_SUPPORTED – transA is not HIPSPARSE_OPERATION_NON_TRANSPOSE, transB is HIPSPARSE_OPERATION_CONJUGATE_TRANSPOSE, or hipsparseMatrixType_t is not HIPSPARSE_MATRIX_TYPE_GENERAL.

int main(int argc, char* argv[])
{
    // hipSPARSE handle
    hipsparseHandle_t handle;
    HIPSPARSE_CHECK(hipsparseCreate(&handle));

    //     1 2 0 3 0 0
    // A = 0 4 5 0 0 0
    //     0 0 0 7 8 0
    //     0 0 1 2 4 1

    const int                  blockDim = 2;
    const int                  mb       = 2;
    const int                  kb       = 3;
    const int                  nnzb     = 4;
    const hipsparseDirection_t dir      = HIPSPARSE_DIRECTION_ROW;

    std::vector<int>   hbsrRowPtr = {0, 2, 4};
    std::vector<int>   hbsrColInd = {0, 1, 1, 2};
    std::vector<float> hbsrVal    = {1, 2, 0, 4, 0, 3, 5, 0, 0, 7, 1, 2, 8, 0, 4, 1};

    // Set dimension n of B
    const int n = 3;
    const int m = mb * blockDim;
    const int k = kb * blockDim;

    // Allocate and generate dense matrix B (k x n)
    std::vector<float> hB = {1.0f,
                             2.0f,
                             3.0f,
                             4.0f,
                             5.0f,
                             6.0f,
                             7.0f,
                             8.0f,
                             9.0f,
                             10.0f,
                             11.0f,
                             12.0f,
                             13.0f,
                             14.0f,
                             15.0f,
                             16.0f,
                             17.0f,
                             18.0f};

    int*   dbsrRowPtr = NULL;
    int*   dbsrColInd = NULL;
    float* dbsrVal    = NULL;
    HIP_CHECK(hipMalloc((void**)&dbsrRowPtr, sizeof(int) * (mb + 1)));
    HIP_CHECK(hipMalloc((void**)&dbsrColInd, sizeof(int) * nnzb));
    HIP_CHECK(hipMalloc((void**)&dbsrVal, sizeof(float) * nnzb * blockDim * blockDim));
    HIP_CHECK(
        hipMemcpy(dbsrRowPtr, hbsrRowPtr.data(), sizeof(int) * (mb + 1), hipMemcpyHostToDevice));
    HIP_CHECK(hipMemcpy(dbsrColInd, hbsrColInd.data(), sizeof(int) * nnzb, hipMemcpyHostToDevice));
    HIP_CHECK(hipMemcpy(dbsrVal,
                        hbsrVal.data(),
                        sizeof(float) * nnzb * blockDim * blockDim,
                        hipMemcpyHostToDevice));

    // Copy B to the device
    float* dB;
    HIP_CHECK(hipMalloc((void**)&dB, sizeof(float) * k * n));
    HIP_CHECK(hipMemcpy(dB, hB.data(), sizeof(float) * k * n, hipMemcpyHostToDevice));

    // alpha and beta
    float alpha = 1.0f;
    float beta  = 0.0f;

    // Allocate memory for the resulting matrix C
    float* dC;
    HIP_CHECK(hipMalloc((void**)&dC, sizeof(float) * m * n));

    // Matrix descriptor
    hipsparseMatDescr_t descr;
    HIPSPARSE_CHECK(hipsparseCreateMatDescr(&descr));

    // Perform the matrix multiplication
    HIPSPARSE_CHECK(hipsparseSbsrmm(handle,
                                    dir,
                                    HIPSPARSE_OPERATION_NON_TRANSPOSE,
                                    HIPSPARSE_OPERATION_NON_TRANSPOSE,
                                    mb,
                                    n,
                                    kb,
                                    nnzb,
                                    &alpha,
                                    descr,
                                    dbsrVal,
                                    dbsrRowPtr,
                                    dbsrColInd,
                                    blockDim,
                                    dB,
                                    k,
                                    &beta,
                                    dC,
                                    m));

    // Copy results to host
    std::vector<float> hC(m * n);
    HIP_CHECK(hipMemcpy(hC.data(), dC, sizeof(float) * m * n, hipMemcpyDeviceToHost));

    std::cout << "hC" << std::endl;
    for(int i = 0; i < m * n; i++)
    {
        std::cout << hC[i] << " ";
    }
    std::cout << std::endl;

    HIP_CHECK(hipFree(dbsrRowPtr));
    HIP_CHECK(hipFree(dbsrColInd));
    HIP_CHECK(hipFree(dbsrVal));
    HIP_CHECK(hipFree(dB));
    HIP_CHECK(hipFree(dC));

    return 0;
}

int main(int argc, char* argv[])
{
    // hipSPARSE handle
    hipsparseHandle_t handle;
    HIPSPARSE_CHECK(hipsparseCreate(&handle));

    //     1 2 0 3 0 0
    // A = 0 4 5 0 0 0
    //     0 0 0 7 8 0
    //     0 0 1 2 4 1

    const int                  blockDim = 2;
    const int                  mb       = 2;
    const int                  kb       = 3;
    const int                  nnzb     = 4;
    const hipsparseDirection_t dir      = HIPSPARSE_DIRECTION_ROW;

    int   hbsrRowPtr[] = {0, 2, 4};
    int   hbsrColInd[] = {0, 1, 1, 2};
    float hbsrVal[]    = {1, 2, 0, 4, 0, 3, 5, 0, 0, 7, 1, 2, 8, 0, 4, 1};

    // Set dimension n of B
    const int n = 3;
    const int m = mb * blockDim;
    const int k = kb * blockDim;

    // Allocate and generate dense matrix B (k x n)
    float hB[] = {1.0,
                  2.0,
                  3.0,
                  4.0,
                  5.0,
                  6.0,
                  7.0,
                  8.0,
                  9.0,
                  10.0,
                  11.0,
                  12.0,
                  13.0,
                  14.0,
                  15.0,
                  16.0,
                  17.0,
                  18.0};

    int*   dbsrRowPtr = NULL;
    int*   dbsrColInd = NULL;
    float* dbsrVal    = NULL;
    HIP_CHECK(hipMalloc((void**)&dbsrRowPtr, sizeof(int) * (mb + 1)));
    HIP_CHECK(hipMalloc((void**)&dbsrColInd, sizeof(int) * nnzb));
    HIP_CHECK(hipMalloc((void**)&dbsrVal, sizeof(float) * nnzb * blockDim * blockDim));
    HIP_CHECK(hipMemcpy(dbsrRowPtr, hbsrRowPtr, sizeof(int) * (mb + 1), hipMemcpyHostToDevice));
    HIP_CHECK(hipMemcpy(dbsrColInd, hbsrColInd, sizeof(int) * nnzb, hipMemcpyHostToDevice));
    HIP_CHECK(hipMemcpy(
        dbsrVal, hbsrVal, sizeof(float) * nnzb * blockDim * blockDim, hipMemcpyHostToDevice));

    // Copy B to the device
    float* dB;
    HIP_CHECK(hipMalloc((void**)&dB, sizeof(float) * k * n));
    HIP_CHECK(hipMemcpy(dB, hB, sizeof(float) * k * n, hipMemcpyHostToDevice));

    // alpha and beta
    float alpha = 1.0;
    float beta  = 0.0;

    // Allocate memory for the resulting matrix C
    float* dC;
    HIP_CHECK(hipMalloc((void**)&dC, sizeof(float) * m * n));

    // Matrix descriptor
    hipsparseMatDescr_t descr;
    HIPSPARSE_CHECK(hipsparseCreateMatDescr(&descr));

    // Perform the matrix multiplication
    HIPSPARSE_CHECK(hipsparseSbsrmm(handle,
                                    dir,
                                    HIPSPARSE_OPERATION_NON_TRANSPOSE,
                                    HIPSPARSE_OPERATION_NON_TRANSPOSE,
                                    mb,
                                    n,
                                    kb,
                                    nnzb,
                                    &alpha,
                                    descr,
                                    dbsrVal,
                                    dbsrRowPtr,
                                    dbsrColInd,
                                    blockDim,
                                    dB,
                                    k,
                                    &beta,
                                    dC,
                                    m));

    // Copy results to host
    float hC[m * n];
    HIP_CHECK(hipMemcpy(hC, dC, sizeof(float) * m * n, hipMemcpyDeviceToHost));

    printf("hC\n");
    for(int i = 0; i < m * n; i++)
    {
        printf("%f ", hC[i]);
    }
    printf("\n");

    HIP_CHECK(hipFree(dbsrRowPtr));
    HIP_CHECK(hipFree(dbsrColInd));
    HIP_CHECK(hipFree(dbsrVal));
    HIP_CHECK(hipFree(dB));
    HIP_CHECK(hipFree(dC));

    return 0;
}

program example_hipsparse_bsrmm
    use iso_c_binding
    implicit none

    ! HIP
    interface
        function hipMalloc(ptr, size) &
                bind(c, name = 'hipMalloc')
            use iso_c_binding
            implicit none
            integer(c_int) :: hipMalloc
            type(c_ptr) :: ptr
            integer(c_size_t), value :: size
        end function hipMalloc

        function hipFree(ptr) &
                bind(c, name = 'hipFree')
            use iso_c_binding
            implicit none
            integer(c_int) :: hipFree
            type(c_ptr), value :: ptr
        end function hipFree

        function hipMemcpy(dst, src, size, kind) &
                bind(c, name = 'hipMemcpy')
            use iso_c_binding
            implicit none
            integer(c_int) :: hipMemcpy
            type(c_ptr), value :: dst
            type(c_ptr), intent(in), value :: src
            integer(c_size_t), value :: size
            integer(c_int), value :: kind
        end function hipMemcpy
    end interface

    integer, parameter :: hipMemcpyHostToDevice = 1
    integer, parameter :: hipMemcpyDeviceToHost = 2

    ! hipSPARSE
    interface
        function hipsparseCreate(handle) &
                bind(c, name = 'hipsparseCreate')
            use iso_c_binding
            implicit none
            integer(c_int) :: hipsparseCreate
            type(c_ptr) :: handle
        end function hipsparseCreate

        function hipsparseDestroy(handle) &
                bind(c, name = 'hipsparseDestroy')
            use iso_c_binding
            implicit none
            integer(c_int) :: hipsparseDestroy
            type(c_ptr), value :: handle
        end function hipsparseDestroy

        function hipsparseCreateMatDescr(descr) &
                bind(c, name = 'hipsparseCreateMatDescr')
            use iso_c_binding
            implicit none
            integer(c_int) :: hipsparseCreateMatDescr
            type(c_ptr) :: descr
        end function hipsparseCreateMatDescr

        function hipsparseDestroyMatDescr(descr) &
                bind(c, name = 'hipsparseDestroyMatDescr')
            use iso_c_binding
            implicit none
            integer(c_int) :: hipsparseDestroyMatDescr
            type(c_ptr), value :: descr
        end function hipsparseDestroyMatDescr

        function hipsparseSbsrmm(handle, dirA, transA, transB, mb, n, kb, nnzb, alpha, descrA, &
                                 bsrSortedValA, bsrSortedRowPtrA, bsrSortedColIndA, blockDim, &
                                 B, ldb, beta, C, ldc) &
                bind(c, name = 'hipsparseSbsrmm')
            use iso_c_binding
            implicit none
            integer(c_int) :: hipsparseSbsrmm
            type(c_ptr), value :: handle
            integer(c_int), value :: dirA
            integer(c_int), value :: transA
            integer(c_int), value :: transB
            integer(c_int), value :: mb
            integer(c_int), value :: n
            integer(c_int), value :: kb
            integer(c_int), value :: nnzb
            type(c_ptr), intent(in), value :: alpha
            type(c_ptr), value :: descrA
            type(c_ptr), intent(in), value :: bsrSortedValA
            type(c_ptr), intent(in), value :: bsrSortedRowPtrA
            type(c_ptr), intent(in), value :: bsrSortedColIndA
            integer(c_int), value :: blockDim
            type(c_ptr), intent(in), value :: B
            integer(c_int), value :: ldb
            type(c_ptr), intent(in), value :: beta
            type(c_ptr), value :: C
            integer(c_int), value :: ldc
        end function hipsparseSbsrmm
    end interface

    integer, parameter :: HIPSPARSE_DIRECTION_ROW = 0
    integer, parameter :: HIPSPARSE_OPERATION_NON_TRANSPOSE = 0

    ! Variables
    type(c_ptr) :: handle
    type(c_ptr) :: descr
    integer :: i, stat
    
    ! Block sparse matrix A (2x3 blocks of size 2x2)
    integer, parameter :: blockDim = 2
    integer, parameter :: mb = 2
    integer, parameter :: kb = 3
    integer, parameter :: nnzb = 4
    integer, parameter :: m = mb * blockDim  ! 4
    integer, parameter :: k = kb * blockDim  ! 6
    
    integer, dimension(mb+1), target :: hbsrRowPtr = (/0, 2, 4/)
    integer, dimension(nnzb), target :: hbsrColInd = (/0, 1, 1, 2/)
    real(c_float), dimension(nnzb*blockDim*blockDim), target :: hbsrVal = (/ &
        1.0, 2.0, 0.0, 4.0, &  ! block (0,0)
        0.0, 3.0, 5.0, 0.0, &  ! block (0,1)
        0.0, 7.0, 1.0, 2.0, &  ! block (1,1)
        8.0, 0.0, 4.0, 1.0 /)  ! block (1,2)
    
    ! Dense matrix B (6x3) - column major
    integer, parameter :: n = 3
    real(c_float), dimension(k*n), target :: hB = (/ &
        1.0, 2.0, 3.0, 4.0, 5.0, 6.0, &
        7.0, 8.0, 9.0, 10.0, 11.0, 12.0, &
        13.0, 14.0, 15.0, 16.0, 17.0, 18.0 /)
    
    ! Result matrix C (4x3)
    real(c_float), dimension(m*n), target :: hC
    
    ! Scalar values
    real(c_float), target :: alpha = 1.0
    real(c_float), target :: beta = 0.0
    
    ! Device pointers
    type(c_ptr) :: dbsrRowPtr
    type(c_ptr) :: dbsrColInd
    type(c_ptr) :: dbsrVal
    type(c_ptr) :: dB
    type(c_ptr) :: dC

    ! Create hipSPARSE handle
    stat = hipsparseCreate(handle)
    if (stat /= 0) then
        write(*,*) 'Error: hipsparseCreate failed'
        stop
    end if

    ! Create matrix descriptor
    stat = hipsparseCreateMatDescr(descr)
    if (stat /= 0) then
        write(*,*) 'Error: hipsparseCreateMatDescr failed'
        stop
    end if

    ! Allocate device memory for BSR matrix A
    stat = hipMalloc(dbsrRowPtr, int((mb + 1) * 4, c_size_t))
    if (stat /= 0) then
        write(*,*) 'Error: hipMalloc dbsrRowPtr failed'
        stop
    end if

    stat = hipMalloc(dbsrColInd, int(nnzb * 4, c_size_t))
    if (stat /= 0) then
        write(*,*) 'Error: hipMalloc dbsrColInd failed'
        stop
    end if

    stat = hipMalloc(dbsrVal, int(nnzb * blockDim * blockDim * 4, c_size_t))
    if (stat /= 0) then
        write(*,*) 'Error: hipMalloc dbsrVal failed'
        stop
    end if

    ! Allocate device memory for dense matrices B and C
    stat = hipMalloc(dB, int(k * n * 4, c_size_t))
    if (stat /= 0) then
        write(*,*) 'Error: hipMalloc dB failed'
        stop
    end if

    stat = hipMalloc(dC, int(m * n * 4, c_size_t))
    if (stat /= 0) then
        write(*,*) 'Error: hipMalloc dC failed'
        stop
    end if

    ! Copy data to device
    stat = hipMemcpy(dbsrRowPtr, c_loc(hbsrRowPtr), int((mb + 1) * 4, c_size_t), hipMemcpyHostToDevice)
    if (stat /= 0) then
        write(*,*) 'Error: hipMemcpy dbsrRowPtr failed'
        stop
    end if

    stat = hipMemcpy(dbsrColInd, c_loc(hbsrColInd), int(nnzb * 4, c_size_t), hipMemcpyHostToDevice)
    if (stat /= 0) then
        write(*,*) 'Error: hipMemcpy dbsrColInd failed'
        stop
    end if

    stat = hipMemcpy(dbsrVal, c_loc(hbsrVal), int(nnzb * blockDim * blockDim * 4, c_size_t), hipMemcpyHostToDevice)
    if (stat /= 0) then
        write(*,*) 'Error: hipMemcpy dbsrVal failed'
        stop
    end if

    stat = hipMemcpy(dB, c_loc(hB), int(k * n * 4, c_size_t), hipMemcpyHostToDevice)
    if (stat /= 0) then
        write(*,*) 'Error: hipMemcpy dB failed'
        stop
    end if

    ! Perform block sparse matrix-matrix multiplication: C = alpha * A * B + beta * C
    stat = hipsparseSbsrmm(handle, &
                           HIPSPARSE_DIRECTION_ROW, &
                           HIPSPARSE_OPERATION_NON_TRANSPOSE, &
                           HIPSPARSE_OPERATION_NON_TRANSPOSE, &
                           mb, &
                           n, &
                           kb, &
                           nnzb, &
                           c_loc(alpha), &
                           descr, &
                           dbsrVal, &
                           dbsrRowPtr, &
                           dbsrColInd, &
                           blockDim, &
                           dB, &
                           k, &
                           c_loc(beta), &
                           dC, &
                           m)
    if (stat /= 0) then
        write(*,*) 'Error: hipsparseSbsrmm failed'
        stop
    end if

    ! Copy result back to host
    stat = hipMemcpy(c_loc(hC), dC, int(m * n * 4, c_size_t), hipMemcpyDeviceToHost)
    if (stat /= 0) then
        write(*,*) 'Error: hipMemcpy hC failed'
        stop
    end if

    ! Print result
    write(*,*) 'hC:'
    do i = 1, m * n
        write(*,*) hC(i)
    end do

    ! Clean up
    stat = hipFree(dbsrRowPtr)
    stat = hipFree(dbsrColInd)
    stat = hipFree(dbsrVal)
    stat = hipFree(dB)
    stat = hipFree(dC)

    stat = hipsparseDestroyMatDescr(descr)
    stat = hipsparseDestroy(handle)

end program example_hipsparse_bsrmm

hipsparseXcsrmm()#

hipsparseStatus_t hipsparseScsrmm(hipsparseHandle_t handle, hipsparseOperation_t transA, int m, int n, int k, int nnz, const float *alpha, const hipsparseMatDescr_t descrA, const float *csrSortedValA, const int *csrSortedRowPtrA, const int *csrSortedColIndA, const float *B, int ldb, const float *beta, float *C, int ldc)#

hipsparseStatus_t hipsparseDcsrmm(hipsparseHandle_t handle, hipsparseOperation_t transA, int m, int n, int k, int nnz, const double *alpha, const hipsparseMatDescr_t descrA, const double *csrSortedValA, const int *csrSortedRowPtrA, const int *csrSortedColIndA, const double *B, int ldb, const double *beta, double *C, int ldc)#

hipsparseStatus_t hipsparseCcsrmm(hipsparseHandle_t handle, hipsparseOperation_t transA, int m, int n, int k, int nnz, const hipComplex *alpha, const hipsparseMatDescr_t descrA, const hipComplex *csrSortedValA, const int *csrSortedRowPtrA, const int *csrSortedColIndA, const hipComplex *B, int ldb, const hipComplex *beta, hipComplex *C, int ldc)#

hipsparseStatus_t hipsparseZcsrmm(hipsparseHandle_t handle, hipsparseOperation_t transA, int m, int n, int k, int nnz, const hipDoubleComplex *alpha, const hipsparseMatDescr_t descrA, const hipDoubleComplex *csrSortedValA, const int *csrSortedRowPtrA, const int *csrSortedColIndA, const hipDoubleComplex *B, int ldb, const hipDoubleComplex *beta, hipDoubleComplex *C, int ldc)#

Sparse matrix dense matrix multiplication using CSR storage format.

hipsparseXcsrmm multiplies the scalar \(\alpha\) with a sparse \(m \times k\) matrix \(A\), defined in CSR storage format, and the column-oriented dense \(k \times n\) matrix \(B\) and adds the result to the column-oriented dense \(m \times n\) matrix \(C\) that is multiplied by the scalar \(\beta\), such that

\[ C := \alpha \cdot op(A) \cdot B + \beta \cdot C, \]

with

\[\begin{split} op(A) = \left\{ \begin{array}{ll} A, & \text{if transA == HIPSPARSE_OPERATION_NON_TRANSPOSE} \\ A^T, & \text{if transA == HIPSPARSE_OPERATION_TRANSPOSE} \\ A^H, & \text{if transA == HIPSPARSE_OPERATION_CONJUGATE_TRANSPOSE} \end{array} \right. \end{split}\]

for(i = 0; i < ldc; ++i)
{
    for(j = 0; j < n; ++j)
    {
        C[i][j] = beta * C[i][j];

        for(k = csrRowPtr[i]; k < csrRowPtr[i + 1]; ++k)
        {
            C[i][j] += alpha * csrVal[k] * B[csrColInd[k]][j];
        }
    }
}

Deprecated:: This function is deprecated when using the CUDA backend (CUDA 10.0+) and will be removed in CUDA 11.0. This deprecation does not apply to the ROCm backend.

Note

This function is non blocking and executed asynchronously with respect to the host. It may return before the actual computation has finished.

Parameters:

handle – [in] handle to the hipsparse library context queue.
transA – [in] matrix \(A\) operation type.
m – [in] number of rows of the sparse CSR matrix \(A\). Must be non-negative.
n – [in] number of columns of the dense matrix \(op(B)\) and \(C\). Must be non-negative.
k – [in] number of columns of the sparse CSR matrix \(A\). Must be non-negative.
nnz – [in] number of non-zero entries of the sparse CSR matrix \(A\). Must be non-negative.
alpha – [in] scalar \(\alpha\).
descrA – [in] descriptor of the sparse CSR matrix \(A\). Currently, only HIPSPARSE_MATRIX_TYPE_GENERAL is supported.
csrSortedValA – [in] array of nnz elements of the sparse CSR matrix \(A\).
csrSortedRowPtrA – [in] array of m+1 elements that point to the start of every row of the sparse CSR matrix \(A\).
csrSortedColIndA – [in] array of nnz elements containing the column indices of the sparse CSR matrix \(A\).
B – [in] array of dimension ldb*n ( \(op(B) == B\)), ldb*k otherwise.
ldb – [in] leading dimension of \(B\), must be at least \(\max{(1, k)}\) ( \(op(B) == B\)), \(\max{(1, n)}\) otherwise.
beta – [in] scalar \(\beta\).
C – [inout] array of dimension ldc*n.
ldc – [in] leading dimension of \(C\), must be at least \(\max{(1, m)}\) ( \(op(A) == A\)), \(\max{(1, k)}\) otherwise.

Return values:

HIPSPARSE_STATUS_SUCCESS – the operation completed successfully.
HIPSPARSE_STATUS_NOT_INITIALIZED – handle is not initialized.
HIPSPARSE_STATUS_INVALID_VALUE – handle, descrA, alpha or beta is nullptr, m, n, k or nnz is negative, ldb or ldc is invalid, or csrSortedValA, csrSortedRowPtrA, csrSortedColIndA, B or C is nullptr.
HIPSPARSE_STATUS_ARCH_MISMATCH – the device is not supported.
HIPSPARSE_STATUS_NOT_SUPPORTED – hipsparseMatrixType_t is not HIPSPARSE_MATRIX_TYPE_GENERAL.

int main(int argc, char* argv[])
{
    // hipSPARSE handle
    hipsparseHandle_t handle;
    HIPSPARSE_CHECK(hipsparseCreate(&handle));

    //     1 2 0 3 0 0
    // A = 0 4 5 0 0 0
    //     0 0 0 7 8 0
    //     0 0 1 2 4 1

    const int                  m   = 4;
    const int                  k   = 6;
    const int                  nnz = 11;
    const hipsparseDirection_t dir = HIPSPARSE_DIRECTION_ROW;

    std::vector<int>   hcsrRowPtr = {0, 3, 5, 7, 11};
    std::vector<int>   hcsrColInd = {0, 1, 3, 1, 2, 3, 4, 2, 3, 4, 5};
    std::vector<float> hcsrVal    = {1, 2, 3, 4, 5, 7, 8, 1, 2, 4, 1};

    // Set dimension n of B
    const int n = 3;

    // Allocate and generate dense matrix B (k x n)
    std::vector<float> hB = {1.0f,
                             2.0f,
                             3.0f,
                             4.0f,
                             5.0f,
                             6.0f,
                             7.0f,
                             8.0f,
                             9.0f,
                             10.0f,
                             11.0f,
                             12.0f,
                             13.0f,
                             14.0f,
                             15.0f,
                             16.0f,
                             17.0f,
                             18.0f};

    int*   dcsrRowPtr = NULL;
    int*   dcsrColInd = NULL;
    float* dcsrVal    = NULL;
    HIP_CHECK(hipMalloc((void**)&dcsrRowPtr, sizeof(int) * (m + 1)));
    HIP_CHECK(hipMalloc((void**)&dcsrColInd, sizeof(int) * nnz));
    HIP_CHECK(hipMalloc((void**)&dcsrVal, sizeof(float) * nnz));
    HIP_CHECK(
        hipMemcpy(dcsrRowPtr, hcsrRowPtr.data(), sizeof(int) * (m + 1), hipMemcpyHostToDevice));
    HIP_CHECK(hipMemcpy(dcsrColInd, hcsrColInd.data(), sizeof(int) * nnz, hipMemcpyHostToDevice));
    HIP_CHECK(hipMemcpy(dcsrVal, hcsrVal.data(), sizeof(float) * nnz, hipMemcpyHostToDevice));

    // Copy B to the device
    float* dB;
    HIP_CHECK(hipMalloc((void**)&dB, sizeof(float) * k * n));
    HIP_CHECK(hipMemcpy(dB, hB.data(), sizeof(float) * k * n, hipMemcpyHostToDevice));

    // alpha and beta
    float alpha = 1.0f;
    float beta  = 0.0f;

    // Allocate memory for the resulting matrix C
    float* dC;
    HIP_CHECK(hipMalloc((void**)&dC, sizeof(float) * m * n));

    // Matrix descriptor
    hipsparseMatDescr_t descr;
    HIPSPARSE_CHECK(hipsparseCreateMatDescr(&descr));

    // Perform the matrix multiplication
    HIPSPARSE_CHECK(hipsparseScsrmm(handle,
                                    HIPSPARSE_OPERATION_NON_TRANSPOSE,
                                    m,
                                    n,
                                    k,
                                    nnz,
                                    &alpha,
                                    descr,
                                    dcsrVal,
                                    dcsrRowPtr,
                                    dcsrColInd,
                                    dB,
                                    k,
                                    &beta,
                                    dC,
                                    m));

    // Copy results to host
    std::vector<float> hC(6 * 3);
    HIP_CHECK(hipMemcpy(hC.data(), dC, sizeof(float) * m * n, hipMemcpyDeviceToHost));

    std::cout << "hC" << std::endl;
    for(int i = 0; i < m * n; i++)
    {
        std::cout << hC[i] << " ";
    }
    std::cout << std::endl;

    HIP_CHECK(hipFree(dcsrRowPtr));
    HIP_CHECK(hipFree(dcsrColInd));
    HIP_CHECK(hipFree(dcsrVal));
    HIP_CHECK(hipFree(dB));
    HIP_CHECK(hipFree(dC));

    HIPSPARSE_CHECK(hipsparseDestroy(handle));

    return 0;
}

int main(int argc, char* argv[])
{
    // hipSPARSE handle
    hipsparseHandle_t handle;
    HIPSPARSE_CHECK(hipsparseCreate(&handle));

    //     1 2 0 3 0 0
    // A = 0 4 5 0 0 0
    //     0 0 0 7 8 0
    //     0 0 1 2 4 1

    const int                  m   = 4;
    const int                  k   = 6;
    const int                  nnz = 11;
    const hipsparseDirection_t dir = HIPSPARSE_DIRECTION_ROW;

    int   hcsrRowPtr[] = {0, 3, 5, 7, 11};
    int   hcsrColInd[] = {0, 1, 3, 1, 2, 3, 4, 2, 3, 4, 5};
    float hcsrVal[]    = {1, 2, 3, 4, 5, 7, 8, 1, 2, 4, 1};

    // Set dimension n of B
    const int n = 3;

    // Allocate and generate dense matrix B (k x n)
    float hB[] = {1.0,
                  2.0,
                  3.0,
                  4.0,
                  5.0,
                  6.0,
                  7.0,
                  8.0,
                  9.0,
                  10.0,
                  11.0,
                  12.0,
                  13.0,
                  14.0,
                  15.0,
                  16.0,
                  17.0,
                  18.0};

    int*   dcsrRowPtr = NULL;
    int*   dcsrColInd = NULL;
    float* dcsrVal    = NULL;
    HIP_CHECK(hipMalloc((void**)&dcsrRowPtr, sizeof(int) * (m + 1)));
    HIP_CHECK(hipMalloc((void**)&dcsrColInd, sizeof(int) * nnz));
    HIP_CHECK(hipMalloc((void**)&dcsrVal, sizeof(float) * nnz));
    HIP_CHECK(hipMemcpy(dcsrRowPtr, hcsrRowPtr, sizeof(int) * (m + 1), hipMemcpyHostToDevice));
    HIP_CHECK(hipMemcpy(dcsrColInd, hcsrColInd, sizeof(int) * nnz, hipMemcpyHostToDevice));
    HIP_CHECK(hipMemcpy(dcsrVal, hcsrVal, sizeof(float) * nnz, hipMemcpyHostToDevice));

    // Copy B to the device
    float* dB;
    HIP_CHECK(hipMalloc((void**)&dB, sizeof(float) * k * n));
    HIP_CHECK(hipMemcpy(dB, hB, sizeof(float) * k * n, hipMemcpyHostToDevice));

    // alpha and beta
    float alpha = 1.0;
    float beta  = 0.0;

    // Allocate memory for the resulting matrix C
    float* dC;
    HIP_CHECK(hipMalloc((void**)&dC, sizeof(float) * m * n));

    // Matrix descriptor
    hipsparseMatDescr_t descr;
    HIPSPARSE_CHECK(hipsparseCreateMatDescr(&descr));

    // Perform the matrix multiplication
    HIPSPARSE_CHECK(hipsparseScsrmm(handle,
                                    HIPSPARSE_OPERATION_NON_TRANSPOSE,
                                    m,
                                    n,
                                    k,
                                    nnz,
                                    &alpha,
                                    descr,
                                    dcsrVal,
                                    dcsrRowPtr,
                                    dcsrColInd,
                                    dB,
                                    k,
                                    &beta,
                                    dC,
                                    m));

    // Copy results to host
    float hC[6 * 3];
    HIP_CHECK(hipMemcpy(hC, dC, sizeof(float) * m * n, hipMemcpyDeviceToHost));

    printf("hC\n");
    for(int i = 0; i < m * n; i++)
    {
        printf("%f ", hC[i]);
    }
    printf("\n");

    HIP_CHECK(hipFree(dcsrRowPtr));
    HIP_CHECK(hipFree(dcsrColInd));
    HIP_CHECK(hipFree(dcsrVal));
    HIP_CHECK(hipFree(dB));
    HIP_CHECK(hipFree(dC));

    HIPSPARSE_CHECK(hipsparseDestroy(handle));

    return 0;
}

program example_hipsparse_csrmm
    use iso_c_binding
    implicit none

    ! HIP
    interface
        function hipMalloc(ptr, size) &
                bind(c, name = 'hipMalloc')
            use iso_c_binding
            implicit none
            integer(c_int) :: hipMalloc
            type(c_ptr) :: ptr
            integer(c_size_t), value :: size
        end function hipMalloc

        function hipFree(ptr) &
                bind(c, name = 'hipFree')
            use iso_c_binding
            implicit none
            integer(c_int) :: hipFree
            type(c_ptr), value :: ptr
        end function hipFree

        function hipMemcpy(dst, src, size, kind) &
                bind(c, name = 'hipMemcpy')
            use iso_c_binding
            implicit none
            integer(c_int) :: hipMemcpy
            type(c_ptr), value :: dst
            type(c_ptr), intent(in), value :: src
            integer(c_size_t), value :: size
            integer(c_int), value :: kind
        end function hipMemcpy

        function hipMemset(dst, val, size) &
                bind(c, name = 'hipMemset')
            use iso_c_binding
            implicit none
            integer(c_int) :: hipMemset
            type(c_ptr), value :: dst
            integer(c_int), value :: val
            integer(c_size_t), value :: size
        end function hipMemset

        function hipDeviceSynchronize() &
                bind(c, name = 'hipDeviceSynchronize')
            use iso_c_binding
            implicit none
            integer(c_int) :: hipDeviceSynchronize
        end function hipDeviceSynchronize
    end interface

    integer, parameter :: hipMemcpyHostToDevice = 1
    integer, parameter :: hipMemcpyDeviceToHost = 2

    ! hipSPARSE
    interface
        function hipsparseCreate(handle) &
                bind(c, name = 'hipsparseCreate')
            use iso_c_binding
            implicit none
            integer(c_int) :: hipsparseCreate
            type(c_ptr) :: handle
        end function hipsparseCreate

        function hipsparseDestroy(handle) &
                bind(c, name = 'hipsparseDestroy')
            use iso_c_binding
            implicit none
            integer(c_int) :: hipsparseDestroy
            type(c_ptr), value :: handle
        end function hipsparseDestroy

        function hipsparseCreateMatDescr(descr) &
                bind(c, name = 'hipsparseCreateMatDescr')
            use iso_c_binding
            implicit none
            integer(c_int) :: hipsparseCreateMatDescr
            type(c_ptr) :: descr
        end function hipsparseCreateMatDescr

        function hipsparseDestroyMatDescr(descr) &
                bind(c, name = 'hipsparseDestroyMatDescr')
            use iso_c_binding
            implicit none
            integer(c_int) :: hipsparseDestroyMatDescr
            type(c_ptr), value :: descr
        end function hipsparseDestroyMatDescr

        function hipsparseScsrmm(handle, transA, m, n, k, nnz, alpha, descrA, csrSortedValA, &
                                 csrSortedRowPtrA, csrSortedColIndA, B, ldb, beta, C, ldc) &
                bind(c, name = 'hipsparseScsrmm')
            use iso_c_binding
            implicit none
            integer(c_int) :: hipsparseScsrmm
            type(c_ptr), value :: handle
            integer(c_int), value :: transA
            integer(c_int), value :: m
            integer(c_int), value :: n
            integer(c_int), value :: k
            integer(c_int), value :: nnz
            type(c_ptr), intent(in), value :: alpha
            type(c_ptr), value :: descrA
            type(c_ptr), intent(in), value :: csrSortedValA
            type(c_ptr), intent(in), value :: csrSortedRowPtrA
            type(c_ptr), intent(in), value :: csrSortedColIndA
            type(c_ptr), intent(in), value :: B
            integer(c_int), value :: ldb
            type(c_ptr), intent(in), value :: beta
            type(c_ptr), value :: C
            integer(c_int), value :: ldc
        end function hipsparseScsrmm
    end interface

    integer, parameter :: HIPSPARSE_OPERATION_NON_TRANSPOSE = 0

    ! Variables
    type(c_ptr) :: handle
    type(c_ptr) :: descr
    integer :: i, j, stat
    
    ! Matrix A (4x6) in CSR format
    !     1 2 0 3 0 0
    ! A = 0 4 5 0 0 0
    !     0 0 0 7 8 0
    !     0 0 1 2 4 1
    integer, parameter :: m = 4
    integer, parameter :: k = 6
    integer, parameter :: nnz = 11
    
    integer, dimension(m+1), target :: hcsrRowPtr = (/0, 3, 5, 7, 11/)
    integer, dimension(nnz), target :: hcsrColInd = (/0, 1, 3, 1, 2, 3, 4, 2, 3, 4, 5/)
    real(c_float), dimension(nnz), target :: hcsrVal = (/1.0, 2.0, 3.0, 4.0, 5.0, 7.0, 8.0, 1.0, 2.0, 4.0, 1.0/)
    
    ! Dense matrix B (6x3) - column major
    integer, parameter :: n = 3
    real(c_float), dimension(k*n), target :: hB = (/ &
        1.0, 2.0, 3.0, 4.0, 5.0, 6.0, &
        7.0, 8.0, 9.0, 10.0, 11.0, 12.0, &
        13.0, 14.0, 15.0, 16.0, 17.0, 18.0 /)
    
    ! Result matrix C (4x3)
    real(c_float), dimension(m*n), target :: hC
    
    ! Scalar values
    real(c_float), target :: alpha = 1.0
    real(c_float), target :: beta = 0.0
    
    ! Device pointers
    type(c_ptr) :: dcsrRowPtr
    type(c_ptr) :: dcsrColInd
    type(c_ptr) :: dcsrVal
    type(c_ptr) :: dB
    type(c_ptr) :: dC

    ! Create hipSPARSE handle
    stat = hipsparseCreate(handle)
    if (stat /= 0) then
        write(*,*) 'Error: hipsparseCreate failed'
        stop
    end if

    ! Create matrix descriptor
    stat = hipsparseCreateMatDescr(descr)
    if (stat /= 0) then
        write(*,*) 'Error: hipsparseCreateMatDescr failed'
        stop
    end if

    ! Allocate device memory for CSR matrix A
    stat = hipMalloc(dcsrRowPtr, int((m + 1) * 4, c_size_t))
    if (stat /= 0) then
        write(*,*) 'Error: hipMalloc dcsrRowPtr failed'
        stop
    end if

    stat = hipMalloc(dcsrColInd, int(nnz * 4, c_size_t))
    if (stat /= 0) then
        write(*,*) 'Error: hipMalloc dcsrColInd failed'
        stop
    end if

    stat = hipMalloc(dcsrVal, int(nnz * 4, c_size_t))
    if (stat /= 0) then
        write(*,*) 'Error: hipMalloc dcsrVal failed'
        stop
    end if

    ! Allocate device memory for dense matrices B and C
    stat = hipMalloc(dB, int(k * n * 4, c_size_t))
    if (stat /= 0) then
        write(*,*) 'Error: hipMalloc dB failed'
        stop
    end if

    stat = hipMalloc(dC, int(m * n * 4, c_size_t))
    if (stat /= 0) then
        write(*,*) 'Error: hipMalloc dC failed'
        stop
    end if

    ! Copy data to device
    stat = hipMemcpy(dcsrRowPtr, c_loc(hcsrRowPtr), int((m + 1) * 4, c_size_t), hipMemcpyHostToDevice)
    if (stat /= 0) then
        write(*,*) 'Error: hipMemcpy dcsrRowPtr failed'
        stop
    end if

    stat = hipMemcpy(dcsrColInd, c_loc(hcsrColInd), int(nnz * 4, c_size_t), hipMemcpyHostToDevice)
    if (stat /= 0) then
        write(*,*) 'Error: hipMemcpy dcsrColInd failed'
        stop
    end if

    stat = hipMemcpy(dcsrVal, c_loc(hcsrVal), int(nnz * 4, c_size_t), hipMemcpyHostToDevice)
    if (stat /= 0) then
        write(*,*) 'Error: hipMemcpy dcsrVal failed'
        stop
    end if

    stat = hipMemcpy(dB, c_loc(hB), int(k * n * 4, c_size_t), hipMemcpyHostToDevice)
    if (stat /= 0) then
        write(*,*) 'Error: hipMemcpy dB failed'
        stop
    end if

    ! Perform matrix-matrix multiplication: C = alpha * A * B + beta * C
    stat = hipsparseScsrmm(handle, &
                           HIPSPARSE_OPERATION_NON_TRANSPOSE, &
                           m, &
                           n, &
                           k, &
                           nnz, &
                           c_loc(alpha), &
                           descr, &
                           dcsrVal, &
                           dcsrRowPtr, &
                           dcsrColInd, &
                           dB, &
                           k, &
                           c_loc(beta), &
                           dC, &
                           m)
    if (stat /= 0) then
        write(*,*) 'Error: hipsparseScsrmm failed'
        stop
    end if

    ! Copy result back to host
    stat = hipMemcpy(c_loc(hC), dC, int(m * n * 4, c_size_t), hipMemcpyDeviceToHost)
    if (stat /= 0) then
        write(*,*) 'Error: hipMemcpy hC failed'
        stop
    end if

    ! Print result
    write(*,*) 'hC:'
    do i = 1, m * n
        write(*,*) hC(i)
    end do

    ! Clean up
    stat = hipFree(dcsrRowPtr)
    stat = hipFree(dcsrColInd)
    stat = hipFree(dcsrVal)
    stat = hipFree(dB)
    stat = hipFree(dC)

    stat = hipsparseDestroyMatDescr(descr)
    stat = hipsparseDestroy(handle)

end program example_hipsparse_csrmm

hipsparseXcsrmm2()#

hipsparseStatus_t hipsparseScsrmm2(hipsparseHandle_t handle, hipsparseOperation_t transA, hipsparseOperation_t transB, int m, int n, int k, int nnz, const float *alpha, const hipsparseMatDescr_t descrA, const float *csrSortedValA, const int *csrSortedRowPtrA, const int *csrSortedColIndA, const float *B, int ldb, const float *beta, float *C, int ldc)#

hipsparseStatus_t hipsparseDcsrmm2(hipsparseHandle_t handle, hipsparseOperation_t transA, hipsparseOperation_t transB, int m, int n, int k, int nnz, const double *alpha, const hipsparseMatDescr_t descrA, const double *csrSortedValA, const int *csrSortedRowPtrA, const int *csrSortedColIndA, const double *B, int ldb, const double *beta, double *C, int ldc)#

hipsparseStatus_t hipsparseCcsrmm2(hipsparseHandle_t handle, hipsparseOperation_t transA, hipsparseOperation_t transB, int m, int n, int k, int nnz, const hipComplex *alpha, const hipsparseMatDescr_t descrA, const hipComplex *csrSortedValA, const int *csrSortedRowPtrA, const int *csrSortedColIndA, const hipComplex *B, int ldb, const hipComplex *beta, hipComplex *C, int ldc)#

hipsparseStatus_t hipsparseZcsrmm2(hipsparseHandle_t handle, hipsparseOperation_t transA, hipsparseOperation_t transB, int m, int n, int k, int nnz, const hipDoubleComplex *alpha, const hipsparseMatDescr_t descrA, const hipDoubleComplex *csrSortedValA, const int *csrSortedRowPtrA, const int *csrSortedColIndA, const hipDoubleComplex *B, int ldb, const hipDoubleComplex *beta, hipDoubleComplex *C, int ldc)#

Sparse matrix dense matrix multiplication using CSR storage format.

hipsparseXcsrmm2 multiplies the scalar \(\alpha\) with a sparse \(m \times k\) matrix \(A\), defined in CSR storage format, and the column-oriented dense \(k \times n\) matrix \(B\) and adds the result to the column-oriented dense \(m \times n\) matrix \(C\) that is multiplied by the scalar \(\beta\), such that

\[ C := \alpha \cdot op(A) \cdot op(B) + \beta \cdot C, \]

with

\[\begin{split} op(A) = \left\{ \begin{array}{ll} A, & \text{if transA == HIPSPARSE_OPERATION_NON_TRANSPOSE} \\ A^T, & \text{if transA == HIPSPARSE_OPERATION_TRANSPOSE} \\ A^H, & \text{if transA == HIPSPARSE_OPERATION_CONJUGATE_TRANSPOSE} \end{array} \right. \end{split}\]

and

\[\begin{split} op(B) = \left\{ \begin{array}{ll} B, & \text{if transB == HIPSPARSE_OPERATION_NON_TRANSPOSE} \\ B^T, & \text{if transB == HIPSPARSE_OPERATION_TRANSPOSE} \\ B^H, & \text{if transB == HIPSPARSE_OPERATION_CONJUGATE_TRANSPOSE} \end{array} \right. \end{split}\]

for(i = 0; i < ldc; ++i)
{
    for(j = 0; j < n; ++j)
    {
        C[i][j] = beta * C[i][j];

        for(k = csrRowPtr[i]; k < csrRowPtr[i + 1]; ++k)
        {
            C[i][j] += alpha * csrVal[k] * B[csrColInd[k]][j];
        }
    }
}

Note

This function is non blocking and executed asynchronously with respect to the host. It may return before the actual computation has finished.

Parameters:

handle – [in] handle to the hipsparse library context queue.
transA – [in] matrix \(A\) operation type.
transB – [in] matrix \(B\) operation type.
m – [in] number of rows of the sparse CSR matrix \(A\).
n – [in] number of columns of the dense matrix \(op(B)\) and \(C\).
k – [in] number of columns of the sparse CSR matrix \(A\).
nnz – [in] number of non-zero entries of the sparse CSR matrix \(A\).
alpha – [in] scalar \(\alpha\).
descrA – [in] descriptor of the sparse CSR matrix \(A\). Currently, only HIPSPARSE_MATRIX_TYPE_GENERAL is supported.
csrSortedValA – [in] array of nnz elements of the sparse CSR matrix \(A\).
csrSortedRowPtrA – [in] array of m+1 elements that point to the start of every row of the sparse CSR matrix \(A\).
csrSortedColIndA – [in] array of nnz elements containing the column indices of the sparse CSR matrix \(A\).
B – [in] array of dimension ldb*n ( \(op(B) == B\)), ldb*k otherwise.
ldb – [in] leading dimension of \(B\), must be at least \(\max{(1, k)}\) ( \(op(B) == B\)), \(\max{(1, n)}\) otherwise.
beta – [in] scalar \(\beta\).
C – [inout] array of dimension ldc*n.
ldc – [in] leading dimension of \(C\), must be at least \(\max{(1, m)}\) ( \(op(A) == A\)), \(\max{(1, k)}\) otherwise.

Return values:

HIPSPARSE_STATUS_SUCCESS – the operation completed successfully.
HIPSPARSE_STATUS_INVALID_VALUE – handle, m, n, k, nnz, ldb, ldc descrA, alpha, csrSortedValA, csrSortedRowPtrA, csrSortedColIndA, B, beta or C is invalid.
HIPSPARSE_STATUS_ARCH_MISMATCH – the device is not supported.
HIPSPARSE_STATUS_NOT_SUPPORTED – hipsparseMatrixType_t != HIPSPARSE_MATRIX_TYPE_GENERAL.

hipsparseXbsrsm2_zeroPivot()#

hipsparseStatus_t hipsparseXbsrsm2_zeroPivot(hipsparseHandle_t handle, bsrsm2Info_t info, int *position)#

hipsparseXbsrsm2_zeroPivot returns HIPSPARSE_STATUS_ZERO_PIVOT, if either a structural or numerical zero has been found during hipsparseXbsrsm2_analysis() or hipsparseXbsrsm2_solve() computation. The first zero pivot \(j\) at \(A_{j,j}\) is stored in position, using same index base as the BSR matrix.

position can be in host or device memory. If no zero pivot has been found, position is set to -1 and HIPSPARSE_STATUS_SUCCESS is returned instead.

Deprecated:: This function is deprecated when using the CUDA backend (CUDA 12.0+) and will be removed in CUDA 13.0. This deprecation does not apply to the ROCm backend.

Note

hipsparseXbsrsm2_zeroPivot is a blocking function. It might influence performance negatively.

Parameters:

handle – [in] handle to the hipsparse library context queue.
info – [in] structure that holds the information collected during the analysis step.
position – [inout] pointer to zero pivot \(j\), can be in host or device memory.

Return values:

HIPSPARSE_STATUS_SUCCESS – the operation completed successfully.
HIPSPARSE_STATUS_NOT_INITIALIZED – handle is not initialized.
HIPSPARSE_STATUS_INVALID_VALUE – handle, info or position is nullptr.
HIPSPARSE_STATUS_INTERNAL_ERROR – an internal error occurred.
HIPSPARSE_STATUS_ZERO_PIVOT – zero pivot has been found.

hipsparseXbsrsm2_bufferSize()#

hipsparseStatus_t hipsparseSbsrsm2_bufferSize(hipsparseHandle_t handle, hipsparseDirection_t dirA, hipsparseOperation_t transA, hipsparseOperation_t transX, int mb, int nrhs, int nnzb, const hipsparseMatDescr_t descrA, float *bsrSortedValA, const int *bsrSortedRowPtrA, const int *bsrSortedColIndA, int blockDim, bsrsm2Info_t info, int *pBufferSizeInBytes)#

hipsparseStatus_t hipsparseDbsrsm2_bufferSize(hipsparseHandle_t handle, hipsparseDirection_t dirA, hipsparseOperation_t transA, hipsparseOperation_t transX, int mb, int nrhs, int nnzb, const hipsparseMatDescr_t descrA, double *bsrSortedValA, const int *bsrSortedRowPtrA, const int *bsrSortedColIndA, int blockDim, bsrsm2Info_t info, int *pBufferSizeInBytes)#

hipsparseStatus_t hipsparseCbsrsm2_bufferSize(hipsparseHandle_t handle, hipsparseDirection_t dirA, hipsparseOperation_t transA, hipsparseOperation_t transX, int mb, int nrhs, int nnzb, const hipsparseMatDescr_t descrA, hipComplex *bsrSortedValA, const int *bsrSortedRowPtrA, const int *bsrSortedColIndA, int blockDim, bsrsm2Info_t info, int *pBufferSizeInBytes)#

hipsparseStatus_t hipsparseZbsrsm2_bufferSize(hipsparseHandle_t handle, hipsparseDirection_t dirA, hipsparseOperation_t transA, hipsparseOperation_t transX, int mb, int nrhs, int nnzb, const hipsparseMatDescr_t descrA, hipDoubleComplex *bsrSortedValA, const int *bsrSortedRowPtrA, const int *bsrSortedColIndA, int blockDim, bsrsm2Info_t info, int *pBufferSizeInBytes)#

hipsparseXbsrsm2_buffer_size returns the size of the temporary storage buffer in bytes that is required by hipsparseXbsrsm2_analysis() and hipsparseXbsrsm2_solve(). The temporary storage buffer must be allocated by the user.

Parameters:

handle – [in] handle to the hipsparse library context queue.
dirA – [in] matrix storage of BSR blocks.
transA – [in] matrix \(A\) operation type.
transX – [in] matrix \(X\) operation type.
mb – [in] number of block rows of the sparse BSR matrix \(A\).
nrhs – [in] number of columns of the dense matrix \(op(X)\).
nnzb – [in] number of non-zero blocks of the sparse BSR matrix \(A\).
descrA – [in] descriptor of the sparse BSR matrix \(A\).
bsrSortedValA – [in] array of nnzb blocks of the sparse BSR matrix.
bsrSortedRowPtrA – [in] array of mb+1 elements that point to the start of every block row of the sparse BSR matrix.
bsrSortedColIndA – [in] array of nnzb containing the block column indices of the sparse BSR matrix.
blockDim – [in] block dimension of the sparse BSR matrix.
info – [in] structure that holds the information collected during the analysis step.
pBufferSizeInBytes – [out] number of bytes of the temporary storage buffer required by hipsparseXbsrsm2_analysis() and hipsparseXbsrsm2_solve().

Return values:

HIPSPARSE_STATUS_SUCCESS – the operation completed successfully.
HIPSPARSE_STATUS_INVALID_VALUE – handle, mb, nrhs, nnzb, blockDim, descrA, bsrSortedValA, bsrSortedRowPtrA, bsrSortedColIndA, info or pBufferSizeInBytes is invalid.
HIPSPARSE_STATUS_INTERNAL_ERROR – an internal error occurred.
HIPSPARSE_STATUS_NOT_SUPPORTED – transA == HIPSPARSE_OPERATION_CONJUGATE_TRANSPOSE, transX == HIPSPARSE_OPERATION_CONJUGATE_TRANSPOSE or hipsparseMatrixType_t != HIPSPARSE_MATRIX_TYPE_GENERAL.

hipsparseXbsrsm2_analysis()#

hipsparseStatus_t hipsparseSbsrsm2_analysis(hipsparseHandle_t handle, hipsparseDirection_t dirA, hipsparseOperation_t transA, hipsparseOperation_t transX, int mb, int nrhs, int nnzb, const hipsparseMatDescr_t descrA, const float *bsrSortedValA, const int *bsrSortedRowPtrA, const int *bsrSortedColIndA, int blockDim, bsrsm2Info_t info, hipsparseSolvePolicy_t policy, void *pBuffer)#

hipsparseStatus_t hipsparseDbsrsm2_analysis(hipsparseHandle_t handle, hipsparseDirection_t dirA, hipsparseOperation_t transA, hipsparseOperation_t transX, int mb, int nrhs, int nnzb, const hipsparseMatDescr_t descrA, const double *bsrSortedValA, const int *bsrSortedRowPtrA, const int *bsrSortedColIndA, int blockDim, bsrsm2Info_t info, hipsparseSolvePolicy_t policy, void *pBuffer)#

hipsparseStatus_t hipsparseCbsrsm2_analysis(hipsparseHandle_t handle, hipsparseDirection_t dirA, hipsparseOperation_t transA, hipsparseOperation_t transX, int mb, int nrhs, int nnzb, const hipsparseMatDescr_t descrA, const hipComplex *bsrSortedValA, const int *bsrSortedRowPtrA, const int *bsrSortedColIndA, int blockDim, bsrsm2Info_t info, hipsparseSolvePolicy_t policy, void *pBuffer)#

hipsparseStatus_t hipsparseZbsrsm2_analysis(hipsparseHandle_t handle, hipsparseDirection_t dirA, hipsparseOperation_t transA, hipsparseOperation_t transX, int mb, int nrhs, int nnzb, const hipsparseMatDescr_t descrA, const hipDoubleComplex *bsrSortedValA, const int *bsrSortedRowPtrA, const int *bsrSortedColIndA, int blockDim, bsrsm2Info_t info, hipsparseSolvePolicy_t policy, void *pBuffer)#

Sparse triangular system solve using BSR storage format.

hipsparseXbsrsm2_analysis performs the analysis step for hipsparseXbsrsm2_solve(). It is expected that this function will be executed only once for a given matrix and particular operation type.

Note

If the matrix sparsity pattern changes, the gathered information will become invalid.

Note

This function is non blocking and executed asynchronously with respect to the host. It may return before the actual computation has finished.

Parameters:

handle – [in] handle to the hipsparse library context queue.
dirA – [in] matrix storage of BSR blocks.
transA – [in] matrix \(A\) operation type.
transX – [in] matrix \(X\) operation type.
mb – [in] number of block rows of the sparse BSR matrix \(A\).
nrhs – [in] number of columns of the dense matrix \(op(X)\).
nnzb – [in] number of non-zero blocks of the sparse BSR matrix \(A\).
descrA – [in] descriptor of the sparse BSR matrix \(A\).
bsrSortedValA – [in] array of nnzb blocks of the sparse BSR matrix \(A\).
bsrSortedRowPtrA – [in] array of mb+1 elements that point to the start of every block row of the sparse BSR matrix \(A\).
bsrSortedColIndA – [in] array of nnzb containing the block column indices of the sparse BSR matrix \(A\).
blockDim – [in] block dimension of the sparse BSR matrix \(A\).
info – [out] structure that holds the information collected during the analysis step.
policy – [in] HIPSPARSE_SOLVE_POLICY_NO_LEVEL or HIPSPARSE_SOLVE_POLICY_USE_LEVEL.
pBuffer – [in] temporary storage buffer allocated by the user.

Return values:

HIPSPARSE_STATUS_SUCCESS – the operation completed successfully.
HIPSPARSE_STATUS_INVALID_VALUE – handle, mb, nrhs, nnzb or blockDim, descrA, bsrSortedValA, bsrSortedRowPtrA, bsrSortedColIndA, info or pBuffer is invalid.
HIPSPARSE_STATUS_INTERNAL_ERROR – an internal error occurred.
HIPSPARSE_STATUS_NOT_SUPPORTED – transA == HIPSPARSE_OPERATION_CONJUGATE_TRANSPOSE, transX == HIPSPARSE_OPERATION_CONJUGATE_TRANSPOSE or hipsparseMatrixType_t != HIPSPARSE_MATRIX_TYPE_GENERAL.

hipsparseXbsrsm2_solve()#

hipsparseStatus_t hipsparseSbsrsm2_solve(hipsparseHandle_t handle, hipsparseDirection_t dirA, hipsparseOperation_t transA, hipsparseOperation_t transX, int mb, int nrhs, int nnzb, const float *alpha, const hipsparseMatDescr_t descrA, const float *bsrSortedValA, const int *bsrSortedRowPtrA, const int *bsrSortedColIndA, int blockDim, bsrsm2Info_t info, const float *B, int ldb, float *X, int ldx, hipsparseSolvePolicy_t policy, void *pBuffer)#

hipsparseStatus_t hipsparseDbsrsm2_solve(hipsparseHandle_t handle, hipsparseDirection_t dirA, hipsparseOperation_t transA, hipsparseOperation_t transX, int mb, int nrhs, int nnzb, const double *alpha, const hipsparseMatDescr_t descrA, const double *bsrSortedValA, const int *bsrSortedRowPtrA, const int *bsrSortedColIndA, int blockDim, bsrsm2Info_t info, const double *B, int ldb, double *X, int ldx, hipsparseSolvePolicy_t policy, void *pBuffer)#

hipsparseStatus_t hipsparseCbsrsm2_solve(hipsparseHandle_t handle, hipsparseDirection_t dirA, hipsparseOperation_t transA, hipsparseOperation_t transX, int mb, int nrhs, int nnzb, const hipComplex *alpha, const hipsparseMatDescr_t descrA, const hipComplex *bsrSortedValA, const int *bsrSortedRowPtrA, const int *bsrSortedColIndA, int blockDim, bsrsm2Info_t info, const hipComplex *B, int ldb, hipComplex *X, int ldx, hipsparseSolvePolicy_t policy, void *pBuffer)#

hipsparseStatus_t hipsparseZbsrsm2_solve(hipsparseHandle_t handle, hipsparseDirection_t dirA, hipsparseOperation_t transA, hipsparseOperation_t transX, int mb, int nrhs, int nnzb, const hipDoubleComplex *alpha, const hipsparseMatDescr_t descrA, const hipDoubleComplex *bsrSortedValA, const int *bsrSortedRowPtrA, const int *bsrSortedColIndA, int blockDim, bsrsm2Info_t info, const hipDoubleComplex *B, int ldb, hipDoubleComplex *X, int ldx, hipsparseSolvePolicy_t policy, void *pBuffer)#

Sparse triangular system solve using BSR storage format.

hipsparseXbsrsm2_solve solves a sparse triangular linear system of a sparse \(m \times m\) matrix, defined in BSR storage format, a column-oriented dense solution matrix \(X\) and the column-oriented dense right-hand side matrix \(B\) that is multiplied by \(\alpha\), such that

\[ op(A) \cdot op(X) = \alpha \cdot op(B), \]

with

\[\begin{split} op(A) = \left\{ \begin{array}{ll} A, & \text{if transA == HIPSPARSE_OPERATION_NON_TRANSPOSE} \\ A^T, & \text{if transA == HIPSPARSE_OPERATION_TRANSPOSE} \\ A^H, & \text{if transA == HIPSPARSE_OPERATION_CONJUGATE_TRANSPOSE} \end{array} \right. \end{split}\]

,

\[\begin{split} op(B) = \left\{ \begin{array}{ll} B, & \text{if transX == HIPSPARSE_OPERATION_NON_TRANSPOSE} \\ B^T, & \text{if transX == HIPSPARSE_OPERATION_TRANSPOSE} \\ B^H, & \text{if transX == HIPSPARSE_OPERATION_CONJUGATE_TRANSPOSE} \end{array} \right. \end{split}\]

and

\[\begin{split} op(X) = \left\{ \begin{array}{ll} X, & \text{if transX == HIPSPARSE_OPERATION_NON_TRANSPOSE} \\ X^T, & \text{if transX == HIPSPARSE_OPERATION_TRANSPOSE} \\ X^H, & \text{if transX == HIPSPARSE_OPERATION_CONJUGATE_TRANSPOSE} \end{array} \right. \end{split}\]

and where \(m = blockDim \times mb\).

Note that as indicated above, the operation type of both \(op(B)\) and \(op(X)\) is specified by the transX parameter and that the operation type of \(B\) and \(X\) must match. For example, if \(op(B)=B\) then \(op(X)=X\). Likewise, if \(op(B)=B^T\) then \(op(X)=X^T\).

Given that the sparse matrix \(A\) is a square matrix, its size is \(m \times m\) regardless of whether \(A\) is transposed or not. The size of the column-oriented dense matrices \(B\) and \(X\) have size that depends on the value of transX:

\[\begin{split} op(B) = \left\{ \begin{array}{ll} ldb \times nrhs, \text{ } ldb \ge m, & \text{if transX == HIPSPARSE_OPERATION_NON_TRANSPOSE} \\ ldb \times m, \text{ } ldb \ge nrhs, & \text{if transX == HIPSPARSE_OPERATION_TRANSPOSE} \\ ldb \times m, \text{ } ldb \ge nrhs, & \text{if transX == HIPSPARSE_OPERATION_CONJUGATE_TRANSPOSE} \end{array} \right. \end{split}\]

and

\[\begin{split} op(X) = \left\{ \begin{array}{ll} ldb \times nrhs, \text{ } ldb \ge m, & \text{if transX == HIPSPARSE_OPERATION_NON_TRANSPOSE} \\ ldb \times m, \text{ } ldb \ge nrhs, & \text{if transX == HIPSPARSE_OPERATION_TRANSPOSE} \\ ldb \times m, \text{ } ldb \ge nrhs, & \text{if transX == HIPSPARSE_OPERATION_CONJUGATE_TRANSPOSE} \end{array} \right. \end{split}\]

hipsparseXbsrsm2_solve requires a user allocated temporary buffer. Its size is returned by hipsparseXbsrsm2_bufferSize(). The size of the required buffer is larger when transA equals HIPSPARSE_OPERATION_TRANSPOSE or HIPSPARSE_OPERATION_CONJUGATE_TRANSPOSE and when transX is HIPSPARSE_OPERATION_NON_TRANSPOSE. The subsequent solve will also be faster when \(A\) is non-transposed and \(B\) is transposed (or conjugate transposed). For example, instead of solving:

\[\begin{split} \left[ \begin{array}{c | c} \begin{array}{c c} a_{00} & a_{01} \\ a_{10} & a_{11} \end{array} & \begin{array}{c c} 0 & 0 \\ 0 & 0 \end{array} \\ \hline \begin{array}{c c} a_{20} & a_{21} \\ a_{30} & a_{31} \end{array} & \begin{array}{c c} a_{22} & a_{23} \\ a_{32} & a_{33} \end{array} \\ \end{array} \right] \cdot \begin{bmatrix} x_{00} & x_{01} \\ x_{10} & x_{11} \\ x_{20} & x_{21} \\ x_{30} & x_{31} \\ \end{bmatrix} = \begin{bmatrix} b_{00} & b_{01} \\ b_{10} & b_{11} \\ b_{20} & b_{21} \\ b_{30} & b_{31} \\ \end{bmatrix} \end{split}\]

Consider solving:

\[\begin{split} \left[ \begin{array}{c | c} \begin{array}{c c} a_{00} & a_{01} \\ a_{10} & a_{11} \end{array} & \begin{array}{c c} 0 & 0 \\ 0 & 0 \end{array} \\ \hline \begin{array}{c c} a_{20} & a_{21} \\ a_{30} & a_{31} \end{array} & \begin{array}{c c} a_{22} & a_{23} \\ a_{32} & a_{33} \end{array} \\ \end{array} \right] \cdot \begin{bmatrix} x_{00} & x_{10} & x_{20} & x_{30} \\ x_{01} & x_{11} & x_{21} & x_{31} \end{bmatrix}^{T} = \begin{bmatrix} b_{00} & b_{10} & b_{20} & b_{30} \\ b_{01} & b_{11} & b_{21} & b_{31} \end{bmatrix}^{T} \end{split}\]

Once the temporary storage buffer has been allocated, analysis meta data is required. It can be obtained by hipsparseSbsrsm2_analysis “hipsparseXbsrsm2_analysis()”. The triangular solve is completed by calling hipsparseXbsrsm2_solve and once all solves are performed, the temporary storage buffer allocated by the user can be freed.

Solving a triangular system involves inverting the diagonal blocks. This means that if the sparse matrix is missing the diagonal block (referred to as a structural zero) or the diagonal block is not invertible (referred to as a numerical zero) then a solution is not possible. hipsparseXbsrsm2_solve tracks the location of the first zero pivot (either numerical or structural zero). The zero pivot status can be checked calling hipsparseXbsrsm2_zeroPivot(). If hipsparseXbsrsm2_zeroPivot() returns HIPSPARSE_STATUS_SUCCESS, then no zero pivot was found and therefore the matrix does not have a structural or numerical zero.

The user can specify that the sparse matrix should be interpreted as having identity blocks on the diagonal by setting the diagonal type on the descriptor descrA to HIPSPARSE_DIAG_TYPE_UNIT using hipsparseSetMatDiagType. If hipsparseDiagType_t == HIPSPARSE_DIAG_TYPE_UNIT, no zero pivot will be reported, even if the diagonal block \(A_{j,j}\) for some \(j\) is not invertible.

The sparse CSR matrix passed to hipsparseXbsrsm2_solve does not actually have to be a triangular matrix. Instead the triangular upper or lower part of the sparse matrix is solved based on hipsparseFillMode_t set on the descriptor descrA. If the fill mode is set to HIPSPARSE_FILL_MODE_LOWER, then the lower triangular matrix is solved. If the fill mode is set to HIPSPARSE_FILL_MODE_UPPER then the upper triangular matrix is solved.

Note

The sparse BSR matrix has to be sorted.

Note

Operation type of B and X must match, if \(op(B)=B, op(X)=X\).

Note

This function is non blocking and executed asynchronously with respect to the host. It may return before the actual computation has finished.

Note

Currently, only transA != HIPSPARSE_OPERATION_CONJUGATE_TRANSPOSE and transX != HIPSPARSE_OPERATION_CONJUGATE_TRANSPOSE is supported.

Parameters:

handle – [in] handle to the hipsparse library context queue.
dirA – [in] matrix storage of BSR blocks.
transA – [in] matrix \(A\) operation type.
transX – [in] matrix \(X\) operation type.
mb – [in] number of block rows of the sparse BSR matrix \(A\).
nrhs – [in] number of columns of the dense matrix \(op(X)\).
nnzb – [in] number of non-zero blocks of the sparse BSR matrix \(A\).
alpha – [in] scalar \(\alpha\).
descrA – [in] descriptor of the sparse BSR matrix \(A\).
bsrSortedValA – [in] array of nnzb blocks of the sparse BSR matrix.
bsrSortedRowPtrA – [in] array of mb+1 elements that point to the start of every block row of the sparse BSR matrix.
bsrSortedColIndA – [in] array of nnzb containing the block column indices of the sparse BSR matrix.
blockDim – [in] block dimension of the sparse BSR matrix.
info – [in] structure that holds the information collected during the analysis step.
B – [in] rhs matrix B with leading dimension ldb.
ldb – [in] leading dimension of rhs matrix \(B\).
X – [out] solution matrix X with leading dimension ldx.
ldx – [in] leading dimension of solution matrix \(X\).
policy – [in] HIPSPARSE_SOLVE_POLICY_NO_LEVEL or HIPSPARSE_SOLVE_POLICY_USE_LEVEL.
pBuffer – [in] temporary storage buffer allocated by the user.

Return values:

HIPSPARSE_STATUS_SUCCESS – the operation completed successfully.
HIPSPARSE_STATUS_INVALID_VALUE – handle, mb, nrhs, nnzb, blockDim, alpha, descrA, bsrSortedValA, bsrSortedRowPtrA, bsrSortedColIndA, B, X info or pBuffer is invalid.
HIPSPARSE_STATUS_INTERNAL_ERROR – an internal error occurred.
HIPSPARSE_STATUS_NOT_SUPPORTED – transA == HIPSPARSE_OPERATION_CONJUGATE_TRANSPOSE, transX == HIPSPARSE_OPERATION_CONJUGATE_TRANSPOSE or hipsparseMatrixType_t != HIPSPARSE_MATRIX_TYPE_GENERAL.

int main(int argc, char* argv[])
{
    // hipSPARSE handle
    hipsparseHandle_t handle;
    HIPSPARSE_CHECK(hipsparseCreate(&handle));

    // A = ( 1.0  0.0  0.0  0.0 )
    //     ( 2.0  3.0  0.0  0.0 )
    //     ( 4.0  5.0  6.0  0.0 )
    //     ( 7.0  0.0  8.0  9.0 )
    //
    // with bsr_dim = 2
    //
    //      -------------------
    //   = | 1.0 0.0 | 0.0 0.0 |
    //     | 2.0 3.0 | 0.0 0.0 |
    //      -------------------
    //     | 4.0 5.0 | 6.0 0.0 |
    //     | 7.0 0.0 | 8.0 9.0 |
    //      -------------------

    // Number of rows and columns
    const int m = 4;

    // Number of block rows and block columns
    const int mb = 2;
    const int nb = 2;

    // BSR block dimension
    const int bsr_dim = 2;

    // Number of right-hand-sides
    const int nrhs = 4;

    // Number of non-zero blocks
    const int nnzb = 3;

    // BSR row pointers
    std::vector<int> hbsrRowPtr = {0, 1, 3};

    // BSR column indices
    std::vector<int> hbsrColInd = {0, 0, 1};

    // BSR values
    std::vector<double> hbsrVal = {1.0, 2.0, 0.0, 3.0, 4.0, 7.0, 5.0, 0.0, 6.0, 8.0, 0.0, 9.0};

    // Storage scheme of the BSR blocks
    hipsparseDirection_t dir = HIPSPARSE_DIRECTION_COLUMN;

    // Transposition of the matrix and rhs matrix
    hipsparseOperation_t transA = HIPSPARSE_OPERATION_NON_TRANSPOSE;
    hipsparseOperation_t transX = HIPSPARSE_OPERATION_NON_TRANSPOSE;

    // Solve policy
    hipsparseSolvePolicy_t solve_policy = HIPSPARSE_SOLVE_POLICY_NO_LEVEL;

    // Scalar alpha and beta
    double alpha = 1.0;

    // rhs and solution matrix
    const int ldb = nb * bsr_dim;
    const int ldx = mb * bsr_dim;

    std::vector<double> hB = {1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 16};
    std::vector<double> hX(ldx * nrhs);

    // Offload data to device
    int*    dbsrRowPtr;
    int*    dbsrColInd;
    double* dbsrVal;
    double* dB;
    double* dX;

    HIP_CHECK(hipMalloc((void**)&dbsrRowPtr, sizeof(int) * (mb + 1)));
    HIP_CHECK(hipMalloc((void**)&dbsrColInd, sizeof(int) * nnzb));
    HIP_CHECK(hipMalloc((void**)&dbsrVal, sizeof(double) * nnzb * bsr_dim * bsr_dim));
    HIP_CHECK(hipMalloc((void**)&dB, sizeof(double) * nb * bsr_dim * nrhs));
    HIP_CHECK(hipMalloc((void**)&dX, sizeof(double) * mb * bsr_dim * nrhs));

    HIP_CHECK(
        hipMemcpy(dbsrRowPtr, hbsrRowPtr.data(), sizeof(int) * (mb + 1), hipMemcpyHostToDevice));
    HIP_CHECK(hipMemcpy(dbsrColInd, hbsrColInd.data(), sizeof(int) * nnzb, hipMemcpyHostToDevice));
    HIP_CHECK(hipMemcpy(
        dbsrVal, hbsrVal.data(), sizeof(double) * nnzb * bsr_dim * bsr_dim, hipMemcpyHostToDevice));
    HIP_CHECK(
        hipMemcpy(dB, hB.data(), sizeof(double) * nb * bsr_dim * nrhs, hipMemcpyHostToDevice));

    // Matrix descriptor
    hipsparseMatDescr_t descr;
    HIPSPARSE_CHECK(hipsparseCreateMatDescr(&descr));

    // Matrix fill mode
    HIPSPARSE_CHECK(hipsparseSetMatFillMode(descr, HIPSPARSE_FILL_MODE_LOWER));

    // Matrix diagonal type
    HIPSPARSE_CHECK(hipsparseSetMatDiagType(descr, HIPSPARSE_DIAG_TYPE_NON_UNIT));

    // Matrix info structure
    bsrsm2Info_t info;
    HIPSPARSE_CHECK(hipsparseCreateBsrsm2Info(&info));

    // Obtain required buffer size
    int buffer_size;
    HIPSPARSE_CHECK(hipsparseDbsrsm2_bufferSize(handle,
                                                dir,
                                                transA,
                                                transX,
                                                mb,
                                                nrhs,
                                                nnzb,
                                                descr,
                                                dbsrVal,
                                                dbsrRowPtr,
                                                dbsrColInd,
                                                bsr_dim,
                                                info,
                                                &buffer_size));

    // Allocate temporary buffer
    void* dbuffer;
    HIP_CHECK(hipMalloc(&dbuffer, buffer_size));

    // Perform analysis step
    HIPSPARSE_CHECK(hipsparseDbsrsm2_analysis(handle,
                                              dir,
                                              transA,
                                              transX,
                                              mb,
                                              nrhs,
                                              nnzb,
                                              descr,
                                              dbsrVal,
                                              dbsrRowPtr,
                                              dbsrColInd,
                                              bsr_dim,
                                              info,
                                              solve_policy,
                                              dbuffer));

    // Call dbsrsm to perform lower triangular solve LX = B
    HIPSPARSE_CHECK(hipsparseDbsrsm2_solve(handle,
                                           dir,
                                           transA,
                                           transX,
                                           mb,
                                           nrhs,
                                           nnzb,
                                           &alpha,
                                           descr,
                                           dbsrVal,
                                           dbsrRowPtr,
                                           dbsrColInd,
                                           bsr_dim,
                                           info,
                                           dB,
                                           ldb,
                                           dX,
                                           ldx,
                                           solve_policy,
                                           dbuffer));

    // Check for zero pivots
    int               pivot;
    hipsparseStatus_t status = hipsparseXbsrsm2_zeroPivot(handle, info, &pivot);

    if(status == HIPSPARSE_STATUS_ZERO_PIVOT)
    {
        std::cout << "Found zero pivot in matrix row " << pivot << std::endl;
    }

    // Copy result back to host
    HIP_CHECK(
        hipMemcpy(hX.data(), dX, sizeof(double) * mb * bsr_dim * nrhs, hipMemcpyDeviceToHost));

    std::cout << "hX" << std::endl;
    for(int i = 0; i < ldx * nrhs; i++)
    {
        std::cout << hX[i] << " ";
    }
    std::cout << std::endl;

    // Clear hipSPARSE
    HIPSPARSE_CHECK(hipsparseDestroyBsrsm2Info(info));
    HIPSPARSE_CHECK(hipsparseDestroyMatDescr(descr));
    HIPSPARSE_CHECK(hipsparseDestroy(handle));

    // Clear device memory
    HIP_CHECK(hipFree(dbsrRowPtr));
    HIP_CHECK(hipFree(dbsrColInd));
    HIP_CHECK(hipFree(dbsrVal));
    HIP_CHECK(hipFree(dB));
    HIP_CHECK(hipFree(dX));
    HIP_CHECK(hipFree(dbuffer));

    return 0;
}

int main(int argc, char* argv[])
{
    // hipSPARSE handle
    hipsparseHandle_t handle;
    HIPSPARSE_CHECK(hipsparseCreate(&handle));

    // A = ( 1.0  0.0  0.0  0.0 )
    //     ( 2.0  3.0  0.0  0.0 )
    //     ( 4.0  5.0  6.0  0.0 )
    //     ( 7.0  0.0  8.0  9.0 )
    //
    // with bsr_dim = 2
    //
    //      -------------------
    //   = | 1.0 0.0 | 0.0 0.0 |
    //     | 2.0 3.0 | 0.0 0.0 |
    //      -------------------
    //     | 4.0 5.0 | 6.0 0.0 |
    //     | 7.0 0.0 | 8.0 9.0 |
    //      -------------------

    // Number of rows and columns
    const int m = 4;

    // Number of block rows and block columns
    const int mb = 2;
    const int nb = 2;

    // BSR block dimension
    const int bsr_dim = 2;

    // Number of right-hand-sides
    const int nrhs = 4;

    // Number of non-zero blocks
    const int nnzb = 3;

    // BSR row pointers
    int hbsrRowPtr[] = {0, 1, 3};

    // BSR column indices
    int hbsrColInd[] = {0, 0, 1};

    // BSR values
    double hbsrVal[] = {1.0, 2.0, 0.0, 3.0, 4.0, 7.0, 5.0, 0.0, 6.0, 8.0, 0.0, 9.0};

    // Storage scheme of the BSR blocks
    hipsparseDirection_t dir = HIPSPARSE_DIRECTION_COLUMN;

    // Transposition of the matrix and rhs matrix
    hipsparseOperation_t transA = HIPSPARSE_OPERATION_NON_TRANSPOSE;
    hipsparseOperation_t transX = HIPSPARSE_OPERATION_NON_TRANSPOSE;

    // Solve policy
    hipsparseSolvePolicy_t solve_policy = HIPSPARSE_SOLVE_POLICY_NO_LEVEL;

    // Scalar alpha and beta
    double alpha = 1.0;

    // rhs and solution matrix
    const int ldb = nb * bsr_dim;
    const int ldx = mb * bsr_dim;

    double hB[] = {1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 16};
    double hX[ldx * nrhs];

    // Offload data to device
    int*    dbsrRowPtr;
    int*    dbsrColInd;
    double* dbsrVal;
    double* dB;
    double* dX;

    HIP_CHECK(hipMalloc((void**)&dbsrRowPtr, sizeof(int) * (mb + 1)));
    HIP_CHECK(hipMalloc((void**)&dbsrColInd, sizeof(int) * nnzb));
    HIP_CHECK(hipMalloc((void**)&dbsrVal, sizeof(double) * nnzb * bsr_dim * bsr_dim));
    HIP_CHECK(hipMalloc((void**)&dB, sizeof(double) * nb * bsr_dim * nrhs));
    HIP_CHECK(hipMalloc((void**)&dX, sizeof(double) * mb * bsr_dim * nrhs));

    HIP_CHECK(hipMemcpy(dbsrRowPtr, hbsrRowPtr, sizeof(int) * (mb + 1), hipMemcpyHostToDevice));
    HIP_CHECK(hipMemcpy(dbsrColInd, hbsrColInd, sizeof(int) * nnzb, hipMemcpyHostToDevice));
    HIP_CHECK(hipMemcpy(
        dbsrVal, hbsrVal, sizeof(double) * nnzb * bsr_dim * bsr_dim, hipMemcpyHostToDevice));
    HIP_CHECK(hipMemcpy(dB, hB, sizeof(double) * nb * bsr_dim * nrhs, hipMemcpyHostToDevice));

    // Matrix descriptor
    hipsparseMatDescr_t descr;
    HIPSPARSE_CHECK(hipsparseCreateMatDescr(&descr));

    // Matrix fill mode
    HIPSPARSE_CHECK(hipsparseSetMatFillMode(descr, HIPSPARSE_FILL_MODE_LOWER));

    // Matrix diagonal type
    HIPSPARSE_CHECK(hipsparseSetMatDiagType(descr, HIPSPARSE_DIAG_TYPE_NON_UNIT));

    // Matrix info structure
    bsrsm2Info_t info;
    HIPSPARSE_CHECK(hipsparseCreateBsrsm2Info(&info));

    // Obtain required buffer size
    int buffer_size;
    HIPSPARSE_CHECK(hipsparseDbsrsm2_bufferSize(handle,
                                                dir,
                                                transA,
                                                transX,
                                                mb,
                                                nrhs,
                                                nnzb,
                                                descr,
                                                dbsrVal,
                                                dbsrRowPtr,
                                                dbsrColInd,
                                                bsr_dim,
                                                info,
                                                &buffer_size));

    // Allocate temporary buffer
    void* dbuffer;
    HIP_CHECK(hipMalloc(&dbuffer, buffer_size));

    // Perform analysis step
    HIPSPARSE_CHECK(hipsparseDbsrsm2_analysis(handle,
                                              dir,
                                              transA,
                                              transX,
                                              mb,
                                              nrhs,
                                              nnzb,
                                              descr,
                                              dbsrVal,
                                              dbsrRowPtr,
                                              dbsrColInd,
                                              bsr_dim,
                                              info,
                                              solve_policy,
                                              dbuffer));

    // Call dbsrsm to perform lower triangular solve LX = B
    HIPSPARSE_CHECK(hipsparseDbsrsm2_solve(handle,
                                           dir,
                                           transA,
                                           transX,
                                           mb,
                                           nrhs,
                                           nnzb,
                                           &alpha,
                                           descr,
                                           dbsrVal,
                                           dbsrRowPtr,
                                           dbsrColInd,
                                           bsr_dim,
                                           info,
                                           dB,
                                           ldb,
                                           dX,
                                           ldx,
                                           solve_policy,
                                           dbuffer));

    // Check for zero pivots
    int               pivot;
    hipsparseStatus_t status = hipsparseXbsrsm2_zeroPivot(handle, info, &pivot);

    if(status == HIPSPARSE_STATUS_ZERO_PIVOT)
    {
        printf("Found zero pivot in matrix row %d\n", pivot);
    }

    // Copy result back to host
    HIP_CHECK(hipMemcpy(hX, dX, sizeof(double) * mb * bsr_dim * nrhs, hipMemcpyDeviceToHost));

    printf("hX\n");
    for(int i = 0; i < ldx * nrhs; i++)
    {
        printf("%f ", hX[i]);
    }
    printf("\n");

    // Clear hipSPARSE
    HIPSPARSE_CHECK(hipsparseDestroyBsrsm2Info(info));
    HIPSPARSE_CHECK(hipsparseDestroyMatDescr(descr));
    HIPSPARSE_CHECK(hipsparseDestroy(handle));

    // Clear device memory
    HIP_CHECK(hipFree(dbsrRowPtr));
    HIP_CHECK(hipFree(dbsrColInd));
    HIP_CHECK(hipFree(dbsrVal));
    HIP_CHECK(hipFree(dB));
    HIP_CHECK(hipFree(dX));
    HIP_CHECK(hipFree(dbuffer));

    return 0;
}

hipsparseXcsrsm2_zeroPivot()#

hipsparseStatus_t hipsparseXcsrsm2_zeroPivot(hipsparseHandle_t handle, csrsm2Info_t info, int *position)#

hipsparseXcsrsm2_zeroPivot returns HIPSPARSE_STATUS_ZERO_PIVOT, if either a structural or numerical zero has been found during hipsparseXcsrsm2_analysis() or hipsparseXcsrsm2_solve() computation. The first zero pivot \(j\) at \(A_{j,j}\) is stored in position, using same index base as the CSR matrix.

position can be in host or device memory. If no zero pivot has been found, position is set to -1 and HIPSPARSE_STATUS_SUCCESS is returned instead.

Deprecated:: This function is deprecated when using the CUDA backend (CUDA 11.0+) and will be removed in CUDA 12.0. This deprecation does not apply to the ROCm backend.

Note

hipsparseXcsrsm2_zeroPivot is a blocking function. It might influence performance negatively.

Parameters:

handle – [in] handle to the hipsparse library context queue.
info – [in] structure that holds the information collected during the analysis step.
position – [inout] pointer to zero pivot \(j\), can be in host or device memory.

Return values:

HIPSPARSE_STATUS_SUCCESS – the operation completed successfully.
HIPSPARSE_STATUS_NOT_INITIALIZED – handle is not initialized.
HIPSPARSE_STATUS_INVALID_VALUE – handle, info or position is nullptr.
HIPSPARSE_STATUS_INTERNAL_ERROR – an internal error occurred.
HIPSPARSE_STATUS_ZERO_PIVOT – zero pivot has been found.

hipsparseXcsrsm2_bufferSizeExt()#

hipsparseStatus_t hipsparseScsrsm2_bufferSizeExt(hipsparseHandle_t handle, int algo, hipsparseOperation_t transA, hipsparseOperation_t transB, int m, int nrhs, int nnz, const float *alpha, const hipsparseMatDescr_t descrA, const float *csrSortedValA, const int *csrSortedRowPtrA, const int *csrSortedColIndA, const float *B, int ldb, csrsm2Info_t info, hipsparseSolvePolicy_t policy, size_t *pBufferSizeInBytes)#

hipsparseStatus_t hipsparseDcsrsm2_bufferSizeExt(hipsparseHandle_t handle, int algo, hipsparseOperation_t transA, hipsparseOperation_t transB, int m, int nrhs, int nnz, const double *alpha, const hipsparseMatDescr_t descrA, const double *csrSortedValA, const int *csrSortedRowPtrA, const int *csrSortedColIndA, const double *B, int ldb, csrsm2Info_t info, hipsparseSolvePolicy_t policy, size_t *pBufferSizeInBytes)#

hipsparseStatus_t hipsparseCcsrsm2_bufferSizeExt(hipsparseHandle_t handle, int algo, hipsparseOperation_t transA, hipsparseOperation_t transB, int m, int nrhs, int nnz, const hipComplex *alpha, const hipsparseMatDescr_t descrA, const hipComplex *csrSortedValA, const int *csrSortedRowPtrA, const int *csrSortedColIndA, const hipComplex *B, int ldb, csrsm2Info_t info, hipsparseSolvePolicy_t policy, size_t *pBufferSizeInBytes)#

hipsparseStatus_t hipsparseZcsrsm2_bufferSizeExt(hipsparseHandle_t handle, int algo, hipsparseOperation_t transA, hipsparseOperation_t transB, int m, int nrhs, int nnz, const hipDoubleComplex *alpha, const hipsparseMatDescr_t descrA, const hipDoubleComplex *csrSortedValA, const int *csrSortedRowPtrA, const int *csrSortedColIndA, const hipDoubleComplex *B, int ldb, csrsm2Info_t info, hipsparseSolvePolicy_t policy, size_t *pBufferSizeInBytes)#

hipsparseXcsrsm2_bufferSizeExt returns the size of the temporary storage buffer in bytes that is required by hipsparseXcsrsm2_analysis() and hipsparseXcsrsm2_solve(). The temporary storage buffer must be allocated by the user.

Parameters:

handle – [in] handle to the hipsparse library context queue.
algo – [in] algorithm to use.
transA – [in] matrix \(A\) operation type.
transB – [in] matrix \(B\) operation type.
m – [in] number of rows of the sparse CSR matrix \(A\).
nrhs – [in] number of columns of the dense matrix \(op(B)\).
nnz – [in] number of non-zero entries of the sparse CSR matrix \(A\).
alpha – [in] scalar \(\alpha\).
descrA – [in] descriptor of the sparse CSR matrix \(A\).
csrSortedValA – [in] array of nnz elements of the sparse CSR matrix \(A\).
csrSortedRowPtrA – [in] array of m+1 elements that point to the start of every row of the sparse CSR matrix \(A\).
csrSortedColIndA – [in] array of nnz elements containing the column indices of the sparse CSR matrix \(A\).
B – [in] array of m \(\times\) nrhs elements of the rhs matrix \(B\).
ldb – [in] leading dimension of rhs matrix \(B\).
info – [in] structure that holds the information collected during the analysis step.
policy – [in] HIPSPARSE_SOLVE_POLICY_NO_LEVEL or HIPSPARSE_SOLVE_POLICY_USE_LEVEL.
pBufferSizeInBytes – [out] number of bytes of the temporary storage buffer required by hipsparseXcsrsm2_analysis() and hipsparseXcsrsm2_solve().

Return values:

HIPSPARSE_STATUS_SUCCESS – the operation completed successfully.
HIPSPARSE_STATUS_INVALID_VALUE – handle, m, nrhs, nnz, alpha, descrA, csrSortedValA, csrSortedRowPtrA, csrSortedColIndA, B, info or pBufferSizeInBytes is invalid.
HIPSPARSE_STATUS_INTERNAL_ERROR – an internal error occurred.
HIPSPARSE_STATUS_NOT_SUPPORTED – transA == HIPSPARSE_OPERATION_CONJUGATE_TRANSPOSE, transB == HIPSPARSE_OPERATION_CONJUGATE_TRANSPOSE or hipsparseMatrixType_t != HIPSPARSE_MATRIX_TYPE_GENERAL.

hipsparseXcsrsm2_analysis()#

hipsparseStatus_t hipsparseScsrsm2_analysis(hipsparseHandle_t handle, int algo, hipsparseOperation_t transA, hipsparseOperation_t transB, int m, int nrhs, int nnz, const float *alpha, const hipsparseMatDescr_t descrA, const float *csrSortedValA, const int *csrSortedRowPtrA, const int *csrSortedColIndA, const float *B, int ldb, csrsm2Info_t info, hipsparseSolvePolicy_t policy, void *pBuffer)#

hipsparseStatus_t hipsparseDcsrsm2_analysis(hipsparseHandle_t handle, int algo, hipsparseOperation_t transA, hipsparseOperation_t transB, int m, int nrhs, int nnz, const double *alpha, const hipsparseMatDescr_t descrA, const double *csrSortedValA, const int *csrSortedRowPtrA, const int *csrSortedColIndA, const double *B, int ldb, csrsm2Info_t info, hipsparseSolvePolicy_t policy, void *pBuffer)#

hipsparseStatus_t hipsparseCcsrsm2_analysis(hipsparseHandle_t handle, int algo, hipsparseOperation_t transA, hipsparseOperation_t transB, int m, int nrhs, int nnz, const hipComplex *alpha, const hipsparseMatDescr_t descrA, const hipComplex *csrSortedValA, const int *csrSortedRowPtrA, const int *csrSortedColIndA, const hipComplex *B, int ldb, csrsm2Info_t info, hipsparseSolvePolicy_t policy, void *pBuffer)#

hipsparseStatus_t hipsparseZcsrsm2_analysis(hipsparseHandle_t handle, int algo, hipsparseOperation_t transA, hipsparseOperation_t transB, int m, int nrhs, int nnz, const hipDoubleComplex *alpha, const hipsparseMatDescr_t descrA, const hipDoubleComplex *csrSortedValA, const int *csrSortedRowPtrA, const int *csrSortedColIndA, const hipDoubleComplex *B, int ldb, csrsm2Info_t info, hipsparseSolvePolicy_t policy, void *pBuffer)#

hipsparseXcsrsm2_analysis performs the analysis step for hipsparseXcsrsm2_solve(). It is expected that this function will be executed only once for a given matrix and particular operation type.

Note

If the matrix sparsity pattern changes, the gathered information will become invalid.

Note

This function is non blocking and executed asynchronously with respect to the host. It may return before the actual computation has finished.

Parameters:

handle – [in] handle to the hipsparse library context queue.
algo – [in] algorithm to use.
transA – [in] matrix \(A\) operation type.
transB – [in] matrix \(B\) operation type.
m – [in] number of rows of the sparse CSR matrix \(A\).
nrhs – [in] number of columns of the dense matrix \(op(B)\).
nnz – [in] number of non-zero entries of the sparse CSR matrix \(A\).
alpha – [in] scalar \(\alpha\).
descrA – [in] descriptor of the sparse CSR matrix \(A\).
csrSortedValA – [in] array of nnz elements of the sparse CSR matrix \(A\).
csrSortedRowPtrA – [in] array of m+1 elements that point to the start of every row of the sparse CSR matrix \(A\).
csrSortedColIndA – [in] array of nnz elements containing the column indices of the sparse CSR matrix \(A\).
B – [in] array of m \(\times\) nrhs elements of the rhs matrix \(B\).
ldb – [in] leading dimension of rhs matrix \(B\).
info – [out] structure that holds the information collected during the analysis step.
policy – [in] HIPSPARSE_SOLVE_POLICY_NO_LEVEL or HIPSPARSE_SOLVE_POLICY_USE_LEVEL.
pBuffer – [in] temporary storage buffer allocated by the user.

Return values:

HIPSPARSE_STATUS_SUCCESS – the operation completed successfully.
HIPSPARSE_STATUS_INVALID_VALUE – handle, m, nrhs, nnz, alpha, descrA, csrSortedValA, csrSortedRowPtrA, csrSortedColIndA, B, info or pBuffer is invalid.
HIPSPARSE_STATUS_INTERNAL_ERROR – an internal error occurred.
HIPSPARSE_STATUS_NOT_SUPPORTED – transA == HIPSPARSE_OPERATION_CONJUGATE_TRANSPOSE, transB == HIPSPARSE_OPERATION_CONJUGATE_TRANSPOSE or hipsparseMatrixType_t != HIPSPARSE_MATRIX_TYPE_GENERAL.

hipsparseXcsrsm2_solve()#

hipsparseStatus_t hipsparseScsrsm2_solve(hipsparseHandle_t handle, int algo, hipsparseOperation_t transA, hipsparseOperation_t transB, int m, int nrhs, int nnz, const float *alpha, const hipsparseMatDescr_t descrA, const float *csrSortedValA, const int *csrSortedRowPtrA, const int *csrSortedColIndA, float *B, int ldb, csrsm2Info_t info, hipsparseSolvePolicy_t policy, void *pBuffer)#

hipsparseStatus_t hipsparseDcsrsm2_solve(hipsparseHandle_t handle, int algo, hipsparseOperation_t transA, hipsparseOperation_t transB, int m, int nrhs, int nnz, const double *alpha, const hipsparseMatDescr_t descrA, const double *csrSortedValA, const int *csrSortedRowPtrA, const int *csrSortedColIndA, double *B, int ldb, csrsm2Info_t info, hipsparseSolvePolicy_t policy, void *pBuffer)#

hipsparseStatus_t hipsparseCcsrsm2_solve(hipsparseHandle_t handle, int algo, hipsparseOperation_t transA, hipsparseOperation_t transB, int m, int nrhs, int nnz, const hipComplex *alpha, const hipsparseMatDescr_t descrA, const hipComplex *csrSortedValA, const int *csrSortedRowPtrA, const int *csrSortedColIndA, hipComplex *B, int ldb, csrsm2Info_t info, hipsparseSolvePolicy_t policy, void *pBuffer)#

hipsparseStatus_t hipsparseZcsrsm2_solve(hipsparseHandle_t handle, int algo, hipsparseOperation_t transA, hipsparseOperation_t transB, int m, int nrhs, int nnz, const hipDoubleComplex *alpha, const hipsparseMatDescr_t descrA, const hipDoubleComplex *csrSortedValA, const int *csrSortedRowPtrA, const int *csrSortedColIndA, hipDoubleComplex *B, int ldb, csrsm2Info_t info, hipsparseSolvePolicy_t policy, void *pBuffer)#

Sparse triangular system solve using CSR storage format.

hipsparseXcsrsm2_solve solves a sparse triangular linear system of a sparse \(m \times m\) matrix, defined in CSR storage format, a column-oriented dense solution matrix \(X\) and the column-oriented dense right-hand side matrix \(B\) that is multiplied by \(\alpha\), such that

\[ op(A) \cdot op(X) = \alpha \cdot op(B), \]

with

\[\begin{split} op(A) = \left\{ \begin{array}{ll} A, & \text{if transA == HIPSPARSE_OPERATION_NON_TRANSPOSE} \\ A^T, & \text{if transA == HIPSPARSE_OPERATION_TRANSPOSE} \\ A^H, & \text{if transA == HIPSPARSE_OPERATION_CONJUGATE_TRANSPOSE} \end{array} \right. \end{split}\]

,

\[\begin{split} op(B) = \left\{ \begin{array}{ll} B, & \text{if transB == HIPSPARSE_OPERATION_NON_TRANSPOSE} \\ B^T, & \text{if transB == HIPSPARSE_OPERATION_TRANSPOSE} \\ B^H, & \text{if transB == HIPSPARSE_OPERATION_CONJUGATE_TRANSPOSE} \end{array} \right. \end{split}\]

and

\[\begin{split} op(X) = \left\{ \begin{array}{ll} X, & \text{if transB == HIPSPARSE_OPERATION_NON_TRANSPOSE} \\ X^T, & \text{if transB == HIPSPARSE_OPERATION_TRANSPOSE} \\ X^H, & \text{if transB == HIPSPARSE_OPERATION_CONJUGATE_TRANSPOSE} \end{array} \right. \end{split}\]

The solution is performed inplace meaning that the matrix \(B\) is overwritten with the solution \(X\) after calling hipsparseXcsrsm2_solve. Given that the sparse matrix \(A\) is a square matrix, its size is \(m \times m\) regardless of whether \(A\) is transposed or not. The size of the column-oriented dense matrices \(B\) and \(X\) have size that depends on the value of transB:

\[\begin{split} op(B)/op(X) = \left\{ \begin{array}{ll} ldb \times nrhs, \text{ } ldb \ge m, & \text{if transB == HIPSPARSE_OPERATION_NON_TRANSPOSE} \\ ldb \times m, \text{ } ldb \ge nrhs, & \text{if transB == HIPSPARSE_OPERATION_TRANSPOSE} \\ ldb \times m, \text{ } ldb \ge nrhs, & \text{if transB == HIPSPARSE_OPERATION_CONJUGATE_TRANSPOSE} \end{array} \right. \end{split}\]

hipsparseXcsrsm2_solve requires a user allocated temporary buffer. Its size is returned by hipsparseXcsrsm2_bufferSizeExt(). The size of the required buffer is larger when transA equals HIPSPARSE_OPERATION_TRANSPOSE or HIPSPARSE_OPERATION_CONJUGATE_TRANSPOSE and when transB is HIPSPARSE_OPERATION_NON_TRANSPOSE. The subsequent solve will also be faster when \(A\) is non-transposed and \(B\) is transposed (or conjugate transposed). For example, instead of solving:

\[\begin{split} \begin{bmatrix} a_{00} & 0 & 0 \\ a_{10} & a_{11} & 0 \\ a_{20} & a_{21} & a_{22} \\ \end{bmatrix} \cdot \begin{bmatrix} x_{00} & x_{01} \\ x_{10} & x_{11} \\ x_{20} & x_{21} \\ \end{bmatrix} = \begin{bmatrix} b_{00} & b_{01} \\ b_{10} & b_{11} \\ b_{20} & b_{21} \\ \end{bmatrix} \end{split}\]

Consider solving:

\[\begin{split} \begin{bmatrix} a_{00} & 0 & 0 \\ a_{10} & a_{11} & 0 \\ a_{20} & a_{21} & a_{22} \end{bmatrix} \cdot \begin{bmatrix} x_{00} & x_{10} & x_{20} \\ x_{01} & x_{11} & x_{21} \end{bmatrix}^{T} = \begin{bmatrix} b_{00} & b_{10} & b_{20} \\ b_{01} & b_{11} & b_{21} \end{bmatrix}^{T} \end{split}\]

Once the temporary storage buffer has been allocated, analysis meta data is required. It can be obtained by hipsparseXcsrsm2_analysis(). The triangular solve is completed by calling hipsparseXcsrsm2_solve and once all solves are performed, the temporary storage buffer allocated by the user can be freed.

Solving a triangular system involves division by the diagonal elements. This means that if the sparse matrix is missing the diagonal entry (referred to as a structural zero) or the diagonal entry is zero (referred to as a numerical zero) then a division by zero would occur. hipsparseXcsrsm2_solve tracks the location of the first zero pivot (either numerical or structural zero). The zero pivot status can be checked calling hipsparseXcsrsm2_zeroPivot(). If hipsparseXcsrsm2_zeroPivot() returns HIPSPARSE_STATUS_SUCCESS, then no zero pivot was found and therefore the matrix does not have a structural or numerical zero.

The user can specify that the sparse matrix should be interpreted as having ones on the diagonal by setting the diagonal type on the descriptor descrA to HIPSPARSE_DIAG_TYPE_UNIT using hipsparseSetMatDiagType. If hipsparseDiagType_t == HIPSPARSE_DIAG_TYPE_UNIT, no zero pivot will be reported, even if \(A_{j,j} = 0\) for some \(j\).

The sparse CSR matrix passed to hipsparseXcsrsm2_solve does not actually have to be a triangular matrix. Instead the triangular upper or lower part of the sparse matrix is solved based on hipsparseFillMode_t set on the descriptor descrA. If the fill mode is set to HIPSPARSE_FILL_MODE_LOWER, then the lower triangular matrix is solved. If the fill mode is set to HIPSPARSE_FILL_MODE_UPPER then the upper triangular matrix is solved.

Note

The sparse CSR matrix has to be sorted. This can be achieved by calling hipsparseXcsrsort().

Note

This function is non blocking and executed asynchronously with respect to the host. It may return before the actual computation has finished.

Note

Currently, only transA != HIPSPARSE_OPERATION_CONJUGATE_TRANSPOSE and transB != HIPSPARSE_OPERATION_CONJUGATE_TRANSPOSE is supported.

Parameters:

handle – [in] handle to the hipsparse library context queue.
algo – [in] algorithm to use.
transA – [in] matrix \(A\) operation type.
transB – [in] matrix \(B\) operation type.
m – [in] number of rows of the sparse CSR matrix \(A\).
nrhs – [in] number of columns of the dense matrix \(op(B)\).
nnz – [in] number of non-zero entries of the sparse CSR matrix \(A\).
alpha – [in] scalar \(\alpha\).
descrA – [in] descriptor of the sparse CSR matrix \(A\).
csrSortedValA – [in] array of nnz elements of the sparse CSR matrix \(A\).
csrSortedRowPtrA – [in] array of m+1 elements that point to the start of every row of the sparse CSR matrix \(A\).
csrSortedColIndA – [in] array of nnz elements containing the column indices of the sparse CSR matrix \(A\).
B – [inout] array of m \(\times\) nrhs elements of the rhs matrix \(B\).
ldb – [in] leading dimension of rhs matrix \(B\).
info – [in] structure that holds the information collected during the analysis step.
policy – [in] HIPSPARSE_SOLVE_POLICY_NO_LEVEL or HIPSPARSE_SOLVE_POLICY_USE_LEVEL.
pBuffer – [in] temporary storage buffer allocated by the user.

Return values:

HIPSPARSE_STATUS_SUCCESS – the operation completed successfully.
HIPSPARSE_STATUS_INVALID_VALUE – handle, m, nrhs, nnz, alpha, descrA, csrSortedValA, csrSortedRowPtrA, csrSortedColIndA, B, info or pBuffer is invalid.
HIPSPARSE_STATUS_INTERNAL_ERROR – an internal error occurred.
HIPSPARSE_STATUS_NOT_SUPPORTED – transA == HIPSPARSE_OPERATION_CONJUGATE_TRANSPOSE, transB == HIPSPARSE_OPERATION_CONJUGATE_TRANSPOSE or hipsparseMatrixType_t != HIPSPARSE_MATRIX_TYPE_GENERAL.

int main(int argc, char* argv[])
{
    // hipSPARSE handle
    hipsparseHandle_t handle;
    HIPSPARSE_CHECK(hipsparseCreate(&handle));

    // A = ( 1.0  0.0  0.0  0.0 )
    //     ( 2.0  3.0  0.0  0.0 )
    //     ( 4.0  5.0  6.0  0.0 )
    //     ( 7.0  0.0  8.0  9.0 )

    // Number of rows and columns
    int m = 4;
    int n = 4;

    // Number of right-hand-sides
    int nrhs = 4;

    // Number of non-zeros
    int nnz = 9;

    // CSR row pointers
    std::vector<int> hcsrRowPtr = {0, 1, 3, 6, 9};

    // CSR column indices
    std::vector<int> hcsrColInd = {0, 0, 1, 0, 1, 2, 0, 2, 3};

    // CSR values
    std::vector<double> hcsrVal = {1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0, 8.0, 9.0};

    // Transposition of the matrix and rhs matrix
    hipsparseOperation_t transA = HIPSPARSE_OPERATION_NON_TRANSPOSE;
    hipsparseOperation_t transB = HIPSPARSE_OPERATION_NON_TRANSPOSE;

    // Solve policy
    hipsparseSolvePolicy_t solve_policy = HIPSPARSE_SOLVE_POLICY_NO_LEVEL;

    // Scalar alpha and beta
    double alpha = 1.0;

    // rhs and solution matrix
    int ldb = n;

    std::vector<double> hB = {1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 16};

    // Offload data to device
    int*    dcsrRowPtr;
    int*    dcsrColInd;
    double* dcsrVal;
    double* dB;

    HIP_CHECK(hipMalloc((void**)&dcsrRowPtr, sizeof(int) * (m + 1)));
    HIP_CHECK(hipMalloc((void**)&dcsrColInd, sizeof(int) * nnz));
    HIP_CHECK(hipMalloc((void**)&dcsrVal, sizeof(double) * nnz));
    HIP_CHECK(hipMalloc((void**)&dB, sizeof(double) * n * nrhs));

    HIP_CHECK(
        hipMemcpy(dcsrRowPtr, hcsrRowPtr.data(), sizeof(int) * (m + 1), hipMemcpyHostToDevice));
    HIP_CHECK(hipMemcpy(dcsrColInd, hcsrColInd.data(), sizeof(int) * nnz, hipMemcpyHostToDevice));
    HIP_CHECK(hipMemcpy(dcsrVal, hcsrVal.data(), sizeof(double) * nnz, hipMemcpyHostToDevice));
    HIP_CHECK(hipMemcpy(dB, hB.data(), sizeof(double) * n * nrhs, hipMemcpyHostToDevice));

    // Matrix descriptor
    hipsparseMatDescr_t descr;
    HIPSPARSE_CHECK(hipsparseCreateMatDescr(&descr));

    // Matrix fill mode
    HIPSPARSE_CHECK(hipsparseSetMatFillMode(descr, HIPSPARSE_FILL_MODE_LOWER));

    // Matrix diagonal type
    HIPSPARSE_CHECK(hipsparseSetMatDiagType(descr, HIPSPARSE_DIAG_TYPE_NON_UNIT));

    // Matrix info structure
    csrsm2Info_t info;
    HIPSPARSE_CHECK(hipsparseCreateCsrsm2Info(&info));

    // Obtain required buffer size
    size_t buffer_size;
    HIPSPARSE_CHECK(hipsparseDcsrsm2_bufferSizeExt(handle,
                                                   0,
                                                   transA,
                                                   transB,
                                                   m,
                                                   nrhs,
                                                   nnz,
                                                   &alpha,
                                                   descr,
                                                   dcsrVal,
                                                   dcsrRowPtr,
                                                   dcsrColInd,
                                                   dB,
                                                   ldb,
                                                   info,
                                                   solve_policy,
                                                   &buffer_size));

    // Allocate temporary buffer
    void* dbuffer;
    HIP_CHECK(hipMalloc(&dbuffer, buffer_size));

    // Perform analysis step
    HIPSPARSE_CHECK(hipsparseDcsrsm2_analysis(handle,
                                              0,
                                              transA,
                                              transB,
                                              m,
                                              nrhs,
                                              nnz,
                                              &alpha,
                                              descr,
                                              dcsrVal,
                                              dcsrRowPtr,
                                              dcsrColInd,
                                              dB,
                                              ldb,
                                              info,
                                              solve_policy,
                                              dbuffer));

    // Call dcsrsm to perform lower triangular solve LB = B
    HIPSPARSE_CHECK(hipsparseDcsrsm2_solve(handle,
                                           0,
                                           transA,
                                           transB,
                                           m,
                                           nrhs,
                                           nnz,
                                           &alpha,
                                           descr,
                                           dcsrVal,
                                           dcsrRowPtr,
                                           dcsrColInd,
                                           dB,
                                           ldb,
                                           info,
                                           solve_policy,
                                           dbuffer));

    // Check for zero pivots
    int               pivot;
    hipsparseStatus_t status = hipsparseXcsrsm2_zeroPivot(handle, info, &pivot);

    if(status == HIPSPARSE_STATUS_ZERO_PIVOT)
    {
        std::cout << "Found zero pivot in matrix row " << pivot << std::endl;
    }

    // Copy result back to host
    HIP_CHECK(hipMemcpy(hB.data(), dB, sizeof(double) * m * nrhs, hipMemcpyDeviceToHost));

    std::cout << "hB" << std::endl;
    for(size_t i = 0; i < hB.size(); i++)
    {
        std::cout << hB[i] << " ";
    }
    std::cout << "" << std::endl;

    // Clear hipSPARSE
    HIPSPARSE_CHECK(hipsparseDestroyCsrsm2Info(info));
    HIPSPARSE_CHECK(hipsparseDestroyMatDescr(descr));
    HIPSPARSE_CHECK(hipsparseDestroy(handle));

    // Clear device memory
    HIP_CHECK(hipFree(dcsrRowPtr));
    HIP_CHECK(hipFree(dcsrColInd));
    HIP_CHECK(hipFree(dcsrVal));
    HIP_CHECK(hipFree(dB));
    HIP_CHECK(hipFree(dbuffer));

    return 0;
}

int main(int argc, char* argv[])
{
    // hipSPARSE handle
    hipsparseHandle_t handle;
    HIPSPARSE_CHECK(hipsparseCreate(&handle));

    // A = ( 1.0  0.0  0.0  0.0 )
    //     ( 2.0  3.0  0.0  0.0 )
    //     ( 4.0  5.0  6.0  0.0 )
    //     ( 7.0  0.0  8.0  9.0 )

    // Number of rows and columns
    int m = 4;
    int n = 4;

    // Number of right-hand-sides
    int nrhs = 4;

    // Number of non-zeros
    int nnz = 9;

    // CSR row pointers
    int hcsrRowPtr[] = {0, 1, 3, 6, 9};

    // CSR column indices
    int hcsrColInd[] = {0, 0, 1, 0, 1, 2, 0, 2, 3};

    // CSR values
    double hcsrVal[] = {1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0, 8.0, 9.0};

    // Transposition of the matrix and rhs matrix
    hipsparseOperation_t transA = HIPSPARSE_OPERATION_NON_TRANSPOSE;
    hipsparseOperation_t transB = HIPSPARSE_OPERATION_NON_TRANSPOSE;

    // Solve policy
    hipsparseSolvePolicy_t solve_policy = HIPSPARSE_SOLVE_POLICY_NO_LEVEL;

    // Scalar alpha and beta
    double alpha = 1.0;

    // rhs and solution matrix
    int ldb = n;

    double hB[] = {1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 16};

    // Offload data to device
    int*    dcsrRowPtr;
    int*    dcsrColInd;
    double* dcsrVal;
    double* dB;

    HIP_CHECK(hipMalloc((void**)&dcsrRowPtr, sizeof(int) * (m + 1)));
    HIP_CHECK(hipMalloc((void**)&dcsrColInd, sizeof(int) * nnz));
    HIP_CHECK(hipMalloc((void**)&dcsrVal, sizeof(double) * nnz));
    HIP_CHECK(hipMalloc((void**)&dB, sizeof(double) * n * nrhs));

    HIP_CHECK(hipMemcpy(dcsrRowPtr, hcsrRowPtr, sizeof(int) * (m + 1), hipMemcpyHostToDevice));
    HIP_CHECK(hipMemcpy(dcsrColInd, hcsrColInd, sizeof(int) * nnz, hipMemcpyHostToDevice));
    HIP_CHECK(hipMemcpy(dcsrVal, hcsrVal, sizeof(double) * nnz, hipMemcpyHostToDevice));
    HIP_CHECK(hipMemcpy(dB, hB, sizeof(double) * n * nrhs, hipMemcpyHostToDevice));

    // Matrix descriptor
    hipsparseMatDescr_t descr;
    HIPSPARSE_CHECK(hipsparseCreateMatDescr(&descr));

    // Matrix fill mode
    HIPSPARSE_CHECK(hipsparseSetMatFillMode(descr, HIPSPARSE_FILL_MODE_LOWER));

    // Matrix diagonal type
    HIPSPARSE_CHECK(hipsparseSetMatDiagType(descr, HIPSPARSE_DIAG_TYPE_NON_UNIT));

    // Matrix info structure
    csrsm2Info_t info;
    HIPSPARSE_CHECK(hipsparseCreateCsrsm2Info(&info));

    // Obtain required buffer size
    size_t buffer_size;
    HIPSPARSE_CHECK(hipsparseDcsrsm2_bufferSizeExt(handle,
                                                   0,
                                                   transA,
                                                   transB,
                                                   m,
                                                   nrhs,
                                                   nnz,
                                                   &alpha,
                                                   descr,
                                                   dcsrVal,
                                                   dcsrRowPtr,
                                                   dcsrColInd,
                                                   dB,
                                                   ldb,
                                                   info,
                                                   solve_policy,
                                                   &buffer_size));

    // Allocate temporary buffer
    void* dbuffer;
    HIP_CHECK(hipMalloc(&dbuffer, buffer_size));

    // Perform analysis step
    HIPSPARSE_CHECK(hipsparseDcsrsm2_analysis(handle,
                                              0,
                                              transA,
                                              transB,
                                              m,
                                              nrhs,
                                              nnz,
                                              &alpha,
                                              descr,
                                              dcsrVal,
                                              dcsrRowPtr,
                                              dcsrColInd,
                                              dB,
                                              ldb,
                                              info,
                                              solve_policy,
                                              dbuffer));

    // Call dcsrsm to perform lower triangular solve LB = B
    HIPSPARSE_CHECK(hipsparseDcsrsm2_solve(handle,
                                           0,
                                           transA,
                                           transB,
                                           m,
                                           nrhs,
                                           nnz,
                                           &alpha,
                                           descr,
                                           dcsrVal,
                                           dcsrRowPtr,
                                           dcsrColInd,
                                           dB,
                                           ldb,
                                           info,
                                           solve_policy,
                                           dbuffer));

    // Check for zero pivots
    int               pivot;
    hipsparseStatus_t status = hipsparseXcsrsm2_zeroPivot(handle, info, &pivot);

    if(status == HIPSPARSE_STATUS_ZERO_PIVOT)
    {
        printf("Found zero pivot in matrix row %d\n", pivot);
    }

    // Copy result back to host
    HIP_CHECK(hipMemcpy(hB, dB, sizeof(double) * m * nrhs, hipMemcpyDeviceToHost));

    printf("hB\n");
    for(int i = 0; i < m * nrhs; i++)
    {
        printf("%f ", hB[i]);
    }
    printf("\n");

    // Clear hipSPARSE
    HIPSPARSE_CHECK(hipsparseDestroyCsrsm2Info(info));
    HIPSPARSE_CHECK(hipsparseDestroyMatDescr(descr));
    HIPSPARSE_CHECK(hipsparseDestroy(handle));

    // Clear device memory
    HIP_CHECK(hipFree(dcsrRowPtr));
    HIP_CHECK(hipFree(dcsrColInd));
    HIP_CHECK(hipFree(dcsrVal));
    HIP_CHECK(hipFree(dB));
    HIP_CHECK(hipFree(dbuffer));

    return 0;
}

program example_hipsparse_csrsm2
    use iso_c_binding
    implicit none

    ! HIP
    interface
        function hipMalloc(ptr, size) &
                bind(c, name = 'hipMalloc')
            use iso_c_binding
            implicit none
            integer(c_int) :: hipMalloc
            type(c_ptr) :: ptr
            integer(c_size_t), value :: size
        end function hipMalloc

        function hipFree(ptr) &
                bind(c, name = 'hipFree')
            use iso_c_binding
            implicit none
            integer(c_int) :: hipFree
            type(c_ptr), value :: ptr
        end function hipFree

        function hipMemcpy(dst, src, size, kind) &
                bind(c, name = 'hipMemcpy')
            use iso_c_binding
            implicit none
            integer(c_int) :: hipMemcpy
            type(c_ptr), value :: dst
            type(c_ptr), intent(in), value :: src
            integer(c_size_t), value :: size
            integer(c_int), value :: kind
        end function hipMemcpy
    end interface

    integer, parameter :: hipMemcpyHostToDevice = 1
    integer, parameter :: hipMemcpyDeviceToHost = 2

    ! hipSPARSE
    interface
        function hipsparseCreate(handle) &
                bind(c, name = 'hipsparseCreate')
            use iso_c_binding
            implicit none
            integer(c_int) :: hipsparseCreate
            type(c_ptr) :: handle
        end function hipsparseCreate

        function hipsparseDestroy(handle) &
                bind(c, name = 'hipsparseDestroy')
            use iso_c_binding
            implicit none
            integer(c_int) :: hipsparseDestroy
            type(c_ptr), value :: handle
        end function hipsparseDestroy

        function hipsparseCreateMatDescr(descr) &
                bind(c, name = 'hipsparseCreateMatDescr')
            use iso_c_binding
            implicit none
            integer(c_int) :: hipsparseCreateMatDescr
            type(c_ptr) :: descr
        end function hipsparseCreateMatDescr

        function hipsparseDestroyMatDescr(descr) &
                bind(c, name = 'hipsparseDestroyMatDescr')
            use iso_c_binding
            implicit none
            integer(c_int) :: hipsparseDestroyMatDescr
            type(c_ptr), value :: descr
        end function hipsparseDestroyMatDescr

        function hipsparseSetMatFillMode(descr, fillMode) &
                bind(c, name = 'hipsparseSetMatFillMode')
            use iso_c_binding
            implicit none
            integer(c_int) :: hipsparseSetMatFillMode
            type(c_ptr), value :: descr
            integer(c_int), value :: fillMode
        end function hipsparseSetMatFillMode

        function hipsparseSetMatDiagType(descr, diagType) &
                bind(c, name = 'hipsparseSetMatDiagType')
            use iso_c_binding
            implicit none
            integer(c_int) :: hipsparseSetMatDiagType
            type(c_ptr), value :: descr
            integer(c_int), value :: diagType
        end function hipsparseSetMatDiagType

        function hipsparseCreateCsrsm2Info(info) &
                bind(c, name = 'hipsparseCreateCsrsm2Info')
            use iso_c_binding
            implicit none
            integer(c_int) :: hipsparseCreateCsrsm2Info
            type(c_ptr) :: info
        end function hipsparseCreateCsrsm2Info

        function hipsparseDestroyCsrsm2Info(info) &
                bind(c, name = 'hipsparseDestroyCsrsm2Info')
            use iso_c_binding
            implicit none
            integer(c_int) :: hipsparseDestroyCsrsm2Info
            type(c_ptr), value :: info
        end function hipsparseDestroyCsrsm2Info

        function hipsparseDcsrsm2_bufferSizeExt(handle, algo, transA, transB, m, nrhs, nnz, &
                                                alpha, descr, csrSortedValA, csrSortedRowPtrA, &
                                                csrSortedColIndA, B, ldb, info, policy, pBufferSize) &
                bind(c, name = 'hipsparseDcsrsm2_bufferSizeExt')
            use iso_c_binding
            implicit none
            integer(c_int) :: hipsparseDcsrsm2_bufferSizeExt
            type(c_ptr), value :: handle
            integer(c_int), value :: algo
            integer(c_int), value :: transA
            integer(c_int), value :: transB
            integer(c_int), value :: m
            integer(c_int), value :: nrhs
            integer(c_int), value :: nnz
            type(c_ptr), intent(in), value :: alpha
            type(c_ptr), value :: descr
            type(c_ptr), intent(in), value :: csrSortedValA
            type(c_ptr), intent(in), value :: csrSortedRowPtrA
            type(c_ptr), intent(in), value :: csrSortedColIndA
            type(c_ptr), value :: B
            integer(c_int), value :: ldb
            type(c_ptr), value :: info
            integer(c_int), value :: policy
            type(c_ptr), value :: pBufferSize
        end function hipsparseDcsrsm2_bufferSizeExt

        function hipsparseDcsrsm2_analysis(handle, algo, transA, transB, m, nrhs, nnz, alpha, &
                                           descr, csrSortedValA, csrSortedRowPtrA, csrSortedColIndA, &
                                           B, ldb, info, policy, pBuffer) &
                bind(c, name = 'hipsparseDcsrsm2_analysis')
            use iso_c_binding
            implicit none
            integer(c_int) :: hipsparseDcsrsm2_analysis
            type(c_ptr), value :: handle
            integer(c_int), value :: algo
            integer(c_int), value :: transA
            integer(c_int), value :: transB
            integer(c_int), value :: m
            integer(c_int), value :: nrhs
            integer(c_int), value :: nnz
            type(c_ptr), intent(in), value :: alpha
            type(c_ptr), value :: descr
            type(c_ptr), intent(in), value :: csrSortedValA
            type(c_ptr), intent(in), value :: csrSortedRowPtrA
            type(c_ptr), intent(in), value :: csrSortedColIndA
            type(c_ptr), value :: B
            integer(c_int), value :: ldb
            type(c_ptr), value :: info
            integer(c_int), value :: policy
            type(c_ptr), value :: pBuffer
        end function hipsparseDcsrsm2_analysis

        function hipsparseDcsrsm2_solve(handle, algo, transA, transB, m, nrhs, nnz, alpha, descr, &
                                        csrSortedValA, csrSortedRowPtrA, csrSortedColIndA, B, ldb, &
                                        info, policy, pBuffer) &
                bind(c, name = 'hipsparseDcsrsm2_solve')
            use iso_c_binding
            implicit none
            integer(c_int) :: hipsparseDcsrsm2_solve
            type(c_ptr), value :: handle
            integer(c_int), value :: algo
            integer(c_int), value :: transA
            integer(c_int), value :: transB
            integer(c_int), value :: m
            integer(c_int), value :: nrhs
            integer(c_int), value :: nnz
            type(c_ptr), intent(in), value :: alpha
            type(c_ptr), value :: descr
            type(c_ptr), intent(in), value :: csrSortedValA
            type(c_ptr), intent(in), value :: csrSortedRowPtrA
            type(c_ptr), intent(in), value :: csrSortedColIndA
            type(c_ptr), value :: B
            integer(c_int), value :: ldb
            type(c_ptr), value :: info
            integer(c_int), value :: policy
            type(c_ptr), value :: pBuffer
        end function hipsparseDcsrsm2_solve

        function hipsparseXcsrsm2_zeroPivot(handle, info, position) &
                bind(c, name = 'hipsparseXcsrsm2_zeroPivot')
            use iso_c_binding
            implicit none
            integer(c_int) :: hipsparseXcsrsm2_zeroPivot
            type(c_ptr), value :: handle
            type(c_ptr), value :: info
            type(c_ptr), value :: position
        end function hipsparseXcsrsm2_zeroPivot
    end interface

    integer, parameter :: HIPSPARSE_OPERATION_NON_TRANSPOSE = 0
    integer, parameter :: HIPSPARSE_FILL_MODE_LOWER = 0
    integer, parameter :: HIPSPARSE_DIAG_TYPE_NON_UNIT = 0
    integer, parameter :: HIPSPARSE_SOLVE_POLICY_NO_LEVEL = 0
    integer, parameter :: HIPSPARSE_STATUS_ZERO_PIVOT = 6

    ! Variables
    type(c_ptr) :: handle
    type(c_ptr) :: descr
    type(c_ptr) :: info
    integer :: i, stat
    integer, target :: pivot
    integer(c_size_t), target :: buffer_size
    
    ! Lower triangular matrix A (4x4) in CSR format
    ! A = ( 1.0  0.0  0.0  0.0 )
    !     ( 2.0  3.0  0.0  0.0 )
    !     ( 4.0  5.0  6.0  0.0 )
    !     ( 7.0  0.0  8.0  9.0 )
    integer, parameter :: m = 4
    integer, parameter :: n = 4
    integer, parameter :: nrhs = 4
    integer, parameter :: nnz = 9
    
    integer, dimension(m+1), target :: hcsrRowPtr = (/0, 1, 3, 6, 9/)
    integer, dimension(nnz), target :: hcsrColInd = (/0, 0, 1, 0, 1, 2, 0, 2, 3/)
    real(c_double), dimension(nnz), target :: hcsrVal = (/1.0d0, 2.0d0, 3.0d0, 4.0d0, 5.0d0, 6.0d0, 7.0d0, 8.0d0, 9.0d0/)
    
    ! Right-hand side matrix B (4x4) - column major
    integer, parameter :: ldb = n
    real(c_double), dimension(n*nrhs), target :: hB = (/ &
        1.0d0, 2.0d0, 3.0d0, 4.0d0, &
        5.0d0, 6.0d0, 7.0d0, 8.0d0, &
        9.0d0, 10.0d0, 11.0d0, 12.0d0, &
        13.0d0, 14.0d0, 15.0d0, 16.0d0 /)
    
    ! Scalar values
    real(c_double), target :: alpha = 1.0d0
    
    ! Device pointers
    type(c_ptr) :: dcsrRowPtr
    type(c_ptr) :: dcsrColInd
    type(c_ptr) :: dcsrVal
    type(c_ptr) :: dB
    type(c_ptr) :: dbuffer

    ! Create hipSPARSE handle
    stat = hipsparseCreate(handle)
    if (stat /= 0) then
        write(*,*) 'Error: hipsparseCreate failed'
        stop
    end if

    ! Create matrix descriptor
    stat = hipsparseCreateMatDescr(descr)
    if (stat /= 0) then
        write(*,*) 'Error: hipsparseCreateMatDescr failed'
        stop
    end if

    ! Set matrix fill mode (lower triangular)
    stat = hipsparseSetMatFillMode(descr, HIPSPARSE_FILL_MODE_LOWER)
    if (stat /= 0) then
        write(*,*) 'Error: hipsparseSetMatFillMode failed'
        stop
    end if

    ! Set matrix diagonal type
    stat = hipsparseSetMatDiagType(descr, HIPSPARSE_DIAG_TYPE_NON_UNIT)
    if (stat /= 0) then
        write(*,*) 'Error: hipsparseSetMatDiagType failed'
        stop
    end if

    ! Create csrsm2 info
    stat = hipsparseCreateCsrsm2Info(info)
    if (stat /= 0) then
        write(*,*) 'Error: hipsparseCreateCsrsm2Info failed'
        stop
    end if

    ! Allocate device memory
    stat = hipMalloc(dcsrRowPtr, int((m + 1) * 4, c_size_t))
    if (stat /= 0) then
        write(*,*) 'Error: hipMalloc dcsrRowPtr failed'
        stop
    end if

    stat = hipMalloc(dcsrColInd, int(nnz * 4, c_size_t))
    if (stat /= 0) then
        write(*,*) 'Error: hipMalloc dcsrColInd failed'
        stop
    end if

    stat = hipMalloc(dcsrVal, int(nnz * 8, c_size_t))
    if (stat /= 0) then
        write(*,*) 'Error: hipMalloc dcsrVal failed'
        stop
    end if

    stat = hipMalloc(dB, int(n * nrhs * 8, c_size_t))
    if (stat /= 0) then
        write(*,*) 'Error: hipMalloc dB failed'
        stop
    end if

    ! Copy data to device
    stat = hipMemcpy(dcsrRowPtr, c_loc(hcsrRowPtr), int((m + 1) * 4, c_size_t), hipMemcpyHostToDevice)
    if (stat /= 0) then
        write(*,*) 'Error: hipMemcpy dcsrRowPtr failed'
        stop
    end if

    stat = hipMemcpy(dcsrColInd, c_loc(hcsrColInd), int(nnz * 4, c_size_t), hipMemcpyHostToDevice)
    if (stat /= 0) then
        write(*,*) 'Error: hipMemcpy dcsrColInd failed'
        stop
    end if

    stat = hipMemcpy(dcsrVal, c_loc(hcsrVal), int(nnz * 8, c_size_t), hipMemcpyHostToDevice)
    if (stat /= 0) then
        write(*,*) 'Error: hipMemcpy dcsrVal failed'
        stop
    end if

    stat = hipMemcpy(dB, c_loc(hB), int(n * nrhs * 8, c_size_t), hipMemcpyHostToDevice)
    if (stat /= 0) then
        write(*,*) 'Error: hipMemcpy dB failed'
        stop
    end if

    ! Obtain required buffer size
    stat = hipsparseDcsrsm2_bufferSizeExt(handle, &
                                          0, &
                                          HIPSPARSE_OPERATION_NON_TRANSPOSE, &
                                          HIPSPARSE_OPERATION_NON_TRANSPOSE, &
                                          m, &
                                          nrhs, &
                                          nnz, &
                                          c_loc(alpha), &
                                          descr, &
                                          dcsrVal, &
                                          dcsrRowPtr, &
                                          dcsrColInd, &
                                          dB, &
                                          ldb, &
                                          info, &
                                          HIPSPARSE_SOLVE_POLICY_NO_LEVEL, &
                                          c_loc(buffer_size))
    if (stat /= 0) then
        write(*,*) 'Error: hipsparseDcsrsm2_bufferSizeExt failed'
        stop
    end if

    ! Allocate temporary buffer
    stat = hipMalloc(dbuffer, buffer_size)
    if (stat /= 0) then
        write(*,*) 'Error: hipMalloc dbuffer failed'
        stop
    end if

    ! Perform analysis step
    stat = hipsparseDcsrsm2_analysis(handle, &
                                     0, &
                                     HIPSPARSE_OPERATION_NON_TRANSPOSE, &
                                     HIPSPARSE_OPERATION_NON_TRANSPOSE, &
                                     m, &
                                     nrhs, &
                                     nnz, &
                                     c_loc(alpha), &
                                     descr, &
                                     dcsrVal, &
                                     dcsrRowPtr, &
                                     dcsrColInd, &
                                     dB, &
                                     ldb, &
                                     info, &
                                     HIPSPARSE_SOLVE_POLICY_NO_LEVEL, &
                                     dbuffer)
    if (stat /= 0) then
        write(*,*) 'Error: hipsparseDcsrsm2_analysis failed'
        stop
    end if

    ! Perform solve: solve Ax = B for x
    stat = hipsparseDcsrsm2_solve(handle, &
                                  0, &
                                  HIPSPARSE_OPERATION_NON_TRANSPOSE, &
                                  HIPSPARSE_OPERATION_NON_TRANSPOSE, &
                                  m, &
                                  nrhs, &
                                  nnz, &
                                  c_loc(alpha), &
                                  descr, &
                                  dcsrVal, &
                                  dcsrRowPtr, &
                                  dcsrColInd, &
                                  dB, &
                                  ldb, &
                                  info, &
                                  HIPSPARSE_SOLVE_POLICY_NO_LEVEL, &
                                  dbuffer)
    if (stat /= 0) then
        write(*,*) 'Error: hipsparseDcsrsm2_solve failed'
        stop
    end if

    ! Check for zero pivots
    stat = hipsparseXcsrsm2_zeroPivot(handle, info, c_loc(pivot))
    if (stat == HIPSPARSE_STATUS_ZERO_PIVOT) then
        write(*,*) 'Found zero pivot in matrix row ', pivot
    end if

    ! Copy result back to host
    stat = hipMemcpy(c_loc(hB), dB, int(m * nrhs * 8, c_size_t), hipMemcpyDeviceToHost)
    if (stat /= 0) then
        write(*,*) 'Error: hipMemcpy hB failed'
        stop
    end if

    ! Print result
    write(*,*) 'Solution:'
    do i = 1, m * nrhs
        write(*,*) hB(i)
    end do

    ! Clean up
    stat = hipFree(dcsrRowPtr)
    stat = hipFree(dcsrColInd)
    stat = hipFree(dcsrVal)
    stat = hipFree(dB)
    stat = hipFree(dbuffer)

    stat = hipsparseDestroyCsrsm2Info(info)
    stat = hipsparseDestroyMatDescr(descr)
    stat = hipsparseDestroy(handle)

end program example_hipsparse_csrsm2

hipsparseXgemmi()#

hipsparseStatus_t hipsparseSgemmi(hipsparseHandle_t handle, int m, int n, int k, int nnz, const float *alpha, const float *A, int lda, const float *cscValB, const int *cscColPtrB, const int *cscRowIndB, const float *beta, float *C, int ldc)#

hipsparseStatus_t hipsparseDgemmi(hipsparseHandle_t handle, int m, int n, int k, int nnz, const double *alpha, const double *A, int lda, const double *cscValB, const int *cscColPtrB, const int *cscRowIndB, const double *beta, double *C, int ldc)#

hipsparseStatus_t hipsparseCgemmi(hipsparseHandle_t handle, int m, int n, int k, int nnz, const hipComplex *alpha, const hipComplex *A, int lda, const hipComplex *cscValB, const int *cscColPtrB, const int *cscRowIndB, const hipComplex *beta, hipComplex *C, int ldc)#

hipsparseStatus_t hipsparseZgemmi(hipsparseHandle_t handle, int m, int n, int k, int nnz, const hipDoubleComplex *alpha, const hipDoubleComplex *A, int lda, const hipDoubleComplex *cscValB, const int *cscColPtrB, const int *cscRowIndB, const hipDoubleComplex *beta, hipDoubleComplex *C, int ldc)#

Dense matrix sparse matrix multiplication using CSC storage format.

hipsparseXgemmi multiplies the scalar \(\alpha\) with a dense column-oriented \(m \times k\) matrix \(A\) and the sparse \(k \times n\) matrix \(B\), defined in CSC storage format and adds the result to the dense column-oriented \(m \times n\) matrix \(C\) that is multiplied by the scalar \(\beta\), such that

\[ C := \alpha \cdot A \cdot B + \beta \cdot C \]

Deprecated:: This function is deprecated when using the CUDA backend (CUDA 11.0+) and will be removed in CUDA 12.0. This deprecation does not apply to the ROCm backend.

Note

This function is non blocking and executed asynchronously with respect to the host. It may return before the actual computation has finished.

Parameters:

handle – [in] handle to the hipsparse library context queue.
m – [in] number of rows of the dense matrix \(A\). Must be non-negative.
n – [in] number of columns of the sparse CSC matrix \(op(B)\) and \(C\). Must be non-negative.
k – [in] number of columns of the dense matrix \(A\). Must be non-negative.
nnz – [in] number of non-zero entries of the sparse CSC matrix \(B\). Must be non-negative.
alpha – [in] scalar \(\alpha\).
A – [in] array of dimension \(lda \times k\) ( \(op(A) == A\)) or \(lda \times m\) ( \(op(A) == A^T\) or \(op(A) == A^H\)).
lda – [in] leading dimension of \(A\), must be at least \(m\) ( \(op(A) == A\)) or \(k\) ( \(op(A) == A^T\) or \(op(A) == A^H\)).
cscValB – [in] array of nnz elements of the sparse CSC matrix \(B\).
cscColPtrB – [in] array of n+1 elements that point to the start of every column of the sparse CSC matrix \(B\).
cscRowIndB – [in] array of nnz elements containing the column indices of the sparse CSC matrix \(B\).
beta – [in] scalar \(\beta\).
C – [inout] array of dimension \(ldc \times n\) that holds the values of \(C\).
ldc – [in] leading dimension of \(C\), must be at least \(m\).

Return values:

HIPSPARSE_STATUS_SUCCESS – the operation completed successfully.
HIPSPARSE_STATUS_NOT_INITIALIZED – handle is not initialized.
HIPSPARSE_STATUS_INVALID_VALUE – handle, alpha or beta is nullptr, m, n, k or nnz is negative, lda or ldc is invalid, or A, cscValB, cscColPtrB, cscRowIndB or C is nullptr.

int main(int argc, char* argv[])
{
    // A, B, and C are m×k, k×n, and m×n
    const int m = 3, n = 5, k = 4;
    const int lda = m, ldc = m;
    const int nnz_A = m * k, nnz_B = 10, nnz_C = m * n;

    // alpha and beta
    float alpha = 0.5f;
    float beta  = 0.25f;

    std::vector<int>   hcscColPtr = {0, 2, 5, 7, 8, 10};
    std::vector<int>   hcscRowInd = {0, 2, 0, 1, 3, 1, 3, 2, 0, 2};
    std::vector<float> hcsc_val   = {1, 6, 2, 4, 9, 5, 2, 7, 3, 8};

    std::vector<float> hA(nnz_A, 1.0f);
    std::vector<float> hC(nnz_C, 1.0f);

    int*   dcscColPtr;
    int*   dcscRowInd;
    float* dcsc_val;
    HIP_CHECK(hipMalloc((void**)&dcscColPtr, sizeof(int) * (n + 1)));
    HIP_CHECK(hipMalloc((void**)&dcscRowInd, sizeof(int) * nnz_B));
    HIP_CHECK(hipMalloc((void**)&dcsc_val, sizeof(float) * nnz_B));

    HIP_CHECK(
        hipMemcpy(dcscColPtr, hcscColPtr.data(), sizeof(int) * (n + 1), hipMemcpyHostToDevice));
    HIP_CHECK(hipMemcpy(dcscRowInd, hcscRowInd.data(), sizeof(int) * nnz_B, hipMemcpyHostToDevice));
    HIP_CHECK(hipMemcpy(dcsc_val, hcsc_val.data(), sizeof(float) * nnz_B, hipMemcpyHostToDevice));

    hipsparseHandle_t handle;
    HIPSPARSE_CHECK(hipsparseCreate(&handle));

    // Allocate memory for the matrix A
    float* dA;
    HIP_CHECK(hipMalloc((void**)&dA, sizeof(float) * nnz_A));
    HIP_CHECK(hipMemcpy(dA, hA.data(), sizeof(float) * nnz_A, hipMemcpyHostToDevice));

    // Allocate memory for the resulting matrix C
    float* dC;
    HIP_CHECK(hipMalloc((void**)&dC, sizeof(float) * nnz_C));
    HIP_CHECK(hipMemcpy(dC, hC.data(), sizeof(float) * nnz_C, hipMemcpyHostToDevice));

    // Perform operation
    HIPSPARSE_CHECK(hipsparseSgemmi(
        handle, m, n, k, nnz_B, &alpha, dA, lda, dcsc_val, dcscColPtr, dcscRowInd, &beta, dC, ldc));

    // Copy device to host
    HIP_CHECK(hipMemcpy(hC.data(), dC, sizeof(float) * nnz_C, hipMemcpyDeviceToHost));

    std::cout << "hC" << std::endl;
    for(int i = 0; i < nnz_C; i++)
    {
        std::cout << hC[i] << " ";
    }
    std::cout << std::endl;

    // Destroy matrix descriptors and handles
    HIPSPARSE_CHECK(hipsparseDestroy(handle));

    HIP_CHECK(hipFree(dcscColPtr));
    HIP_CHECK(hipFree(dcscRowInd));
    HIP_CHECK(hipFree(dcsc_val));
    HIP_CHECK(hipFree(dA));
    HIP_CHECK(hipFree(dC));

    return 0;
}

int main(int argc, char* argv[])
{
    // A, B, and C are m×k, k×n, and m×n
    const int m = 3, n = 5, k = 4;
    const int lda = m, ldc = m;
    const int nnz_A = m * k, nnz_B = 10, nnz_C = m * n;

    // alpha and beta
    float alpha = 0.5;
    float beta  = 0.25;

    int   hcscColPtr[] = {0, 2, 5, 7, 8, 10};
    int   hcscRowInd[] = {0, 2, 0, 1, 3, 1, 3, 2, 0, 2};
    float hcsc_val[]   = {1, 6, 2, 4, 9, 5, 2, 7, 3, 8};

    float* hA = (float*)malloc(nnz_A * sizeof(float));
    float* hC = (float*)malloc(nnz_C * sizeof(float));
    
    // Initialize A and C to 1.0
    for(int i = 0; i < nnz_A; i++)
    {
        hA[i] = 1.0f;
    }
    for(int i = 0; i < nnz_C; i++)
    {
        hC[i] = 1.0f;
    }

    int*   dcscColPtr;
    int*   dcscRowInd;
    float* dcsc_val;
    HIP_CHECK(hipMalloc((void**)&dcscColPtr, sizeof(int) * (n + 1)));
    HIP_CHECK(hipMalloc((void**)&dcscRowInd, sizeof(int) * nnz_B));
    HIP_CHECK(hipMalloc((void**)&dcsc_val, sizeof(float) * nnz_B));

    HIP_CHECK(hipMemcpy(dcscColPtr, hcscColPtr, sizeof(int) * (n + 1), hipMemcpyHostToDevice));
    HIP_CHECK(hipMemcpy(dcscRowInd, hcscRowInd, sizeof(int) * nnz_B, hipMemcpyHostToDevice));
    HIP_CHECK(hipMemcpy(dcsc_val, hcsc_val, sizeof(float) * nnz_B, hipMemcpyHostToDevice));

    hipsparseHandle_t handle;
    HIPSPARSE_CHECK(hipsparseCreate(&handle));

    // Allocate memory for the matrix A
    float* dA;
    HIP_CHECK(hipMalloc((void**)&dA, sizeof(float) * nnz_A));
    HIP_CHECK(hipMemcpy(dA, hA, sizeof(float) * nnz_A, hipMemcpyHostToDevice));

    // Allocate memory for the resulting matrix C
    float* dC;
    HIP_CHECK(hipMalloc((void**)&dC, sizeof(float) * nnz_C));
    HIP_CHECK(hipMemcpy(dC, hC, sizeof(float) * nnz_C, hipMemcpyHostToDevice));

    // Perform operation
    HIPSPARSE_CHECK(hipsparseSgemmi(
        handle, m, n, k, nnz_B, &alpha, dA, lda, dcsc_val, dcscColPtr, dcscRowInd, &beta, dC, ldc));

    // Copy device to host
    HIP_CHECK(hipMemcpy(hC, dC, sizeof(float) * nnz_C, hipMemcpyDeviceToHost));

    printf("hC\n");
    for(int i = 0; i < nnz_C; i++)
    {
        printf("%f ", hC[i]);
    }
    printf("\n");

    // Destroy matrix descriptors and handles
    HIPSPARSE_CHECK(hipsparseDestroy(handle));

    HIP_CHECK(hipFree(dcscColPtr));
    HIP_CHECK(hipFree(dcscRowInd));
    HIP_CHECK(hipFree(dcsc_val));
    HIP_CHECK(hipFree(dA));
    HIP_CHECK(hipFree(dC));

    return 0;
}

program example_hipsparse_gemmi
    use iso_c_binding
    implicit none

    ! HIP
    interface
        function hipMalloc(ptr, size) &
                bind(c, name = 'hipMalloc')
            use iso_c_binding
            implicit none
            integer(c_int) :: hipMalloc
            type(c_ptr) :: ptr
            integer(c_size_t), value :: size
        end function hipMalloc

        function hipFree(ptr) &
                bind(c, name = 'hipFree')
            use iso_c_binding
            implicit none
            integer(c_int) :: hipFree
            type(c_ptr), value :: ptr
        end function hipFree

        function hipMemcpy(dst, src, size, kind) &
                bind(c, name = 'hipMemcpy')
            use iso_c_binding
            implicit none
            integer(c_int) :: hipMemcpy
            type(c_ptr), value :: dst
            type(c_ptr), intent(in), value :: src
            integer(c_size_t), value :: size
            integer(c_int), value :: kind
        end function hipMemcpy
    end interface

    integer, parameter :: hipMemcpyHostToDevice = 1
    integer, parameter :: hipMemcpyDeviceToHost = 2

    ! hipSPARSE
    interface
        function hipsparseCreate(handle) &
                bind(c, name = 'hipsparseCreate')
            use iso_c_binding
            implicit none
            integer(c_int) :: hipsparseCreate
            type(c_ptr) :: handle
        end function hipsparseCreate

        function hipsparseDestroy(handle) &
                bind(c, name = 'hipsparseDestroy')
            use iso_c_binding
            implicit none
            integer(c_int) :: hipsparseDestroy
            type(c_ptr), value :: handle
        end function hipsparseDestroy

        function hipsparseSgemmi(handle, m, n, k, nnz, alpha, A, lda, cscValB, cscColPtrB, &
                                 cscRowIndB, beta, C, ldc) &
                bind(c, name = 'hipsparseSgemmi')
            use iso_c_binding
            implicit none
            integer(c_int) :: hipsparseSgemmi
            type(c_ptr), value :: handle
            integer(c_int), value :: m
            integer(c_int), value :: n
            integer(c_int), value :: k
            integer(c_int), value :: nnz
            type(c_ptr), intent(in), value :: alpha
            type(c_ptr), intent(in), value :: A
            integer(c_int), value :: lda
            type(c_ptr), intent(in), value :: cscValB
            type(c_ptr), intent(in), value :: cscColPtrB
            type(c_ptr), intent(in), value :: cscRowIndB
            type(c_ptr), intent(in), value :: beta
            type(c_ptr), value :: C
            integer(c_int), value :: ldc
        end function hipsparseSgemmi
    end interface

    ! Variables
    type(c_ptr) :: handle
    integer :: i, stat
    
    ! C = alpha * A * B + beta * C
    ! A is dense (m x k), B is sparse CSC (k x n), C is dense (m x n)
    integer, parameter :: m = 3
    integer, parameter :: n = 5
    integer, parameter :: k = 4
    integer, parameter :: lda = m
    integer, parameter :: ldc = m
    integer, parameter :: nnz_A = m * k
    integer, parameter :: nnz_B = 10
    integer, parameter :: nnz_C = m * n
    
    ! Sparse matrix B in CSC format
    integer, dimension(n+1), target :: hcscColPtr = (/0, 2, 5, 7, 8, 10/)
    integer, dimension(nnz_B), target :: hcscRowInd = (/0, 2, 0, 1, 3, 1, 3, 2, 0, 2/)
    real(c_float), dimension(nnz_B), target :: hcsc_val = (/1.0, 6.0, 2.0, 4.0, 9.0, 5.0, 2.0, 7.0, 3.0, 8.0/)
    
    ! Dense matrices A and C
    real(c_float), dimension(nnz_A), target :: hA
    real(c_float), dimension(nnz_C), target :: hC
    
    ! Scalar values
    real(c_float), target :: alpha = 0.5
    real(c_float), target :: beta = 0.25
    
    ! Device pointers
    type(c_ptr) :: dcscColPtr
    type(c_ptr) :: dcscRowInd
    type(c_ptr) :: dcsc_val
    type(c_ptr) :: dA
    type(c_ptr) :: dC
    
    ! Initialize dense matrices with 1.0
    do i = 1, nnz_A
        hA(i) = 1.0
    end do
    
    do i = 1, nnz_C
        hC(i) = 1.0
    end do

    ! Create hipSPARSE handle
    stat = hipsparseCreate(handle)
    if (stat /= 0) then
        write(*,*) 'Error: hipsparseCreate failed'
        stop
    end if

    ! Allocate device memory for CSC matrix B
    stat = hipMalloc(dcscColPtr, int((n + 1) * 4, c_size_t))
    if (stat /= 0) then
        write(*,*) 'Error: hipMalloc dcscColPtr failed'
        stop
    end if

    stat = hipMalloc(dcscRowInd, int(nnz_B * 4, c_size_t))
    if (stat /= 0) then
        write(*,*) 'Error: hipMalloc dcscRowInd failed'
        stop
    end if

    stat = hipMalloc(dcsc_val, int(nnz_B * 4, c_size_t))
    if (stat /= 0) then
        write(*,*) 'Error: hipMalloc dcsc_val failed'
        stop
    end if

    ! Allocate device memory for dense matrices A and C
    stat = hipMalloc(dA, int(nnz_A * 4, c_size_t))
    if (stat /= 0) then
        write(*,*) 'Error: hipMalloc dA failed'
        stop
    end if

    stat = hipMalloc(dC, int(nnz_C * 4, c_size_t))
    if (stat /= 0) then
        write(*,*) 'Error: hipMalloc dC failed'
        stop
    end if

    ! Copy data to device
    stat = hipMemcpy(dcscColPtr, c_loc(hcscColPtr), int((n + 1) * 4, c_size_t), hipMemcpyHostToDevice)
    if (stat /= 0) then
        write(*,*) 'Error: hipMemcpy dcscColPtr failed'
        stop
    end if

    stat = hipMemcpy(dcscRowInd, c_loc(hcscRowInd), int(nnz_B * 4, c_size_t), hipMemcpyHostToDevice)
    if (stat /= 0) then
        write(*,*) 'Error: hipMemcpy dcscRowInd failed'
        stop
    end if

    stat = hipMemcpy(dcsc_val, c_loc(hcsc_val), int(nnz_B * 4, c_size_t), hipMemcpyHostToDevice)
    if (stat /= 0) then
        write(*,*) 'Error: hipMemcpy dcsc_val failed'
        stop
    end if

    stat = hipMemcpy(dA, c_loc(hA), int(nnz_A * 4, c_size_t), hipMemcpyHostToDevice)
    if (stat /= 0) then
        write(*,*) 'Error: hipMemcpy dA failed'
        stop
    end if

    stat = hipMemcpy(dC, c_loc(hC), int(nnz_C * 4, c_size_t), hipMemcpyHostToDevice)
    if (stat /= 0) then
        write(*,*) 'Error: hipMemcpy dC failed'
        stop
    end if

    ! Perform operation: C = alpha * A * B + beta * C
    stat = hipsparseSgemmi(handle, &
                           m, &
                           n, &
                           k, &
                           nnz_B, &
                           c_loc(alpha), &
                           dA, &
                           lda, &
                           dcsc_val, &
                           dcscColPtr, &
                           dcscRowInd, &
                           c_loc(beta), &
                           dC, &
                           ldc)
    if (stat /= 0) then
        write(*,*) 'Error: hipsparseSgemmi failed'
        stop
    end if

    ! Copy result back to host
    stat = hipMemcpy(c_loc(hC), dC, int(nnz_C * 4, c_size_t), hipMemcpyDeviceToHost)
    if (stat /= 0) then
        write(*,*) 'Error: hipMemcpy hC failed'
        stop
    end if

    ! Print result
    write(*,*) 'hC:'
    do i = 1, nnz_C
        write(*,*) hC(i)
    end do

    ! Clean up
    stat = hipFree(dcscColPtr)
    stat = hipFree(dcscRowInd)
    stat = hipFree(dcsc_val)
    stat = hipFree(dA)
    stat = hipFree(dC)

    stat = hipsparseDestroy(handle)

end program example_hipsparse_gemmi

Sparse level 3 functions

Contents

Sparse level 3 functions#

hipsparseXbsrmm()#

hipsparseXcsrmm()#

hipsparseXcsrmm2()#

hipsparseXbsrsm2_zeroPivot()#

hipsparseXbsrsm2_bufferSize()#

hipsparseXbsrsm2_analysis()#

hipsparseXbsrsm2_solve()#

hipsparseXcsrsm2_zeroPivot()#

hipsparseXcsrsm2_bufferSizeExt()#

hipsparseXcsrsm2_analysis()#

hipsparseXcsrsm2_solve()#

hipsparseXgemmi()#