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.

rocsparse_bsrmm()

rocsparse_status rocsparse_sbsrmm(rocsparse_handle handle, rocsparse_direction dir, rocsparse_operation trans_A, rocsparse_operation trans_B, rocsparse_int mb, rocsparse_int n, rocsparse_int kb, rocsparse_int nnzb, const float *alpha, const rocsparse_mat_descr descr, const float *bsr_val, const rocsparse_int *bsr_row_ptr, const rocsparse_int *bsr_col_ind, rocsparse_int block_dim, const float *B, rocsparse_int ldb, const float *beta, float *C, rocsparse_int ldc)

rocsparse_status rocsparse_dbsrmm(rocsparse_handle handle, rocsparse_direction dir, rocsparse_operation trans_A, rocsparse_operation trans_B, rocsparse_int mb, rocsparse_int n, rocsparse_int kb, rocsparse_int nnzb, const double *alpha, const rocsparse_mat_descr descr, const double *bsr_val, const rocsparse_int *bsr_row_ptr, const rocsparse_int *bsr_col_ind, rocsparse_int block_dim, const double *B, rocsparse_int ldb, const double *beta, double *C, rocsparse_int ldc)

rocsparse_status rocsparse_cbsrmm(rocsparse_handle handle, rocsparse_direction dir, rocsparse_operation trans_A, rocsparse_operation trans_B, rocsparse_int mb, rocsparse_int n, rocsparse_int kb, rocsparse_int nnzb, const rocsparse_float_complex *alpha, const rocsparse_mat_descr descr, const rocsparse_float_complex *bsr_val, const rocsparse_int *bsr_row_ptr, const rocsparse_int *bsr_col_ind, rocsparse_int block_dim, const rocsparse_float_complex *B, rocsparse_int ldb, const rocsparse_float_complex *beta, rocsparse_float_complex *C, rocsparse_int ldc)

rocsparse_status rocsparse_zbsrmm(rocsparse_handle handle, rocsparse_direction dir, rocsparse_operation trans_A, rocsparse_operation trans_B, rocsparse_int mb, rocsparse_int n, rocsparse_int kb, rocsparse_int nnzb, const rocsparse_double_complex *alpha, const rocsparse_mat_descr descr, const rocsparse_double_complex *bsr_val, const rocsparse_int *bsr_row_ptr, const rocsparse_int *bsr_col_ind, rocsparse_int block_dim, const rocsparse_double_complex *B, rocsparse_int ldb, const rocsparse_double_complex *beta, rocsparse_double_complex *C, rocsparse_int ldc)

Sparse matrix dense matrix multiplication using BSR storage format.

rocsparse_bsrmm 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 trans_A == rocsparse_operation_none} \\ \end{array} \right. \end{split}\]

and

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

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

Example

This example multiplies a BSR matrix with a column-oriented dense matrix.

int main()
{
    //     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

    rocsparse_int       block_dim = 2;
    rocsparse_int       mb        = 2;
    rocsparse_int       kb        = 3;
    rocsparse_int       nnzb      = 4;
    rocsparse_direction dir       = rocsparse_direction_row;

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

    std::vector<rocsparse_int> hbsr_row_ptr = {0, 2, 4};
    std::vector<rocsparse_int> hbsr_col_ind = {0, 1, 1, 2};
    std::vector<float>         hbsr_val     = {1, 2, 0, 4, 0, 3, 5, 0, 0, 7, 1, 2, 8, 0, 4, 1};

    // Set dimension n of B
    rocsparse_int n = 64;
    rocsparse_int m = mb * block_dim;
    rocsparse_int k = kb * block_dim;

    // Allocate and generate column-oriented dense matrix B
    std::vector<float> hB(k * n);
    for(rocsparse_int i = 0; i < k * n; ++i)
    {
        hB[i] = static_cast<float>(rand()) / static_cast<float>(RAND_MAX);
    }
    std::vector<float> hC(m * n);

    int*   dbsr_row_ptr;
    int*   dbsr_col_ind;
    float* dbsr_val;
    float* dB;
    float* dC;
    HIP_CHECK(hipMalloc(&dbsr_row_ptr, sizeof(int) * (mb + 1)));
    HIP_CHECK(hipMalloc(&dbsr_col_ind, sizeof(int) * nnzb));
    HIP_CHECK(hipMalloc(&dbsr_val, sizeof(float) * nnzb * block_dim * block_dim));
    HIP_CHECK(hipMalloc(&dB, sizeof(float) * k * n));
    HIP_CHECK(hipMalloc(&dC, sizeof(float) * m * n));

    HIP_CHECK(hipMemcpy(
        dbsr_row_ptr, hbsr_row_ptr.data(), sizeof(int) * (mb + 1), hipMemcpyHostToDevice));
    HIP_CHECK(
        hipMemcpy(dbsr_col_ind, hbsr_col_ind.data(), sizeof(int) * nnzb, hipMemcpyHostToDevice));
    HIP_CHECK(hipMemcpy(dbsr_val,
                        hbsr_val.data(),
                        sizeof(float) * nnzb * block_dim * block_dim,
                        hipMemcpyHostToDevice));
    HIP_CHECK(hipMemcpy(dB, hB.data(), sizeof(float) * k * n, hipMemcpyHostToDevice));
    HIP_CHECK(hipMemcpy(dC, hC.data(), sizeof(float) * m * n, hipMemcpyHostToDevice));

    // Create rocSPARSE handle
    rocsparse_handle handle;
    ROCSPARSE_CHECK(rocsparse_create_handle(&handle));

    // Create matrix descriptor
    rocsparse_mat_descr descr;
    ROCSPARSE_CHECK(rocsparse_create_mat_descr(&descr));

    // Perform the matrix multiplication
    ROCSPARSE_CHECK(rocsparse_sbsrmm(handle,
                                     dir,
                                     rocsparse_operation_none,
                                     rocsparse_operation_none,
                                     mb,
                                     n,
                                     kb,
                                     nnzb,
                                     &alpha,
                                     descr,
                                     dbsr_val,
                                     dbsr_row_ptr,
                                     dbsr_col_ind,
                                     block_dim,
                                     dB,
                                     k,
                                     &beta,
                                     dC,
                                     m));

    HIP_CHECK(hipMemcpy(hC.data(), dC, sizeof(float) * m * n, hipMemcpyDeviceToHost));

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

    // Clean up
    ROCSPARSE_CHECK(rocsparse_destroy_mat_descr(descr));
    ROCSPARSE_CHECK(rocsparse_destroy_handle(handle));

    // Clear up on device
    HIP_CHECK(hipFree(dbsr_row_ptr));
    HIP_CHECK(hipFree(dbsr_col_ind));
    HIP_CHECK(hipFree(dbsr_val));
    HIP_CHECK(hipFree(dB));
    HIP_CHECK(hipFree(dC));

    return 0;
}

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 trans_A == rocsparse_operation_none is supported.

Note

This routine supports execution in a hipGraph context.

Parameters:

• handle – [in] handle to the rocsparse library context queue.

• dir – [in] the storage format of the blocks. Can be rocsparse_direction_row or rocsparse_direction_column.

• trans_A – [in] matrix \(A\) operation type. Currently, only rocsparse_operation_none is supported.

• trans_B – [in] matrix \(B\) operation type. Currently, only rocsparse_operation_none and rocsparse_operation_transpose are supported.

• mb – [in] number of block rows of the sparse BSR matrix \(A\).

• n – [in] number of columns of the column-oriented dense matrix \(op(B)\) and \(C\).

• kb – [in] number of block columns of the sparse BSR matrix \(A\).

• nnzb – [in] number of non-zero blocks of the sparse BSR matrix \(A\).

• alpha – [in] scalar \(\alpha\).

• descr – [in] descriptor of the sparse BSR matrix \(A\). Currently, only rocsparse_matrix_type_general is supported.

• bsr_val – [in] array of nnzb*block_dim*block_dim elements of the sparse BSR matrix \(A\).

• bsr_row_ptr – [in] array of mb+1 elements that point to the start of every block row of the sparse BSR matrix \(A\).

• bsr_col_ind – [in] array of nnzb elements containing the block column indices of the sparse BSR matrix \(A\).

• block_dim – [in] size of the blocks in the sparse BSR matrix.

• B – [in] column-oriented dense matrix of dimension \(ldb \times n\) ( \(op(B) == B\)), \(ldb \times k\) otherwise.

• ldb – [in] leading dimension of \(B\), must be at least \(\max{(1, k)}\) ( \( op(B) == B\)) where \(k = block\_dim \times kb\), \(\max{(1, n)}\) otherwise.

• beta – [in] scalar \(\beta\).

• C – [inout] column-oriented dense matrix of dimension \(ldc \times n\).

• ldc – [in] leading dimension of \(C\), must be at least \(\max{(1, m)}\) ( \( op(A) == A\)) where \(m = block\_dim \times mb\), \(\max{(1, k)}\) where \(k = block\_dim \times kb\) otherwise.

Return values:

• rocsparse_status_success – the operation completed successfully.

• rocsparse_status_invalid_handle – the library context was not initialized.

• rocsparse_status_invalid_size – mb, n, kb, nnzb, ldb or ldc is invalid.

• rocsparse_status_invalid_pointer – descr, alpha, bsr_val, bsr_row_ptr, bsr_col_ind, B, beta or C pointer is invalid.

• rocsparse_status_arch_mismatch – the device is not supported.

• rocsparse_status_not_implemented – trans_A != rocsparse_operation_none or trans_B == rocsparse_operation_conjugate_transpose or rocsparse_matrix_type != rocsparse_matrix_type_general.

rocsparse_gebsrmm()

rocsparse_status rocsparse_sgebsrmm(rocsparse_handle handle, rocsparse_direction dir, rocsparse_operation trans_A, rocsparse_operation trans_B, rocsparse_int mb, rocsparse_int n, rocsparse_int kb, rocsparse_int nnzb, const float *alpha, const rocsparse_mat_descr descr, const float *bsr_val, const rocsparse_int *bsr_row_ptr, const rocsparse_int *bsr_col_ind, rocsparse_int row_block_dim, rocsparse_int col_block_dim, const float *B, rocsparse_int ldb, const float *beta, float *C, rocsparse_int ldc)

rocsparse_status rocsparse_dgebsrmm(rocsparse_handle handle, rocsparse_direction dir, rocsparse_operation trans_A, rocsparse_operation trans_B, rocsparse_int mb, rocsparse_int n, rocsparse_int kb, rocsparse_int nnzb, const double *alpha, const rocsparse_mat_descr descr, const double *bsr_val, const rocsparse_int *bsr_row_ptr, const rocsparse_int *bsr_col_ind, rocsparse_int row_block_dim, rocsparse_int col_block_dim, const double *B, rocsparse_int ldb, const double *beta, double *C, rocsparse_int ldc)

rocsparse_status rocsparse_cgebsrmm(rocsparse_handle handle, rocsparse_direction dir, rocsparse_operation trans_A, rocsparse_operation trans_B, rocsparse_int mb, rocsparse_int n, rocsparse_int kb, rocsparse_int nnzb, const rocsparse_float_complex *alpha, const rocsparse_mat_descr descr, const rocsparse_float_complex *bsr_val, const rocsparse_int *bsr_row_ptr, const rocsparse_int *bsr_col_ind, rocsparse_int row_block_dim, rocsparse_int col_block_dim, const rocsparse_float_complex *B, rocsparse_int ldb, const rocsparse_float_complex *beta, rocsparse_float_complex *C, rocsparse_int ldc)

rocsparse_status rocsparse_zgebsrmm(rocsparse_handle handle, rocsparse_direction dir, rocsparse_operation trans_A, rocsparse_operation trans_B, rocsparse_int mb, rocsparse_int n, rocsparse_int kb, rocsparse_int nnzb, const rocsparse_double_complex *alpha, const rocsparse_mat_descr descr, const rocsparse_double_complex *bsr_val, const rocsparse_int *bsr_row_ptr, const rocsparse_int *bsr_col_ind, rocsparse_int row_block_dim, rocsparse_int col_block_dim, const rocsparse_double_complex *B, rocsparse_int ldb, const rocsparse_double_complex *beta, rocsparse_double_complex *C, rocsparse_int ldc)

Sparse matrix dense matrix multiplication using GEneral BSR storage format.

rocsparse_gebsrmm multiplies the scalar \(\alpha\) with a sparse \(m \times k\) matrix \(A\), defined in GEneral 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 trans_A == rocsparse_operation_none} \\ \end{array} \right. \end{split}\]

and

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

and where \(k = col\_block\_dim \times kb\) and \(m = row\_block\_dim \times mb\).

Example

This example multiplies a GEneral BSR matrix with a column-oriented dense matrix.

int main()
{
    //     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

    rocsparse_int       row_block_dim = 2;
    rocsparse_int       col_block_dim = 3;
    rocsparse_int       mb            = 2;
    rocsparse_int       kb            = 2;
    rocsparse_int       nnzb          = 4;
    rocsparse_direction dir           = rocsparse_direction_row;

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

    std::vector<rocsparse_int> hbsr_row_ptr = {0, 2, 4};
    std::vector<rocsparse_int> hbsr_col_ind = {0, 1, 0, 1};
    std::vector<float>         hbsr_val
        = {1, 2, 0, 0, 4, 5, 3, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 1, 7, 8, 0, 2, 4, 1};

    // Set dimension n of B
    rocsparse_int n = 64;
    rocsparse_int m = mb * row_block_dim;
    rocsparse_int k = kb * col_block_dim;

    // Allocate and generate column-oriented dense matrix B
    std::vector<float> hB(k * n);
    for(rocsparse_int i = 0; i < k * n; ++i)
    {
        hB[i] = static_cast<float>(rand()) / static_cast<float>(RAND_MAX);
    }
    std::vector<float> hC(m * n);

    int*   dbsr_row_ptr;
    int*   dbsr_col_ind;
    float* dbsr_val;
    float* dB;
    float* dC;
    HIP_CHECK(hipMalloc(&dbsr_row_ptr, sizeof(int) * (mb + 1)));
    HIP_CHECK(hipMalloc(&dbsr_col_ind, sizeof(int) * nnzb));
    HIP_CHECK(hipMalloc(&dbsr_val, sizeof(float) * nnzb * row_block_dim * col_block_dim));
    HIP_CHECK(hipMalloc(&dB, sizeof(float) * k * n));
    HIP_CHECK(hipMalloc(&dC, sizeof(float) * m * n));

    HIP_CHECK(hipMemcpy(
        dbsr_row_ptr, hbsr_row_ptr.data(), sizeof(int) * (mb + 1), hipMemcpyHostToDevice));
    HIP_CHECK(
        hipMemcpy(dbsr_col_ind, hbsr_col_ind.data(), sizeof(int) * nnzb, hipMemcpyHostToDevice));
    HIP_CHECK(hipMemcpy(dbsr_val,
                        hbsr_val.data(),
                        sizeof(float) * nnzb * row_block_dim * col_block_dim,
                        hipMemcpyHostToDevice));
    HIP_CHECK(hipMemcpy(dB, hB.data(), sizeof(float) * k * n, hipMemcpyHostToDevice));
    HIP_CHECK(hipMemcpy(dC, hC.data(), sizeof(float) * m * n, hipMemcpyHostToDevice));

    // Create rocSPARSE handle
    rocsparse_handle handle;
    ROCSPARSE_CHECK(rocsparse_create_handle(&handle));

    // Create matrix descriptor
    rocsparse_mat_descr descr;
    ROCSPARSE_CHECK(rocsparse_create_mat_descr(&descr));

    // Perform the matrix multiplication
    ROCSPARSE_CHECK(rocsparse_sgebsrmm(handle,
                                       dir,
                                       rocsparse_operation_none,
                                       rocsparse_operation_none,
                                       mb,
                                       n,
                                       kb,
                                       nnzb,
                                       &alpha,
                                       descr,
                                       dbsr_val,
                                       dbsr_row_ptr,
                                       dbsr_col_ind,
                                       row_block_dim,
                                       col_block_dim,
                                       dB,
                                       k,
                                       &beta,
                                       dC,
                                       m));

    HIP_CHECK(hipMemcpy(hC.data(), dC, sizeof(float) * m * n, hipMemcpyDeviceToHost));

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

    // Clean up
    ROCSPARSE_CHECK(rocsparse_destroy_mat_descr(descr));
    ROCSPARSE_CHECK(rocsparse_destroy_handle(handle));

    // Clear up on device
    HIP_CHECK(hipFree(dbsr_row_ptr));
    HIP_CHECK(hipFree(dbsr_col_ind));
    HIP_CHECK(hipFree(dbsr_val));
    HIP_CHECK(hipFree(dB));
    HIP_CHECK(hipFree(dC));

    return 0;
}

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 trans_A == rocsparse_operation_none is supported.

Note

This routine supports execution in a hipGraph context.

Parameters:

• handle – [in] handle to the rocsparse library context queue.

• dir – [in] the storage format of the blocks. Can be rocsparse_direction_row or rocsparse_direction_column.

• trans_A – [in] matrix \(A\) operation type. Currently, only rocsparse_operation_none is supported.

• trans_B – [in] matrix \(B\) operation type. Currently, only rocsparse_operation_none and rocsparse_operation_transpose are supported.

• mb – [in] number of block rows of the sparse GEneral BSR matrix \(A\).

• n – [in] number of columns of the column-oriented dense matrix \(op(B)\) and \(C\).

• kb – [in] number of block columns of the sparse GEneral BSR matrix \(A\).

• nnzb – [in] number of non-zero blocks of the sparse GEneral BSR matrix \(A\).

• alpha – [in] scalar \(\alpha\).

• descr – [in] descriptor of the sparse GEneral BSR matrix \(A\). Currently, only rocsparse_matrix_type_general is supported.

• bsr_val – [in] array of nnzb*row_block_dim*col_block_dim elements of the sparse GEneral BSR matrix \(A\).

• bsr_row_ptr – [in] array of mb+1 elements that point to the start of every block row of the sparse GEneral BSR matrix \(A\).

• bsr_col_ind – [in] array of nnzb elements containing the block column indices of the sparse GEneral BSR matrix \(A\).

• row_block_dim – [in] row size of the blocks in the sparse GEneral BSR matrix.

• col_block_dim – [in] column size of the blocks in the sparse GEneral BSR matrix.

• B – [in] column-oriented dense matrix of dimension \(ldb \times n\) ( \(op(B) == B\)), \(ldb \times k\) otherwise.

• ldb – [in] leading dimension of \(B\), must be at least \(\max{(1, k)}\) ( \( op(B) == B\)) where \(k = col\_block\_dim \times kb\), \(\max{(1, n)}\) otherwise.

• beta – [in] scalar \(\beta\).

• C – [inout] column-oriented dense matrix of dimension \(ldc \times n\).

• ldc – [in] leading dimension of \(C\), must be at least \(\max{(1, m)}\) ( \( op(A) == A\)) where \(m = row\_block\_dim \times mb\), \(\max{(1, k)}\) where \(k = col\_block\_dim \times kb\) otherwise.

Return values:

• rocsparse_status_success – the operation completed successfully.

• rocsparse_status_invalid_handle – the library context was not initialized.

• rocsparse_status_invalid_size – mb, n, kb, nnzb, ldb, ldc, row_block_dim or col_block_dim is invalid.

• rocsparse_status_invalid_pointer – descr, alpha, bsr_val, bsr_row_ptr, bsr_col_ind, B, beta or C pointer is invalid.

• rocsparse_status_arch_mismatch – the device is not supported.

• rocsparse_status_not_implemented – trans_A != rocsparse_operation_none or trans_B == rocsparse_operation_conjugate_transpose or rocsparse_matrix_type != rocsparse_matrix_type_general.

rocsparse_csrmm()

rocsparse_status rocsparse_scsrmm(rocsparse_handle handle, rocsparse_operation trans_A, rocsparse_operation trans_B, rocsparse_int m, rocsparse_int n, rocsparse_int k, rocsparse_int nnz, const float *alpha, const rocsparse_mat_descr descr, const float *csr_val, const rocsparse_int *csr_row_ptr, const rocsparse_int *csr_col_ind, const float *B, rocsparse_int ldb, const float *beta, float *C, rocsparse_int ldc)

rocsparse_status rocsparse_dcsrmm(rocsparse_handle handle, rocsparse_operation trans_A, rocsparse_operation trans_B, rocsparse_int m, rocsparse_int n, rocsparse_int k, rocsparse_int nnz, const double *alpha, const rocsparse_mat_descr descr, const double *csr_val, const rocsparse_int *csr_row_ptr, const rocsparse_int *csr_col_ind, const double *B, rocsparse_int ldb, const double *beta, double *C, rocsparse_int ldc)

rocsparse_status rocsparse_ccsrmm(rocsparse_handle handle, rocsparse_operation trans_A, rocsparse_operation trans_B, rocsparse_int m, rocsparse_int n, rocsparse_int k, rocsparse_int nnz, const rocsparse_float_complex *alpha, const rocsparse_mat_descr descr, const rocsparse_float_complex *csr_val, const rocsparse_int *csr_row_ptr, const rocsparse_int *csr_col_ind, const rocsparse_float_complex *B, rocsparse_int ldb, const rocsparse_float_complex *beta, rocsparse_float_complex *C, rocsparse_int ldc)

rocsparse_status rocsparse_zcsrmm(rocsparse_handle handle, rocsparse_operation trans_A, rocsparse_operation trans_B, rocsparse_int m, rocsparse_int n, rocsparse_int k, rocsparse_int nnz, const rocsparse_double_complex *alpha, const rocsparse_mat_descr descr, const rocsparse_double_complex *csr_val, const rocsparse_int *csr_row_ptr, const rocsparse_int *csr_col_ind, const rocsparse_double_complex *B, rocsparse_int ldb, const rocsparse_double_complex *beta, rocsparse_double_complex *C, rocsparse_int ldc)

Sparse matrix dense matrix multiplication using CSR storage format.

rocsparse_csrmm 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 trans_A == rocsparse_operation_none} \\ A^T, & \text{if trans_A == rocsparse_operation_transpose} \\ A^H, & \text{if trans_A == rocsparse_operation_conjugate_transpose} \end{array} \right. \end{split}\]

and

\[\begin{split} op(B) = \left\{ \begin{array}{ll} B, & \text{if trans_B == rocsparse_operation_none} \\ B^T, & \text{if trans_B == rocsparse_operation_transpose} \\ B^H, & \text{if trans_B == rocsparse_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 = csr_row_ptr[i]; k < csr_row_ptr[i + 1]; ++k)
        {
            C[i][j] += alpha * csr_val[k] * B[csr_col_ind[k]][j];
        }
    }
}

Example

This example multiplies a CSR matrix with a column-oriented dense matrix.

int main()
{
    //     1 2 0 3 0
    // A = 0 4 5 0 0
    //     6 0 0 7 8

    rocsparse_int m   = 3;
    rocsparse_int k   = 5;
    rocsparse_int nnz = 8;

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

    std::vector<int>   hcsr_row_ptr = {0, 3, 5, 8};
    std::vector<int>   hcsr_col_ind = {0, 1, 3, 1, 2, 0, 3, 4};
    std::vector<float> hcsr_val     = {1, 2, 3, 4, 5, 6, 7, 8};

    // Set dimension n of B
    rocsparse_int n = 64;

    // Allocate and generate column-oriented dense matrix B
    std::vector<float> hB(k * n);
    for(rocsparse_int i = 0; i < k * n; ++i)
    {
        hB[i] = static_cast<float>(rand()) / static_cast<float>(RAND_MAX);
    }
    std::vector<float> hC(m * n);

    int*   dcsr_row_ptr;
    int*   dcsr_col_ind;
    float* dcsr_val;
    float* dB;
    float* dC;
    HIP_CHECK(hipMalloc(&dcsr_row_ptr, sizeof(int) * (m + 1)));
    HIP_CHECK(hipMalloc(&dcsr_col_ind, sizeof(int) * nnz));
    HIP_CHECK(hipMalloc(&dcsr_val, sizeof(float) * nnz));
    HIP_CHECK(hipMalloc(&dB, sizeof(float) * k * n));
    HIP_CHECK(hipMalloc(&dC, sizeof(float) * m * n));

    HIP_CHECK(
        hipMemcpy(dcsr_row_ptr, hcsr_row_ptr.data(), sizeof(int) * (m + 1), hipMemcpyHostToDevice));
    HIP_CHECK(
        hipMemcpy(dcsr_col_ind, hcsr_col_ind.data(), sizeof(int) * nnz, hipMemcpyHostToDevice));
    HIP_CHECK(hipMemcpy(dcsr_val, hcsr_val.data(), sizeof(float) * nnz, hipMemcpyHostToDevice));
    HIP_CHECK(hipMemcpy(dB, hB.data(), sizeof(float) * k * n, hipMemcpyHostToDevice));
    HIP_CHECK(hipMemcpy(dC, hC.data(), sizeof(float) * m * n, hipMemcpyHostToDevice));

    // Create rocSPARSE handle
    rocsparse_handle handle;
    ROCSPARSE_CHECK(rocsparse_create_handle(&handle));

    // Create matrix descriptor
    rocsparse_mat_descr descr;
    ROCSPARSE_CHECK(rocsparse_create_mat_descr(&descr));

    // Perform the matrix multiplication
    ROCSPARSE_CHECK(rocsparse_scsrmm(handle,
                                     rocsparse_operation_none,
                                     rocsparse_operation_none,
                                     m,
                                     n,
                                     k,
                                     nnz,
                                     &alpha,
                                     descr,
                                     dcsr_val,
                                     dcsr_row_ptr,
                                     dcsr_col_ind,
                                     dB,
                                     k,
                                     &beta,
                                     dC,
                                     m));

    HIP_CHECK(hipMemcpy(hC.data(), dC, sizeof(float) * m * n, hipMemcpyDeviceToHost));

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

    // Clean up
    ROCSPARSE_CHECK(rocsparse_destroy_mat_descr(descr));
    ROCSPARSE_CHECK(rocsparse_destroy_handle(handle));

    // Clear up on device
    HIP_CHECK(hipFree(dcsr_row_ptr));
    HIP_CHECK(hipFree(dcsr_col_ind));
    HIP_CHECK(hipFree(dcsr_val));
    HIP_CHECK(hipFree(dB));
    HIP_CHECK(hipFree(dC));

    return 0;
}

Note

This function does not produce deterministic results when A is transposed.

Note

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

Note

This routine supports execution in a hipGraph context.

Parameters:

• handle – [in] handle to the rocsparse library context queue.

• trans_A – [in] matrix \(A\) operation type.

• trans_B – [in] matrix \(B\) operation type.

• m – [in] number of rows of the sparse CSR matrix \(A\).

• n – [in] number of columns of the column-oriented 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\).

• descr – [in] descriptor of the sparse CSR matrix \(A\). Currently, only rocsparse_matrix_type_general is supported.

• csr_val – [in] array of nnz elements of the sparse CSR matrix \(A\).

• csr_row_ptr – [in] array of m+1 elements that point to the start of every row of the sparse CSR matrix \(A\).

• csr_col_ind – [in] array of nnz elements containing the column indices of the sparse CSR matrix \(A\).

• B – [in] column-oriented dense matrix of dimension \(ldb \times n\) ( \(op(B) == B\)), \(ldb \times 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] column-oriented dense matrix of dimension \(ldc \times n\).

• ldc – [in] leading dimension of \(C\), must be at least \(\max{(1, m)}\) ( \(op(A) == A\)), \(\max{(1, k)}\) otherwise.

Return values:

• rocsparse_status_success – the operation completed successfully.

• rocsparse_status_invalid_handle – the library context was not initialized.

• rocsparse_status_invalid_size – m, n, k, nnz, ldb or ldc is invalid.

• rocsparse_status_invalid_pointer – descr, alpha, csr_val, csr_row_ptr, csr_col_ind, B, beta or C pointer is invalid.

• rocsparse_status_arch_mismatch – the device is not supported.

• rocsparse_status_not_implemented – rocsparse_matrix_type != rocsparse_matrix_type_general.

rocsparse_csrsm_zero_pivot()

rocsparse_status rocsparse_csrsm_zero_pivot(rocsparse_handle handle, rocsparse_mat_info info, rocsparse_int *position)

rocsparse_csrsm_zero_pivot returns rocsparse_status_zero_pivot, if either a structural or numerical zero has been found during rocsparse_Xcsrsm_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 rocsparse_status_success is returned instead.

Note

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

Note

This routine does not support execution in a hipGraph context.

Parameters:

• handle – [in] handle to the rocsparse 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:

• rocsparse_status_success – the operation completed successfully.

• rocsparse_status_invalid_handle – the library context was not initialized.

• rocsparse_status_invalid_pointer – info or position pointer is invalid.

• rocsparse_status_internal_error – an internal error occurred.

• rocsparse_status_zero_pivot – zero pivot has been found.

rocsparse_csrsm_buffer_size()

rocsparse_status rocsparse_scsrsm_buffer_size(rocsparse_handle handle, rocsparse_operation trans_A, rocsparse_operation trans_B, rocsparse_int m, rocsparse_int nrhs, rocsparse_int nnz, const float *alpha, const rocsparse_mat_descr descr, const float *csr_val, const rocsparse_int *csr_row_ptr, const rocsparse_int *csr_col_ind, const float *B, rocsparse_int ldb, rocsparse_mat_info info, rocsparse_solve_policy policy, size_t *buffer_size)

rocsparse_status rocsparse_dcsrsm_buffer_size(rocsparse_handle handle, rocsparse_operation trans_A, rocsparse_operation trans_B, rocsparse_int m, rocsparse_int nrhs, rocsparse_int nnz, const double *alpha, const rocsparse_mat_descr descr, const double *csr_val, const rocsparse_int *csr_row_ptr, const rocsparse_int *csr_col_ind, const double *B, rocsparse_int ldb, rocsparse_mat_info info, rocsparse_solve_policy policy, size_t *buffer_size)

rocsparse_status rocsparse_ccsrsm_buffer_size(rocsparse_handle handle, rocsparse_operation trans_A, rocsparse_operation trans_B, rocsparse_int m, rocsparse_int nrhs, rocsparse_int nnz, const rocsparse_float_complex *alpha, const rocsparse_mat_descr descr, const rocsparse_float_complex *csr_val, const rocsparse_int *csr_row_ptr, const rocsparse_int *csr_col_ind, const rocsparse_float_complex *B, rocsparse_int ldb, rocsparse_mat_info info, rocsparse_solve_policy policy, size_t *buffer_size)

rocsparse_status rocsparse_zcsrsm_buffer_size(rocsparse_handle handle, rocsparse_operation trans_A, rocsparse_operation trans_B, rocsparse_int m, rocsparse_int nrhs, rocsparse_int nnz, const rocsparse_double_complex *alpha, const rocsparse_mat_descr descr, const rocsparse_double_complex *csr_val, const rocsparse_int *csr_row_ptr, const rocsparse_int *csr_col_ind, const rocsparse_double_complex *B, rocsparse_int ldb, rocsparse_mat_info info, rocsparse_solve_policy policy, size_t *buffer_size)

rocsparse_csrsm_buffer_size returns the size of the temporary storage buffer that is required by rocsparse_Xcsrsm_analysis() and rocsparse_Xcsrsm_solve(). The temporary storage buffer must be allocated by the user.

Note

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

Note

This routine supports execution in a hipGraph context.

Parameters:

• handle – [in] handle to the rocsparse library context queue.

• trans_A – [in] matrix A operation type.

• trans_B – [in] matrix B operation type.

• m – [in] number of rows of the sparse CSR matrix A.

• nrhs – [in] number of columns of the column-oriented dense matrix op(B).

• nnz – [in] number of non-zero entries of the sparse CSR matrix A.

• alpha – [in] scalar \(\alpha\).

• descr – [in] descriptor of the sparse CSR matrix A.

• csr_val – [in] array of nnz elements of the sparse CSR matrix A.

• csr_row_ptr – [in] array of m+1 elements that point to the start of every row of the sparse CSR matrix A.

• csr_col_ind – [in] array of nnz elements containing the column indices of the sparse CSR matrix A.

• B – [in] column-oriented dense matrix of dimension 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] rocsparse_solve_policy_auto.

• buffer_size – [out] number of bytes of the temporary storage buffer required by rocsparse_scsrsm_analysis(), rocsparse_dcsrsm_analysis(), rocsparse_ccsrsm_analysis(), rocsparse_zcsrsm_analysis(), rocsparse_scsrsm_solve(), rocsparse_dcsrsm_solve(), rocsparse_ccsrsm_solve() and rocsparse_zcsrsm_solve().

Return values:

• rocsparse_status_success – the operation completed successfully.

• rocsparse_status_invalid_handle – the library context was not initialized.

• rocsparse_status_invalid_size – m, nrhs or nnz is invalid.

• rocsparse_status_invalid_pointer – alpha, descr, csr_val, csr_row_ptr, csr_col_ind, B, info or buffer_size pointer is invalid.

• rocsparse_status_internal_error – an internal error occurred.

• rocsparse_status_not_implemented – trans_A == rocsparse_operation_conjugate_transpose, trans_B == rocsparse_operation_conjugate_transpose or rocsparse_matrix_type != rocsparse_matrix_type_general.

rocsparse_csrsm_analysis()

rocsparse_status rocsparse_scsrsm_analysis(rocsparse_handle handle, rocsparse_operation trans_A, rocsparse_operation trans_B, rocsparse_int m, rocsparse_int nrhs, rocsparse_int nnz, const float *alpha, const rocsparse_mat_descr descr, const float *csr_val, const rocsparse_int *csr_row_ptr, const rocsparse_int *csr_col_ind, const float *B, rocsparse_int ldb, rocsparse_mat_info info, rocsparse_analysis_policy analysis, rocsparse_solve_policy solve, void *temp_buffer)

rocsparse_status rocsparse_dcsrsm_analysis(rocsparse_handle handle, rocsparse_operation trans_A, rocsparse_operation trans_B, rocsparse_int m, rocsparse_int nrhs, rocsparse_int nnz, const double *alpha, const rocsparse_mat_descr descr, const double *csr_val, const rocsparse_int *csr_row_ptr, const rocsparse_int *csr_col_ind, const double *B, rocsparse_int ldb, rocsparse_mat_info info, rocsparse_analysis_policy analysis, rocsparse_solve_policy solve, void *temp_buffer)

rocsparse_status rocsparse_ccsrsm_analysis(rocsparse_handle handle, rocsparse_operation trans_A, rocsparse_operation trans_B, rocsparse_int m, rocsparse_int nrhs, rocsparse_int nnz, const rocsparse_float_complex *alpha, const rocsparse_mat_descr descr, const rocsparse_float_complex *csr_val, const rocsparse_int *csr_row_ptr, const rocsparse_int *csr_col_ind, const rocsparse_float_complex *B, rocsparse_int ldb, rocsparse_mat_info info, rocsparse_analysis_policy analysis, rocsparse_solve_policy solve, void *temp_buffer)

rocsparse_status rocsparse_zcsrsm_analysis(rocsparse_handle handle, rocsparse_operation trans_A, rocsparse_operation trans_B, rocsparse_int m, rocsparse_int nrhs, rocsparse_int nnz, const rocsparse_double_complex *alpha, const rocsparse_mat_descr descr, const rocsparse_double_complex *csr_val, const rocsparse_int *csr_row_ptr, const rocsparse_int *csr_col_ind, const rocsparse_double_complex *B, rocsparse_int ldb, rocsparse_mat_info info, rocsparse_analysis_policy analysis, rocsparse_solve_policy solve, void *temp_buffer)

rocsparse_csrsm_analysis performs the analysis step for rocsparse_Xcsrsm_solve(). It is expected that this function will be executed only once for a given matrix and particular operation type. The analysis meta data can be cleared by rocsparse_csrsm_clear().

rocsparse_csrsm_analysis can share its meta data with rocsparse_Xcsrilu0_analysis(), rocsparse_Xcsric0_analysis(), and rocsparse_Xcsrsv_analysis(). Selecting rocsparse_analysis_policy_reuse policy can greatly improve computation performance of meta data. However, the user needs to make sure that the sparsity pattern remains unchanged. If this cannot be assured, rocsparse_analysis_policy_force has to be used.

Note

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

Note

This function is blocking with respect to the host.

Note

This routine does not support execution in a hipGraph context.

Parameters:

• handle – [in] handle to the rocsparse library context queue.

• trans_A – [in] matrix A operation type.

• trans_B – [in] matrix B operation type.

• m – [in] number of rows of the sparse CSR matrix A.

• nrhs – [in] number of columns of the column-oriented dense matrix op(B).

• nnz – [in] number of non-zero entries of the sparse CSR matrix A.

• alpha – [in] scalar \(\alpha\).

• descr – [in] descriptor of the sparse CSR matrix A.

• csr_val – [in] array of nnz elements of the sparse CSR matrix A.

• csr_row_ptr – [in] array of m+1 elements that point to the start of every row of the sparse CSR matrix A.

• csr_col_ind – [in] array of nnz elements containing the column indices of the sparse CSR matrix A.

• B – [in] column-oriented dense matrix of dimension 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.

• analysis – [in] rocsparse_analysis_policy_reuse or rocsparse_analysis_policy_force.

• solve – [in] rocsparse_solve_policy_auto.

• temp_buffer – [in] temporary storage buffer allocated by the user.

Return values:

• rocsparse_status_success – the operation completed successfully.

• rocsparse_status_invalid_handle – the library context was not initialized.

• rocsparse_status_invalid_size – m, nrhs or nnz is invalid.

• rocsparse_status_invalid_pointer – alpha, descr, csr_val, csr_row_ptr, csr_col_ind, B, info or temp_buffer pointer is invalid.

• rocsparse_status_internal_error – an internal error occurred.

• rocsparse_status_not_implemented – trans_A == rocsparse_operation_conjugate_transpose, trans_B == rocsparse_operation_conjugate_transpose or rocsparse_matrix_type != rocsparse_matrix_type_general.

rocsparse_csrsm_solve()

rocsparse_status rocsparse_scsrsm_solve(rocsparse_handle handle, rocsparse_operation trans_A, rocsparse_operation trans_B, rocsparse_int m, rocsparse_int nrhs, rocsparse_int nnz, const float *alpha, const rocsparse_mat_descr descr, const float *csr_val, const rocsparse_int *csr_row_ptr, const rocsparse_int *csr_col_ind, float *B, rocsparse_int ldb, rocsparse_mat_info info, rocsparse_solve_policy policy, void *temp_buffer)

rocsparse_status rocsparse_dcsrsm_solve(rocsparse_handle handle, rocsparse_operation trans_A, rocsparse_operation trans_B, rocsparse_int m, rocsparse_int nrhs, rocsparse_int nnz, const double *alpha, const rocsparse_mat_descr descr, const double *csr_val, const rocsparse_int *csr_row_ptr, const rocsparse_int *csr_col_ind, double *B, rocsparse_int ldb, rocsparse_mat_info info, rocsparse_solve_policy policy, void *temp_buffer)

rocsparse_status rocsparse_ccsrsm_solve(rocsparse_handle handle, rocsparse_operation trans_A, rocsparse_operation trans_B, rocsparse_int m, rocsparse_int nrhs, rocsparse_int nnz, const rocsparse_float_complex *alpha, const rocsparse_mat_descr descr, const rocsparse_float_complex *csr_val, const rocsparse_int *csr_row_ptr, const rocsparse_int *csr_col_ind, rocsparse_float_complex *B, rocsparse_int ldb, rocsparse_mat_info info, rocsparse_solve_policy policy, void *temp_buffer)

rocsparse_status rocsparse_zcsrsm_solve(rocsparse_handle handle, rocsparse_operation trans_A, rocsparse_operation trans_B, rocsparse_int m, rocsparse_int nrhs, rocsparse_int nnz, const rocsparse_double_complex *alpha, const rocsparse_mat_descr descr, const rocsparse_double_complex *csr_val, const rocsparse_int *csr_row_ptr, const rocsparse_int *csr_col_ind, rocsparse_double_complex *B, rocsparse_int ldb, rocsparse_mat_info info, rocsparse_solve_policy policy, void *temp_buffer)

Sparse triangular system solve using CSR storage format.

rocsparse_csrsm_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 trans_A == rocsparse_operation_none} \\ A^T, & \text{if trans_A == rocsparse_operation_transpose} \\ A^H, & \text{if trans_A == rocsparse_operation_conjugate_transpose} \end{array} \right. \end{split}\]

,

\[\begin{split} op(B) = \left\{ \begin{array}{ll} B, & \text{if trans_B == rocsparse_operation_none} \\ B^T, & \text{if trans_B == rocsparse_operation_transpose} \\ B^H, & \text{if trans_B == rocsparse_operation_conjugate_transpose} \end{array} \right. \end{split}\]

and

\[\begin{split} op(X) = \left\{ \begin{array}{ll} X, & \text{if trans_B == rocsparse_operation_none} \\ X^T, & \text{if trans_B == rocsparse_operation_transpose} \\ X^H, & \text{if trans_B == rocsparse_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 rocsparse_csrsm_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 trans_B:

\[\begin{split} op(B)/op(X) = \left\{ \begin{array}{ll} ldb \times nrhs, \text{ } ldb \ge m, & \text{if trans_B == rocsparse_operation_none} \\ ldb \times m, \text{ } ldb \ge nrhs, & \text{if trans_B == rocsparse_operation_transpose} \\ ldb \times m, \text{ } ldb \ge nrhs, & \text{if trans_B == rocsparse_operation_conjugate_transpose} \end{array} \right. \end{split}\]

rocsparse_csrsm_solve requires a user allocated temporary buffer. Its size is returned by rocsparse_Xcsrsm_buffer_size(). The size of the required buffer is larger when trans_A equals rocsparse_operation_transpose or rocsparse_operation_conjugate_transpose and when trans_B is rocsparse_operation_none. 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 rocsparse_Xcsrsm_analysis().

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. rocsparse_csrsm_solve tracks the location of the first zero pivot (either numerical or structural zero). The zero pivot status can be checked calling rocsparse_csrsm_zero_pivot(). If rocsparse_csrsm_zero_pivot() returns rocsparse_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 descr to rocsparse_diag_type_unit using rocsparse_set_mat_diag_type. If rocsparse_diag_type == rocsparse_diag_type_unit, no zero pivot will be reported, even if \(A_{j,j} = 0\) for some \(j\).

The sparse CSR matrix passed to rocsparse_csrsm_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 rocsparse_fill_mode set on the descriptor descr. If the fill mode is set to rocsparse_fill_mode_lower, then the lower triangular matrix is solved. If the fill mode is set to rocsparse_fill_mode_upper then the upper triangular matrix is solved.

Example

Consider the lower triangular \(m \times m\) matrix \(L\), stored in CSR storage format with unit diagonal. The following example solves \(L \cdot X = B\).

int main()
{
    //     1 2 0 3 0
    // A = 0 4 5 0 0
    //     6 0 0 7 8

    rocsparse_int m   = 3;
    rocsparse_int k   = 5;
    rocsparse_int nnz = 8;

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

    std::vector<int>   hcsr_row_ptr = {0, 3, 5, 8};
    std::vector<int>   hcsr_col_ind = {0, 1, 3, 1, 2, 0, 3, 4};
    std::vector<float> hcsr_val     = {1, 2, 3, 4, 5, 6, 7, 8};

    // Set dimension n of B
    rocsparse_int n = 64;

    // Allocate and generate column-oriented dense matrix B
    std::vector<float> hB(k * n);
    for(rocsparse_int i = 0; i < k * n; ++i)
    {
        hB[i] = static_cast<float>(rand()) / static_cast<float>(RAND_MAX);
    }
    std::vector<float> hC(m * n);

    int*   dcsr_row_ptr;
    int*   dcsr_col_ind;
    float* dcsr_val;
    float* dB;
    float* dC;
    HIP_CHECK(hipMalloc(&dcsr_row_ptr, sizeof(int) * (m + 1)));
    HIP_CHECK(hipMalloc(&dcsr_col_ind, sizeof(int) * nnz));
    HIP_CHECK(hipMalloc(&dcsr_val, sizeof(float) * nnz));
    HIP_CHECK(hipMalloc(&dB, sizeof(float) * k * n));
    HIP_CHECK(hipMalloc(&dC, sizeof(float) * m * n));

    HIP_CHECK(
        hipMemcpy(dcsr_row_ptr, hcsr_row_ptr.data(), sizeof(int) * (m + 1), hipMemcpyHostToDevice));
    HIP_CHECK(
        hipMemcpy(dcsr_col_ind, hcsr_col_ind.data(), sizeof(int) * nnz, hipMemcpyHostToDevice));
    HIP_CHECK(hipMemcpy(dcsr_val, hcsr_val.data(), sizeof(float) * nnz, hipMemcpyHostToDevice));
    HIP_CHECK(hipMemcpy(dB, hB.data(), sizeof(float) * k * n, hipMemcpyHostToDevice));
    HIP_CHECK(hipMemcpy(dC, hC.data(), sizeof(float) * m * n, hipMemcpyHostToDevice));

    // Create rocSPARSE handle
    rocsparse_handle handle;
    ROCSPARSE_CHECK(rocsparse_create_handle(&handle));

    // Create matrix descriptor
    rocsparse_mat_descr descr;
    ROCSPARSE_CHECK(rocsparse_create_mat_descr(&descr));

    // Perform the matrix multiplication
    ROCSPARSE_CHECK(rocsparse_scsrmm(handle,
                                     rocsparse_operation_none,
                                     rocsparse_operation_none,
                                     m,
                                     n,
                                     k,
                                     nnz,
                                     &alpha,
                                     descr,
                                     dcsr_val,
                                     dcsr_row_ptr,
                                     dcsr_col_ind,
                                     dB,
                                     k,
                                     &beta,
                                     dC,
                                     m));

    HIP_CHECK(hipMemcpy(hC.data(), dC, sizeof(float) * m * n, hipMemcpyDeviceToHost));

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

    // Clean up
    ROCSPARSE_CHECK(rocsparse_destroy_mat_descr(descr));
    ROCSPARSE_CHECK(rocsparse_destroy_handle(handle));

    // Clear up on device
    HIP_CHECK(hipFree(dcsr_row_ptr));
    HIP_CHECK(hipFree(dcsr_col_ind));
    HIP_CHECK(hipFree(dcsr_val));
    HIP_CHECK(hipFree(dB));
    HIP_CHECK(hipFree(dC));

    return 0;
}

Note

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

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 trans_A != rocsparse_operation_conjugate_transpose and trans_B != rocsparse_operation_conjugate_transpose is supported.

Note

This routine supports execution in a hipGraph context.

Parameters:

• handle – [in] handle to the rocsparse library context queue.

• trans_A – [in] matrix A operation type.

• trans_B – [in] matrix B operation type.

• m – [in] number of rows of the sparse CSR matrix A.

• nrhs – [in] number of columns of the column-oriented dense matrix op(B).

• nnz – [in] number of non-zero entries of the sparse CSR matrix A.

• alpha – [in] scalar \(\alpha\).

• descr – [in] descriptor of the sparse CSR matrix A.

• csr_val – [in] array of nnz elements of the sparse CSR matrix A.

• csr_row_ptr – [in] array of m+1 elements that point to the start of every row of the sparse CSR matrix A.

• csr_col_ind – [in] array of nnz elements containing the column indices of the sparse CSR matrix A.

• B – [inout] column-oriented dense matrix of dimension 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] rocsparse_solve_policy_auto.

• temp_buffer – [in] temporary storage buffer allocated by the user.

Return values:

• rocsparse_status_success – the operation completed successfully.

• rocsparse_status_invalid_handle – the library context was not initialized.

• rocsparse_status_invalid_size – m, nrhs or nnz is invalid.

• rocsparse_status_invalid_pointer – alpha, descr, csr_val, csr_row_ptr, csr_col_ind, B, info or temp_buffer pointer is invalid.

• rocsparse_status_internal_error – an internal error occurred.

• rocsparse_status_not_implemented – trans_A == rocsparse_operation_conjugate_transpose, trans_B == rocsparse_operation_conjugate_transpose or rocsparse_matrix_type != rocsparse_matrix_type_general.

rocsparse_csrsm_clear()

rocsparse_status rocsparse_csrsm_clear(rocsparse_handle handle, rocsparse_mat_info info)

rocsparse_csrsm_clear deallocates all memory that was allocated by rocsparse_Xcsrsm_analysis(). This is especially useful, if memory is an issue and the analysis data is not required for further computation, e.g. when switching to another sparse matrix format. Calling rocsparse_csrsm_clear is optional. All allocated resources will be cleared when the opaque rocsparse_mat_info struct is destroyed using rocsparse_destroy_mat_info().

Note

This routine does not support execution in a hipGraph context.

Parameters:

• handle – [in] handle to the rocsparse library context queue.

• info – [inout] structure that holds the information collected during the analysis step.

Return values:

• rocsparse_status_success – the operation completed successfully.

• rocsparse_status_invalid_handle – the library context was not initialized.

• rocsparse_status_invalid_pointer – info pointer is invalid.

• rocsparse_status_memory_error – the buffer holding the meta data could not be deallocated.

• rocsparse_status_internal_error – an internal error occurred.

rocsparse_bsrsm_zero_pivot()

rocsparse_status rocsparse_bsrsm_zero_pivot(rocsparse_handle handle, rocsparse_mat_info info, rocsparse_int *position)

rocsparse_bsrsm_zero_pivot returns rocsparse_status_zero_pivot, if either a structural or numerical zero has been found during rocsparse_Xbsrsm_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 rocsparse_status_success is returned instead.

Note

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

Note

This routine does not support execution in a hipGraph context.

Parameters:

• handle – [in] handle to the rocsparse 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:

• rocsparse_status_success – the operation completed successfully.

• rocsparse_status_invalid_handle – the library context was not initialized.

• rocsparse_status_invalid_pointer – info or position pointer is invalid.

• rocsparse_status_internal_error – an internal error occurred.

• rocsparse_status_zero_pivot – zero pivot has been found.

rocsparse_bsrsm_buffer_size()

rocsparse_status rocsparse_sbsrsm_buffer_size(rocsparse_handle handle, rocsparse_direction dir, rocsparse_operation trans_A, rocsparse_operation trans_X, rocsparse_int mb, rocsparse_int nrhs, rocsparse_int nnzb, const rocsparse_mat_descr descr, const float *bsr_val, const rocsparse_int *bsr_row_ptr, const rocsparse_int *bsr_col_ind, rocsparse_int block_dim, rocsparse_mat_info info, size_t *buffer_size)

rocsparse_status rocsparse_dbsrsm_buffer_size(rocsparse_handle handle, rocsparse_direction dir, rocsparse_operation trans_A, rocsparse_operation trans_X, rocsparse_int mb, rocsparse_int nrhs, rocsparse_int nnzb, const rocsparse_mat_descr descr, const double *bsr_val, const rocsparse_int *bsr_row_ptr, const rocsparse_int *bsr_col_ind, rocsparse_int block_dim, rocsparse_mat_info info, size_t *buffer_size)

rocsparse_status rocsparse_cbsrsm_buffer_size(rocsparse_handle handle, rocsparse_direction dir, rocsparse_operation trans_A, rocsparse_operation trans_X, rocsparse_int mb, rocsparse_int nrhs, rocsparse_int nnzb, const rocsparse_mat_descr descr, const rocsparse_float_complex *bsr_val, const rocsparse_int *bsr_row_ptr, const rocsparse_int *bsr_col_ind, rocsparse_int block_dim, rocsparse_mat_info info, size_t *buffer_size)

rocsparse_status rocsparse_zbsrsm_buffer_size(rocsparse_handle handle, rocsparse_direction dir, rocsparse_operation trans_A, rocsparse_operation trans_X, rocsparse_int mb, rocsparse_int nrhs, rocsparse_int nnzb, const rocsparse_mat_descr descr, const rocsparse_double_complex *bsr_val, const rocsparse_int *bsr_row_ptr, const rocsparse_int *bsr_col_ind, rocsparse_int block_dim, rocsparse_mat_info info, size_t *buffer_size)

rocsparse_bsrsm_buffer_size returns the size of the temporary storage buffer that is required by rocsparse_Xbsrsm_analysis() and rocsparse_Xbsrsm_solve(). The temporary storage buffer must be allocated by the user.

Note

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

Note

This routine supports execution in a hipGraph context.

Parameters:

• handle – [in] handle to the rocsparse library context queue.

• dir – [in] matrix storage of BSR blocks.

• trans_A – [in] matrix A operation type.

• trans_X – [in] matrix X operation type.

• mb – [in] number of block rows of the sparse BSR matrix A.

• nrhs – [in] number of columns of the column-oriented dense matrix op(X).

• nnzb – [in] number of non-zero blocks of the sparse BSR matrix A.

• descr – [in] descriptor of the sparse BSR matrix A.

• bsr_val – [in] array of nnzb blocks of the sparse BSR matrix.

• bsr_row_ptr – [in] array of mb+1 elements that point to the start of every block row of the sparse BSR matrix.

• bsr_col_ind – [in] array of nnzb containing the block column indices of the sparse BSR matrix.

• block_dim – [in] block dimension of the sparse BSR matrix.

• info – [in] structure that holds the information collected during the analysis step.

• buffer_size – [out] number of bytes of the temporary storage buffer required by rocsparse_Xbsrsm_analysis() and rocsparse_Xbsrsm_solve().

Return values:

• rocsparse_status_success – the operation completed successfully.

• rocsparse_status_invalid_handle – the library context was not initialized.

• rocsparse_status_invalid_size – mb, nrhs, nnzb or block_dim is invalid.

• rocsparse_status_invalid_pointer – descr, bsr_val, bsr_row_ptr, bsr_col_ind, info or buffer_size pointer is invalid.

• rocsparse_status_internal_error – an internal error occurred.

• rocsparse_status_not_implemented – trans_A == rocsparse_operation_conjugate_transpose, trans_X == rocsparse_operation_conjugate_transpose or rocsparse_matrix_type != rocsparse_matrix_type_general.

rocsparse_bsrsm_analysis()

rocsparse_status rocsparse_sbsrsm_analysis(rocsparse_handle handle, rocsparse_direction dir, rocsparse_operation trans_A, rocsparse_operation trans_X, rocsparse_int mb, rocsparse_int nrhs, rocsparse_int nnzb, const rocsparse_mat_descr descr, const float *bsr_val, const rocsparse_int *bsr_row_ptr, const rocsparse_int *bsr_col_ind, rocsparse_int block_dim, rocsparse_mat_info info, rocsparse_analysis_policy analysis, rocsparse_solve_policy solve, void *temp_buffer)

rocsparse_status rocsparse_dbsrsm_analysis(rocsparse_handle handle, rocsparse_direction dir, rocsparse_operation trans_A, rocsparse_operation trans_X, rocsparse_int mb, rocsparse_int nrhs, rocsparse_int nnzb, const rocsparse_mat_descr descr, const double *bsr_val, const rocsparse_int *bsr_row_ptr, const rocsparse_int *bsr_col_ind, rocsparse_int block_dim, rocsparse_mat_info info, rocsparse_analysis_policy analysis, rocsparse_solve_policy solve, void *temp_buffer)

rocsparse_status rocsparse_cbsrsm_analysis(rocsparse_handle handle, rocsparse_direction dir, rocsparse_operation trans_A, rocsparse_operation trans_X, rocsparse_int mb, rocsparse_int nrhs, rocsparse_int nnzb, const rocsparse_mat_descr descr, const rocsparse_float_complex *bsr_val, const rocsparse_int *bsr_row_ptr, const rocsparse_int *bsr_col_ind, rocsparse_int block_dim, rocsparse_mat_info info, rocsparse_analysis_policy analysis, rocsparse_solve_policy solve, void *temp_buffer)

rocsparse_status rocsparse_zbsrsm_analysis(rocsparse_handle handle, rocsparse_direction dir, rocsparse_operation trans_A, rocsparse_operation trans_X, rocsparse_int mb, rocsparse_int nrhs, rocsparse_int nnzb, const rocsparse_mat_descr descr, const rocsparse_double_complex *bsr_val, const rocsparse_int *bsr_row_ptr, const rocsparse_int *bsr_col_ind, rocsparse_int block_dim, rocsparse_mat_info info, rocsparse_analysis_policy analysis, rocsparse_solve_policy solve, void *temp_buffer)

rocsparse_bsrsm_analysis performs the analysis step for rocsparse_Xbsrsm_solve(). It is expected that this function will be executed only once for a given matrix and particular operation type. The analysis meta data can be cleared by rocsparse_bsrsm_clear().

rocsparse_bsrsm_analysis can share its meta data with rocsparse_Xbsrilu0_analysis(), rocsparse_Xbsric0_analysis(), rocsparse_Xbsrsv_analysis(). Selecting rocsparse_analysis_policy_reuse policy can greatly improve computation performance of meta data. However, the user needs to make sure that the sparsity pattern remains unchanged. If this cannot be assured, rocsparse_analysis_policy_force has to be used.

Note

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

Note

This function is blocking with respect to the host.

Note

This routine does not support execution in a hipGraph context.

Parameters:

• handle – [in] handle to the rocsparse library context queue.

• dir – [in] matrix storage of BSR blocks.

• trans_A – [in] matrix A operation type.

• trans_X – [in] matrix X operation type.

• mb – [in] number of block rows of the sparse BSR matrix A.

• nrhs – [in] number of columns of the column-oriented dense matrix op(X).

• nnzb – [in] number of non-zero blocks of the sparse BSR matrix A.

• descr – [in] descriptor of the sparse BSR matrix A.

• bsr_val – [in] array of nnzb blocks of the sparse BSR matrix A.

• bsr_row_ptr – [in] array of mb+1 elements that point to the start of every block row of the sparse BSR matrix A.

• bsr_col_ind – [in] array of nnzb containing the block column indices of the sparse BSR matrix A.

• block_dim – [in] block dimension of the sparse BSR matrix A.

• info – [out] structure that holds the information collected during the analysis step.

• analysis – [in] rocsparse_analysis_policy_reuse or rocsparse_analysis_policy_force.

• solve – [in] rocsparse_solve_policy_auto.

• temp_buffer – [in] temporary storage buffer allocated by the user.

Return values:

• rocsparse_status_success – the operation completed successfully.

• rocsparse_status_invalid_handle – the library context was not initialized.

• rocsparse_status_invalid_size – mb, nrhs, nnzb or block_dim is invalid.

• rocsparse_status_invalid_pointer – descr, bsr_val, bsr_row_ptr, bsr_col_ind, info or temp_buffer pointer is invalid.

• rocsparse_status_internal_error – an internal error occurred.

• rocsparse_status_not_implemented – trans_A == rocsparse_operation_conjugate_transpose, trans_X == rocsparse_operation_conjugate_transpose or rocsparse_matrix_type != rocsparse_matrix_type_general.

rocsparse_bsrsm_solve()

rocsparse_status rocsparse_sbsrsm_solve(rocsparse_handle handle, rocsparse_direction dir, rocsparse_operation trans_A, rocsparse_operation trans_X, rocsparse_int mb, rocsparse_int nrhs, rocsparse_int nnzb, const float *alpha, const rocsparse_mat_descr descr, const float *bsr_val, const rocsparse_int *bsr_row_ptr, const rocsparse_int *bsr_col_ind, rocsparse_int block_dim, rocsparse_mat_info info, const float *B, rocsparse_int ldb, float *X, rocsparse_int ldx, rocsparse_solve_policy policy, void *temp_buffer)

rocsparse_status rocsparse_dbsrsm_solve(rocsparse_handle handle, rocsparse_direction dir, rocsparse_operation trans_A, rocsparse_operation trans_X, rocsparse_int mb, rocsparse_int nrhs, rocsparse_int nnzb, const double *alpha, const rocsparse_mat_descr descr, const double *bsr_val, const rocsparse_int *bsr_row_ptr, const rocsparse_int *bsr_col_ind, rocsparse_int block_dim, rocsparse_mat_info info, const double *B, rocsparse_int ldb, double *X, rocsparse_int ldx, rocsparse_solve_policy policy, void *temp_buffer)

rocsparse_status rocsparse_cbsrsm_solve(rocsparse_handle handle, rocsparse_direction dir, rocsparse_operation trans_A, rocsparse_operation trans_X, rocsparse_int mb, rocsparse_int nrhs, rocsparse_int nnzb, const rocsparse_float_complex *alpha, const rocsparse_mat_descr descr, const rocsparse_float_complex *bsr_val, const rocsparse_int *bsr_row_ptr, const rocsparse_int *bsr_col_ind, rocsparse_int block_dim, rocsparse_mat_info info, const rocsparse_float_complex *B, rocsparse_int ldb, rocsparse_float_complex *X, rocsparse_int ldx, rocsparse_solve_policy policy, void *temp_buffer)

rocsparse_status rocsparse_zbsrsm_solve(rocsparse_handle handle, rocsparse_direction dir, rocsparse_operation trans_A, rocsparse_operation trans_X, rocsparse_int mb, rocsparse_int nrhs, rocsparse_int nnzb, const rocsparse_double_complex *alpha, const rocsparse_mat_descr descr, const rocsparse_double_complex *bsr_val, const rocsparse_int *bsr_row_ptr, const rocsparse_int *bsr_col_ind, rocsparse_int block_dim, rocsparse_mat_info info, const rocsparse_double_complex *B, rocsparse_int ldb, rocsparse_double_complex *X, rocsparse_int ldx, rocsparse_solve_policy policy, void *temp_buffer)

Sparse triangular system solve using BSR storage format.

rocsparse_bsrsm_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 trans_A == rocsparse_operation_none} \\ A^T, & \text{if trans_A == rocsparse_operation_transpose} \\ A^H, & \text{if trans_A == rocsparse_operation_conjugate_transpose} \end{array} \right. \end{split}\]

,

\[\begin{split} op(B) = \left\{ \begin{array}{ll} B, & \text{if trans_X == rocsparse_operation_none} \\ B^T, & \text{if trans_X == rocsparse_operation_transpose} \\ B^H, & \text{if trans_X == rocsparse_operation_conjugate_transpose} \end{array} \right. \end{split}\]

and

\[\begin{split} op(X) = \left\{ \begin{array}{ll} X, & \text{if trans_X == rocsparse_operation_none} \\ X^T, & \text{if trans_X == rocsparse_operation_transpose} \\ X^H, & \text{if trans_X == rocsparse_operation_conjugate_transpose} \end{array} \right. \end{split}\]

and where \(m = block\_dim \times mb\).

Note that as indicated above, the operation type of both \(op(B)\) and \(op(X)\) is specified by the trans_X 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 trans_X:

\[\begin{split} op(B) = \left\{ \begin{array}{ll} ldb \times nrhs, \text{ } ldb \ge m, & \text{if trans_X == rocsparse_operation_none} \\ ldb \times m, \text{ } ldb \ge nrhs, & \text{if trans_X == rocsparse_operation_transpose} \\ ldb \times m, \text{ } ldb \ge nrhs, & \text{if trans_X == rocsparse_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 trans_X == rocsparse_operation_none} \\ ldb \times m, \text{ } ldb \ge nrhs, & \text{if trans_X == rocsparse_operation_transpose} \\ ldb \times m, \text{ } ldb \ge nrhs, & \text{if trans_X == rocsparse_operation_conjugate_transpose} \end{array} \right. \end{split}\]

rocsparse_bsrsm_solve requires a user allocated temporary buffer. Its size is returned by rocsparse_Xbsrsm_buffer_size(). The size of the required buffer is larger when trans_A equals rocsparse_operation_transpose or rocsparse_operation_conjugate_transpose and when trans_X is rocsparse_operation_none. 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 rocsparse_sbsrsm_analysis “rocsparse_Xbsrsm_analysis()”.

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. rocsparse_bsrsm_solve tracks the location of the first zero pivot (either numerical or structural zero). The zero pivot status can be checked calling rocsparse_bsrsm_zero_pivot(). If rocsparse_bsrsm_zero_pivot() returns rocsparse_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 descr to rocsparse_diag_type_unit using rocsparse_set_mat_diag_type. If rocsparse_diag_type == rocsparse_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 rocsparse_bsrsm_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 rocsparse_fill_mode set on the descriptor descr. If the fill mode is set to rocsparse_fill_mode_lower, then the lower triangular matrix is solved. If the fill mode is set to rocsparse_fill_mode_upper then the upper triangular matrix is solved.

Example

Consider the lower triangular \(m \times m\) matrix \(L\), stored in BSR storage format with non-unit diagonal. The following example solves \(L \cdot X = B\).

int main()
{
    // rocSPARSE handle
    rocsparse_handle handle;
    ROCSPARSE_CHECK(rocsparse_create_handle(&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
    rocsparse_int m = 4;

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

    // BSR block dimension
    rocsparse_int bsr_dim = 2;

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

    // Number of non-zero blocks
    rocsparse_int nnzb = 3;

    // BSR row pointers
    rocsparse_int hbsr_row_ptr[3] = {0, 1, 3};

    // BSR column indices
    rocsparse_int hbsr_col_ind[3] = {0, 0, 1};

    // BSR values
    double hbsr_val[12] = {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
    rocsparse_direction dir = rocsparse_direction_column;

    // Transposition of the matrix and rhs matrix
    rocsparse_operation transA = rocsparse_operation_none;
    rocsparse_operation transX = rocsparse_operation_none;

    // Analysis policy
    rocsparse_analysis_policy analysis_policy = rocsparse_analysis_policy_reuse;

    // Solve policy
    rocsparse_solve_policy solve_policy = rocsparse_solve_policy_auto;

    // Scalar alpha and beta
    double alpha = 3.7;

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

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

    // Offload data to device
    rocsparse_int* dbsr_row_ptr;
    rocsparse_int* dbsr_col_ind;
    double*        dbsr_val;
    double*        dB;
    double*        dX;

    HIP_CHECK(hipMalloc(&dbsr_row_ptr, sizeof(rocsparse_int) * (mb + 1)));
    HIP_CHECK(hipMalloc(&dbsr_col_ind, sizeof(rocsparse_int) * nnzb));
    HIP_CHECK(hipMalloc(&dbsr_val, sizeof(double) * nnzb * bsr_dim * bsr_dim));
    HIP_CHECK(hipMalloc(&dB, sizeof(double) * nb * bsr_dim * nrhs));
    HIP_CHECK(hipMalloc(&dX, sizeof(double) * mb * bsr_dim * nrhs));

    HIP_CHECK(hipMemcpy(
        dbsr_row_ptr, hbsr_row_ptr, sizeof(rocsparse_int) * (mb + 1), hipMemcpyHostToDevice));
    HIP_CHECK(
        hipMemcpy(dbsr_col_ind, hbsr_col_ind, sizeof(rocsparse_int) * nnzb, hipMemcpyHostToDevice));
    HIP_CHECK(hipMemcpy(
        dbsr_val, hbsr_val, sizeof(double) * nnzb * bsr_dim * bsr_dim, hipMemcpyHostToDevice));
    HIP_CHECK(hipMemcpy(dB, hB, sizeof(double) * nb * bsr_dim * nrhs, hipMemcpyHostToDevice));

    // Matrix descriptor
    rocsparse_mat_descr descr;
    ROCSPARSE_CHECK(rocsparse_create_mat_descr(&descr));

    // Matrix fill mode
    ROCSPARSE_CHECK(rocsparse_set_mat_fill_mode(descr, rocsparse_fill_mode_lower));

    // Matrix diagonal type
    ROCSPARSE_CHECK(rocsparse_set_mat_diag_type(descr, rocsparse_diag_type_non_unit));

    // Matrix info structure
    rocsparse_mat_info info;
    ROCSPARSE_CHECK(rocsparse_create_mat_info(&info));

    // Obtain required buffer size
    size_t buffer_size;
    ROCSPARSE_CHECK(rocsparse_dbsrsm_buffer_size(handle,
                                                 dir,
                                                 transA,
                                                 transX,
                                                 mb,
                                                 nrhs,
                                                 nnzb,
                                                 descr,
                                                 dbsr_val,
                                                 dbsr_row_ptr,
                                                 dbsr_col_ind,
                                                 bsr_dim,
                                                 info,
                                                 &buffer_size));

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

    // Perform analysis step
    ROCSPARSE_CHECK(rocsparse_dbsrsm_analysis(handle,
                                              dir,
                                              transA,
                                              transX,
                                              mb,
                                              nrhs,
                                              nnzb,
                                              descr,
                                              dbsr_val,
                                              dbsr_row_ptr,
                                              dbsr_col_ind,
                                              bsr_dim,
                                              info,
                                              analysis_policy,
                                              solve_policy,
                                              temp_buffer));

    // Call dbsrsm to perform lower triangular solve LX = B
    ROCSPARSE_CHECK(rocsparse_dbsrsm_solve(handle,
                                           dir,
                                           transA,
                                           transX,
                                           mb,
                                           nrhs,
                                           nnzb,
                                           &alpha,
                                           descr,
                                           dbsr_val,
                                           dbsr_row_ptr,
                                           dbsr_col_ind,
                                           bsr_dim,
                                           info,
                                           dB,
                                           ldb,
                                           dX,
                                           ldx,
                                           solve_policy,
                                           temp_buffer));

    // Check for zero pivots
    rocsparse_int    pivot;
    rocsparse_status status = rocsparse_bsrsm_zero_pivot(handle, info, &pivot);

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

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

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

    // Clear rocSPARSE
    ROCSPARSE_CHECK(rocsparse_destroy_mat_info(info));
    ROCSPARSE_CHECK(rocsparse_destroy_mat_descr(descr));
    ROCSPARSE_CHECK(rocsparse_destroy_handle(handle));

    // Clear device memory
    HIP_CHECK(hipFree(dbsr_row_ptr));
    HIP_CHECK(hipFree(dbsr_col_ind));
    HIP_CHECK(hipFree(dbsr_val));
    HIP_CHECK(hipFree(dB));
    HIP_CHECK(hipFree(dX));
    HIP_CHECK(hipFree(temp_buffer));

    return 0;
}

Note

The sparse BSR matrix has to be sorted.

Note

Operation type of B and X must match, for example if \(op(B)=B\) then \(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 trans_A != rocsparse_operation_conjugate_transpose and trans_X != rocsparse_operation_conjugate_transpose is supported.

Note

This routine supports execution in a hipGraph context.

Parameters:

• handle – [in] handle to the rocsparse library context queue.

• dir – [in] matrix storage of BSR blocks.

• trans_A – [in] matrix A operation type.

• trans_X – [in] matrix X operation type.

• mb – [in] number of block rows of the sparse BSR matrix A.

• nrhs – [in] number of columns of the column-oriented dense matrix op(X).

• nnzb – [in] number of non-zero blocks of the sparse BSR matrix A.

• alpha – [in] scalar \(\alpha\).

• descr – [in] descriptor of the sparse BSR matrix A.

• bsr_val – [in] array of nnzb blocks of the sparse BSR matrix.

• bsr_row_ptr – [in] array of mb+1 elements that point to the start of every block row of the sparse BSR matrix.

• bsr_col_ind – [in] array of nnzb containing the block column indices of the sparse BSR matrix.

• block_dim – [in] block dimension of the sparse BSR matrix.

• info – [in] structure that holds the information collected during the analysis step.

• B – [in] column-oriented dense matrix B with leading dimension ldb.

• ldb – [in] leading dimension of rhs matrix B.

• X – [out] column-oriented dense solution matrix X with leading dimension ldx.

• ldx – [in] leading dimension of solution matrix X.

• policy – [in] rocsparse_solve_policy_auto.

• temp_buffer – [in] temporary storage buffer allocated by the user.

Return values:

• rocsparse_status_success – the operation completed successfully.

• rocsparse_status_invalid_handle – the library context was not initialized.

• rocsparse_status_invalid_size – mb, nrhs, nnzb or block_dim is invalid.

• rocsparse_status_invalid_pointer – alpha, descr, bsr_val, bsr_row_ptr, bsr_col_ind, B, X info or temp_buffer pointer is invalid.

• rocsparse_status_internal_error – an internal error occurred.

• rocsparse_status_not_implemented – trans_A == rocsparse_operation_conjugate_transpose, trans_X == rocsparse_operation_conjugate_transpose or rocsparse_matrix_type != rocsparse_matrix_type_general.

rocsparse_bsrsm_clear()

rocsparse_status rocsparse_bsrsm_clear(rocsparse_handle handle, rocsparse_mat_info info)

rocsparse_bsrsm_clear deallocates all memory that was allocated by rocsparse_Xbsrsm_analysis(). This is especially useful, if memory is an issue and the analysis data is not required for further computation, e.g. when switching to another sparse matrix format. Calling rocsparse_bsrsm_clear is optional. All allocated resources will be cleared when the opaque rocsparse_mat_info struct is destroyed using rocsparse_destroy_mat_info().

Note

This routine does not support execution in a hipGraph context.

Parameters:

• handle – [in] handle to the rocsparse library context queue.

• info – [inout] structure that holds the information collected during the analysis step.

Return values:

• rocsparse_status_success – the operation completed successfully.

• rocsparse_status_invalid_handle – the library context was not initialized.

• rocsparse_status_invalid_pointer – info pointer is invalid.

• rocsparse_status_memory_error – the buffer holding the meta data could not be deallocated.

• rocsparse_status_internal_error – an internal error occurred.

rocsparse_gemmi()

rocsparse_status rocsparse_sgemmi(rocsparse_handle handle, rocsparse_operation trans_A, rocsparse_operation trans_B, rocsparse_int m, rocsparse_int n, rocsparse_int k, rocsparse_int nnz, const float *alpha, const float *A, rocsparse_int lda, const rocsparse_mat_descr descr, const float *csr_val, const rocsparse_int *csr_row_ptr, const rocsparse_int *csr_col_ind, const float *beta, float *C, rocsparse_int ldc)

rocsparse_status rocsparse_dgemmi(rocsparse_handle handle, rocsparse_operation trans_A, rocsparse_operation trans_B, rocsparse_int m, rocsparse_int n, rocsparse_int k, rocsparse_int nnz, const double *alpha, const double *A, rocsparse_int lda, const rocsparse_mat_descr descr, const double *csr_val, const rocsparse_int *csr_row_ptr, const rocsparse_int *csr_col_ind, const double *beta, double *C, rocsparse_int ldc)

rocsparse_status rocsparse_cgemmi(rocsparse_handle handle, rocsparse_operation trans_A, rocsparse_operation trans_B, rocsparse_int m, rocsparse_int n, rocsparse_int k, rocsparse_int nnz, const rocsparse_float_complex *alpha, const rocsparse_float_complex *A, rocsparse_int lda, const rocsparse_mat_descr descr, const rocsparse_float_complex *csr_val, const rocsparse_int *csr_row_ptr, const rocsparse_int *csr_col_ind, const rocsparse_float_complex *beta, rocsparse_float_complex *C, rocsparse_int ldc)

rocsparse_status rocsparse_zgemmi(rocsparse_handle handle, rocsparse_operation trans_A, rocsparse_operation trans_B, rocsparse_int m, rocsparse_int n, rocsparse_int k, rocsparse_int nnz, const rocsparse_double_complex *alpha, const rocsparse_double_complex *A, rocsparse_int lda, const rocsparse_mat_descr descr, const rocsparse_double_complex *csr_val, const rocsparse_int *csr_row_ptr, const rocsparse_int *csr_col_ind, const rocsparse_double_complex *beta, rocsparse_double_complex *C, rocsparse_int ldc)

Dense matrix sparse matrix multiplication using CSR storage format.

rocsparse_gemmi multiplies the scalar \(\alpha\) with a column-oriented dense \(m \times k\) matrix \(op(A)\) and the sparse \(k \times n\) matrix \(op(B)\), defined in CSR storage format 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 trans_A == rocsparse_operation_none} \\ A^T, & \text{if trans_A == rocsparse_operation_transpose} \\ A^H, & \text{if trans_A == rocsparse_operation_conjugate_transpose} \end{array} \right. \end{split}\]

and

\[\begin{split} op(B) = \left\{ \begin{array}{ll} B, & \text{if trans_B == rocsparse_operation_none} \\ B^T, & \text{if trans_B == rocsparse_operation_transpose} \\ B^H, & \text{if trans_B == rocsparse_operation_conjugate_transpose} \end{array} \right. \end{split}\]

Example

This example multiplies a column-oriented dense matrix with a CSC matrix.

int main()
{
    rocsparse_int m   = 2;
    rocsparse_int n   = 5;
    rocsparse_int k   = 3;
    rocsparse_int nnz = 8;
    rocsparse_int lda = m;
    rocsparse_int ldc = m;

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

    // Matrix A (m x k)
    // (  9.0  10.0  11.0 )
    // ( 12.0  13.0  14.0 )

    // Matrix B (k x n)
    // ( 1.0  2.0  0.0  3.0  0.0 )
    // ( 0.0  4.0  5.0  0.0  0.0 )
    // ( 6.0  0.0  0.0  7.0  8.0 )

    // Matrix C (m x n)
    // ( 15.0  16.0  17.0  18.0  19.0 )
    // ( 20.0  21.0  22.0  23.0  24.0 )

    std::vector<float> hA             = {9.0, 12.0, 10.0, 13.0, 11.0, 14.0};
    std::vector<int>   hcsc_col_ptr_B = {0, 2, 4, 5, 7, 8};
    std::vector<int>   hcsc_row_ind_B = {0, 2, 0, 1, 1, 0, 2, 2};
    std::vector<float> hcsc_val_B     = {1.0, 6.0, 2.0, 4.0, 5.0, 3.0, 7.0, 8.0};
    std::vector<float> hC = {15.0, 20.0, 16.0, 21.0, 17.0, 22.0, 18.0, 23.0, 19.0, 24.0};

    float* dA;
    int*   dcsc_col_ptr_B;
    int*   dcsc_row_ind_B;
    float* dcsc_val_B;
    float* dC;
    HIP_CHECK(hipMalloc(&dA, sizeof(float) * lda * k));
    HIP_CHECK(hipMalloc(&dcsc_col_ptr_B, sizeof(int) * (n + 1)));
    HIP_CHECK(hipMalloc(&dcsc_row_ind_B, sizeof(int) * nnz));
    HIP_CHECK(hipMalloc(&dcsc_val_B, sizeof(float) * nnz));
    HIP_CHECK(hipMalloc(&dC, sizeof(float) * ldc * n));

    HIP_CHECK(hipMemcpy(dA, hA.data(), sizeof(float) * lda * k, hipMemcpyHostToDevice));
    HIP_CHECK(hipMemcpy(
        dcsc_col_ptr_B, hcsc_col_ptr_B.data(), sizeof(int) * (n + 1), hipMemcpyHostToDevice));
    HIP_CHECK(
        hipMemcpy(dcsc_row_ind_B, hcsc_row_ind_B.data(), sizeof(int) * nnz, hipMemcpyHostToDevice));
    HIP_CHECK(hipMemcpy(dcsc_val_B, hcsc_val_B.data(), sizeof(float) * nnz, hipMemcpyHostToDevice));
    HIP_CHECK(hipMemcpy(dC, hC.data(), sizeof(float) * ldc * n, hipMemcpyHostToDevice));

    // Create rocSPARSE handle
    rocsparse_handle handle;
    ROCSPARSE_CHECK(rocsparse_create_handle(&handle));

    // Create matrix descriptor
    rocsparse_mat_descr descr_B;
    ROCSPARSE_CHECK(rocsparse_create_mat_descr(&descr_B));

    // Perform the matrix multiplication
    ROCSPARSE_CHECK(rocsparse_sgemmi(handle,
                                     rocsparse_operation_none,
                                     rocsparse_operation_transpose,
                                     m,
                                     n,
                                     k,
                                     nnz,
                                     &alpha,
                                     dA,
                                     lda,
                                     descr_B,
                                     dcsc_val_B,
                                     dcsc_col_ptr_B,
                                     dcsc_row_ind_B,
                                     &beta,
                                     dC,
                                     ldc));

    HIP_CHECK(hipMemcpy(hC.data(), dC, sizeof(float) * ldc * n, hipMemcpyDeviceToHost));

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

    // Clean up
    ROCSPARSE_CHECK(rocsparse_destroy_mat_descr(descr_B));
    ROCSPARSE_CHECK(rocsparse_destroy_handle(handle));

    // Clear up on device
    HIP_CHECK(hipFree(dA));
    HIP_CHECK(hipFree(dcsc_col_ptr_B));
    HIP_CHECK(hipFree(dcsc_row_ind_B));
    HIP_CHECK(hipFree(dcsc_val_B));
    HIP_CHECK(hipFree(dC));

    return 0;
}

Note

Currently, only trans_A == rocsparse_operation_none is supported.

Note

Currently, only trans_B == rocsparse_operation_transpose is supported.

Note

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

Note

This routine supports execution in a hipGraph context.

Parameters:

• handle – [in] handle to the rocsparse library context queue.

• trans_A – [in] matrix \(A\) operation type.

• trans_B – [in] matrix \(B\) operation type.

• m – [in] number of rows of the column-oriented dense matrix \(A\).

• n – [in] number of columns of the sparse CSR matrix \(op(B)\) and \(C\).

• k – [in] number of columns of the column-oriented dense matrix \(A\).

• nnz – [in] number of non-zero entries of the sparse CSR matrix \(B\).

• 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\)).

• descr – [in] descriptor of the sparse CSR matrix \(B\). Currently, only rocsparse_matrix_type_general is supported.

• csr_val – [in] array of nnz elements of the sparse CSR matrix \(B\).

• csr_row_ptr – [in] array of m+1 elements that point to the start of every row of the sparse CSR matrix \(B\).

• csr_col_ind – [in] array of nnz elements containing the column indices of the sparse CSR matrix \(B\).

• beta – [in] scalar \(\beta\).

• C – [inout] column-oriented dense matrix of dimension \(ldc \times n\) that holds the values of \(C\).

• ldc – [in] leading dimension of \(C\), must be at least \(m\).

Return values:

• rocsparse_status_success – the operation completed successfully.

• rocsparse_status_invalid_handle – the library context was not initialized.

• rocsparse_status_invalid_size – m, n, k, nnz, lda or ldc is invalid.

• rocsparse_status_invalid_pointer – alpha, A, csr_val, csr_row_ptr, csr_col_ind, beta or C pointer is invalid.