Skip to content

Commit

Permalink
[HIPIFY][ROCm#674][rocSPARSE][feature] rocSPARSE support - Step 102 -…
Browse files Browse the repository at this point in the history
… `cusparseSpMM_bufferSize` -> `rocsparse_spmm`

+ [IMP] `rocsparse_spmm` has been changed in 6.0.0, so reflected that in HIPIFY, docs, and tests
+ Updated `SPARSE` synthetic tests, the regenerated hipify-perl, and `SPARSE` `CUDA2HIP` documentation
  • Loading branch information
emankov committed Jan 21, 2024
1 parent cdaa3d2 commit e07a612
Show file tree
Hide file tree
Showing 10 changed files with 127 additions and 8 deletions.
1 change: 1 addition & 0 deletions bin/hipify-perl
Original file line number Diff line number Diff line change
Expand Up @@ -2598,6 +2598,7 @@ sub rocSubstitutions {
subst("cusparseShybmv", "rocsparse_shybmv", "library");
subst("cusparseSnnz", "rocsparse_snnz", "library");
subst("cusparseSnnz_compress", "rocsparse_snnz_compress", "library");
subst("cusparseSpMM_bufferSize", "rocsparse_spmm", "library");
subst("cusparseSpMV", "rocsparse_spmv", "library");
subst("cusparseSpMatGetAttribute", "rocsparse_spmat_get_attribute", "library");
subst("cusparseSpMatGetFormat", "rocsparse_spmat_get_format", "library");
Expand Down
2 changes: 1 addition & 1 deletion docs/tables/CUSPARSE_API_supported_by_HIP_and_ROC.md
Original file line number Diff line number Diff line change
Expand Up @@ -878,7 +878,7 @@
|`cusparseSpMMOp`|11.5| | | | | | | | | | | | | | | |
|`cusparseSpMMOp_createPlan`|11.5| | | | | | | | | | | | | | | |
|`cusparseSpMMOp_destroyPlan`|11.5| | | | | | | | | | | | | | | |
|`cusparseSpMM_bufferSize`|10.1| |12.0| |`hipsparseSpMM_bufferSize`|4.2.0| |6.0.0| | | | | | | | |
|`cusparseSpMM_bufferSize`|10.1| |12.0| |`hipsparseSpMM_bufferSize`|4.2.0| |6.0.0| | |`rocsparse_spmm`|4.2.0| |6.0.0| | |
|`cusparseSpMM_preprocess`|11.2| |12.0| |`hipsparseSpMM_preprocess`|4.5.0| |6.0.0| | | | | | | | |
|`cusparseSpMV`|10.2| |12.0| |`hipsparseSpMV`|4.1.0| |6.0.0| | |`rocsparse_spmv`|4.1.0| | | | |
|`cusparseSpMV_bufferSize`|10.2| |12.0| |`hipsparseSpMV_bufferSize`|4.1.0| |6.0.0| | | | | | | | |
Expand Down
2 changes: 1 addition & 1 deletion docs/tables/CUSPARSE_API_supported_by_ROC.md
Original file line number Diff line number Diff line change
Expand Up @@ -878,7 +878,7 @@
|`cusparseSpMMOp`|11.5| | | | | | | | | |
|`cusparseSpMMOp_createPlan`|11.5| | | | | | | | | |
|`cusparseSpMMOp_destroyPlan`|11.5| | | | | | | | | |
|`cusparseSpMM_bufferSize`|10.1| |12.0| | | | | | | |
|`cusparseSpMM_bufferSize`|10.1| |12.0| |`rocsparse_spmm`|4.2.0| |6.0.0| | |
|`cusparseSpMM_preprocess`|11.2| |12.0| | | | | | | |
|`cusparseSpMV`|10.2| |12.0| |`rocsparse_spmv`|4.1.0| | | | |
|`cusparseSpMV_bufferSize`|10.2| |12.0| | | | | | | |
Expand Down
8 changes: 7 additions & 1 deletion src/CUDA2HIP_SPARSE_API_functions.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -825,8 +825,10 @@ const std::map<llvm::StringRef, hipCounter> CUDA_SPARSE_FUNCTION_MAP {
{"cusparseSpSV_updateMatrix", {"hipsparseSpSV_updateMatrix", "", CONV_LIB_FUNC, API_SPARSE, 15, UNSUPPORTED}},

// Sparse Matrix * Matrix Multiplication
// TODO: hipification cusparseSpMM into rocsparse_spmm needs additional variable declared and allocated
{"cusparseSpMM", {"hipsparseSpMM", "", CONV_LIB_FUNC, API_SPARSE, 15, ROC_UNSUPPORTED}},
{"cusparseSpMM_bufferSize", {"hipsparseSpMM_bufferSize", "", CONV_LIB_FUNC, API_SPARSE, 15, ROC_UNSUPPORTED}},
{"cusparseSpMM_bufferSize", {"hipsparseSpMM_bufferSize", "rocsparse_spmm", CONV_LIB_FUNC, API_SPARSE, 15}},
// TODO: hipification cusparseSpMM_preprocess into rocsparse_spmm needs additional variable declared and allocated
{"cusparseSpMM_preprocess", {"hipsparseSpMM_preprocess", "", CONV_LIB_FUNC, API_SPARSE, 15, ROC_UNSUPPORTED}},
{"cusparseSpMMOp", {"hipsparseSpMMOp", "", CONV_LIB_FUNC, API_SPARSE, 15, UNSUPPORTED}},
{"cusparseSpMMOp_createPlan", {"hipsparseSpMMOp_createPlan", "", CONV_LIB_FUNC, API_SPARSE, 15, UNSUPPORTED}},
Expand All @@ -835,6 +837,7 @@ const std::map<llvm::StringRef, hipCounter> CUDA_SPARSE_FUNCTION_MAP {
// Sparse Triangular Matrix Solve
{"cusparseSpSM_createDescr", {"hipsparseSpSM_createDescr", "", CONV_LIB_FUNC, API_SPARSE, 15, ROC_UNSUPPORTED}},
{"cusparseSpSM_destroyDescr", {"hipsparseSpSM_destroyDescr", "", CONV_LIB_FUNC, API_SPARSE, 15, ROC_UNSUPPORTED}},
// NTOE: Additional calculations are needed after calling rocsparse_spsm
{"cusparseSpSM_bufferSize", {"hipsparseSpSM_bufferSize", "", CONV_LIB_FUNC, API_SPARSE, 15, ROC_UNSUPPORTED}},
{"cusparseSpSM_analysis", {"hipsparseSpSM_analysis", "", CONV_LIB_FUNC, API_SPARSE, 15, ROC_UNSUPPORTED}},
{"cusparseSpSM_solve", {"hipsparseSpSM_solve", "", CONV_LIB_FUNC, API_SPARSE, 15, ROC_UNSUPPORTED}},
Expand Down Expand Up @@ -867,6 +870,7 @@ const std::map<llvm::StringRef, hipCounter> CUDA_SPARSE_FUNCTION_MAP {
{"cusparseSparseToDense_bufferSize", {"hipsparseSparseToDense_bufferSize", "rocsparse_sparse_to_dense", CONV_LIB_FUNC, API_SPARSE, 15}},
{"cusparseDenseToSparse_bufferSize", {"hipsparseDenseToSparse_bufferSize", "rocsparse_dense_to_sparse", CONV_LIB_FUNC, API_SPARSE, 15}},
{"cusparseDenseToSparse_analysis", {"hipsparseDenseToSparse_analysis", "rocsparse_dense_to_sparse", CONV_LIB_FUNC, API_SPARSE, 15}},
// TODO: hipification cusparseDenseToSparse_convert into rocsparse_dense_to_sparse needs additional variable declared and allocated
{"cusparseDenseToSparse_convert", {"hipsparseDenseToSparse_convert", "", CONV_LIB_FUNC, API_SPARSE, 15, ROC_UNSUPPORTED}},

// Sampled Dense-dense Matrix Multiplication
Expand Down Expand Up @@ -2425,6 +2429,7 @@ const std::map<llvm::StringRef, hipAPIversions> HIP_SPARSE_FUNCTION_VER_MAP {
{"rocsparse_csr2csc_buffer_size", {HIP_1090, HIP_0, HIP_0 }},
{"rocsparse_sparse_to_dense", {HIP_4010, HIP_0, HIP_0 }},
{"rocsparse_dense_to_sparse", {HIP_4010, HIP_0, HIP_0 }},
{"rocsparse_spmm", {HIP_4020, HIP_0, HIP_0 }},
};

const std::map<llvm::StringRef, cudaAPIChangedVersions> CUDA_SPARSE_FUNCTION_CHANGED_VER_MAP {
Expand Down Expand Up @@ -2529,6 +2534,7 @@ const std::map<llvm::StringRef, hipAPIChangedVersions> HIP_SPARSE_FUNCTION_CHANG
{"rocsparse_dnmat_get_strided_batch", {HIP_6000}},
{"rocsparse_sparse_to_dense", {HIP_6000}},
{"rocsparse_dense_to_sparse", {HIP_6000}},
{"rocsparse_spmm", {HIP_6000}},
};

const std::map<unsigned int, llvm::StringRef> CUDA_SPARSE_API_SECTION_MAP {
Expand Down
14 changes: 13 additions & 1 deletion src/HipifyAction.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -208,6 +208,7 @@ const std::string sCusparseSparseToDense = "cusparseSparseToDense";
const std::string sCusparseSparseToDense_bufferSize = "cusparseSparseToDense_bufferSize";
const std::string sCusparseDenseToSparse_bufferSize = "cusparseDenseToSparse_bufferSize";
const std::string sCusparseDenseToSparse_analysis = "cusparseDenseToSparse_analysis";
const std::string sCusparseSpMM_bufferSize = "cusparseSpMM_bufferSize";

// CUDA_OVERLOADED
const std::string sCudaEventCreate = "cudaEventCreate";
Expand Down Expand Up @@ -1671,6 +1672,16 @@ std::map<std::string, ArgCastStruct> FuncArgCasts {
false
}
},
{sCusparseSpMM_bufferSize,
{
{
{10, {e_add_const_argument, cw_None, "rocsparse_spmm_stage_compute"}},
{12, {e_add_const_argument, cw_None, "nullptr"}}
},
true,
false
}
},
};

void HipifyAction::RewriteString(StringRef s, clang::SourceLocation start) {
Expand Down Expand Up @@ -2519,7 +2530,8 @@ std::unique_ptr<clang::ASTConsumer> HipifyAction::CreateASTConsumer(clang::Compi
sCusparseSparseToDense,
sCusparseSparseToDense_bufferSize,
sCusparseDenseToSparse_bufferSize,
sCusparseDenseToSparse_analysis
sCusparseDenseToSparse_analysis,
sCusparseSpMM_bufferSize
)
)
)
Expand Down
3 changes: 3 additions & 0 deletions tests/lit.cfg
Original file line number Diff line number Diff line change
Expand Up @@ -77,6 +77,9 @@ if config.cuda_version_major < 10:
config.excludes.append('simple_mechs.cu')
config.excludes.append('cusparse2rocsparse_10000.cu')

if config.cuda_version_major < 10 or (config.cuda_version_major == 10 and config.cuda_version_minor < 1) or config.cuda_version_major >= 12:
config.excludes.append('cusparse2rocsparse_10010_12000.cu')

if config.cuda_version_major <= 10:
config.excludes.append('headers_test_12_SOLVER_10010.cu')

Expand Down
1 change: 0 additions & 1 deletion tests/unit_tests/synthetic/libraries/cusparse2hipsparse.cu
Original file line number Diff line number Diff line change
Expand Up @@ -1942,7 +1942,6 @@ int main() {
// CHECK: status_t = hipsparseDnMatGetStridedBatch(dnMatDescr_t, &batchCount, &batchStride);
status_t = cusparseDnMatGetStridedBatch(dnMatDescr_t, &batchCount, &batchStride);

// TODO: Mark as C-Changed in 12.0.0
// CUDA: cusparseStatus_t CUSPARSEAPI cusparseSpMM_bufferSize(cusparseHandle_t handle, cusparseOperation_t opA, cusparseOperation_t opB, const void* alpha, const cusparseSpMatDescr_t matA, const cusparseDnMatDescr_t matB, const void* beta, cusparseDnMatDescr_t matC, cudaDataType computeType, cusparseSpMMAlg_t alg, size_t* bufferSize);
// HIP: hipsparseStatus_t hipsparseSpMM_bufferSize(hipsparseHandle_t handle, hipsparseOperation_t opA, hipsparseOperation_t opB, const void* alpha, const hipsparseSpMatDescr_t matA, const hipsparseDnMatDescr_t matB, const void* beta, const hipsparseDnMatDescr_t matC, hipDataType computeType, hipsparseSpMMAlg_t alg, size_t* bufferSize);
// CHECK: status_t = hipsparseSpMM_bufferSize(handle_t, opA, opB, alpha, spmatA, dnmatB, beta, dnmatC, dataType, spMMAlg_t, &bufferSize);
Expand Down
Original file line number Diff line number Diff line change
@@ -0,0 +1,84 @@
// RUN: %run_test hipify "%s" "%t" %hipify_args 4 --skip-excluded-preprocessor-conditional-blocks --experimental --roc --use-hip-data-types %clang_args -ferror-limit=500

// CHECK: #include <hip/hip_runtime.h>
#include <cuda_runtime.h>
// CHECK: #include "hip/hip_complex.h"
#include "cuComplex.h"
#include <stdio.h>
// CHECK: #include "rocsparse.h"
#include "cusparse.h"
// CHECK-NOT: #include "rocsparse.h"

int main() {
printf("18.1. cuSPARSE API to rocSPARSE API synthetic test\n");

// CHECK: _rocsparse_handle *handle = nullptr;
// CHECK-NEXT: rocsparse_handle handle_t;
cusparseContext *handle = nullptr;
cusparseHandle_t handle_t;

// CHECK: rocsparse_status status_t;
cusparseStatus_t status_t;

// CHECK: _rocsparse_mat_descr *matDescr = nullptr;
// CHECK-NEXT: rocsparse_mat_descr matDescr_t, matDescr_t_2, matDescr_A, matDescr_B, matDescr_C, matDescr_D;
cusparseMatDescr *matDescr = nullptr;
cusparseMatDescr_t matDescr_t, matDescr_t_2, matDescr_A, matDescr_B, matDescr_C, matDescr_D;

// CHECK: rocsparse_action action_t;
cusparseAction_t action_t;

// CHECK: rocsparse_index_base indexBase_t;
cusparseIndexBase_t indexBase_t;

int m = 0;
int n = 0;
int innz = 0;
int csrRowPtrA = 0;
int csrRowPtrB = 0;
int csrRowPtrC = 0;
int cscRowIndA = 0;
int csrColIndA = 0;
int csrColIndB = 0;
int csrColIndC = 0;
int cscColPtrA = 0;
size_t bufferSize = 0;
void *pcsrVal = nullptr;
void *pcscVal = nullptr;
void *alpha = nullptr;
void *beta = nullptr;

// CHECK: rocsparse_operation opA, opB, opX;
cusparseOperation_t opA, opB, opX;

#if CUDA_VERSION >= 8000
// TODO: [#899] There should be rocsparse_datatype instead of hipDataType
cudaDataType_t dataType_t;
cudaDataType dataType;
#endif

#if CUDA_VERSION >= 10010
// TODO: cusparseCsr2CscAlg_t has no analogue in rocSPARSE. The deletion of declaration and usage is needed to be implemented
cusparseCsr2CscAlg_t Csr2CscAlg_t;

#if (CUDA_VERSION < 11000 && !defined(_WIN32)) || CUDA_VERSION >= 11000
// CHECK: rocsparse_spmat_descr spMatDescr_t, spmatA, spmatB, spmatC;
cusparseSpMatDescr_t spMatDescr_t, spmatA, spmatB, spmatC;

// CHECK: rocsparse_dnmat_descr dnMatDescr_t, dnmatA, dnmatB, dnmatC;
cusparseDnMatDescr_t dnMatDescr_t, dnmatA, dnmatB, dnmatC;

// CHECK: rocsparse_spmm_alg spMMAlg_t;
cusparseSpMMAlg_t spMMAlg_t;

#if CUDA_VERSION < 12000
// CUDA: cusparseStatus_t CUSPARSEAPI cusparseSpMM_bufferSize(cusparseHandle_t handle, cusparseOperation_t opA, cusparseOperation_t opB, const void* alpha, const cusparseSpMatDescr_t matA, const cusparseDnMatDescr_t matB, const void* beta, cusparseDnMatDescr_t matC, cudaDataType computeType, cusparseSpMMAlg_t alg, size_t* bufferSize);
// ROC: ROCSPARSE_EXPORT rocsparse_status rocsparse_spmm(rocsparse_handle handle, rocsparse_operation trans_A, rocsparse_operation trans_B, const void* alpha, const rocsparse_spmat_descr mat_A, const rocsparse_dnmat_descr mat_B, const void* beta, const rocsparse_dnmat_descr mat_C, rocsparse_datatype compute_type, rocsparse_spmm_alg alg, rocsparse_spmm_stage stage, size_t* buffer_size, void* temp_buffer);
// CHECK: status_t = rocsparse_spmm(handle_t, opA, opB, alpha, spmatA, dnmatB, beta, dnmatC, dataType, spMMAlg_t, rocsparse_spmm_stage_compute, &bufferSize, nullptr);
status_t = cusparseSpMM_bufferSize(handle_t, opA, opB, alpha, spmatA, dnmatB, beta, dnmatC, dataType, spMMAlg_t, &bufferSize);
#endif
#endif
#endif

return 0;
}
16 changes: 16 additions & 0 deletions tests/unit_tests/synthetic/libraries/cusparse2rocsparse_12000.cu
Original file line number Diff line number Diff line change
Expand Up @@ -69,6 +69,8 @@ int main() {
float csrSortedValD = 0.f;
void *pBuffer = nullptr;
void *tempBuffer = nullptr;
void *alpha = nullptr;
void *beta = nullptr;

// TODO: should be rocsparse_double_complex
// TODO: add to TypeOverloads cuDoubleComplex -> rocsparse_double_complex under a new option --sparse
Expand All @@ -80,12 +82,21 @@ int main() {
// CHECK: rocblas_float_complex complex, complexA, complexAlpha, complexB, complexBeta, complexC, complexF, complexX, complexY, complexbsrValA, complexbsrSortedValC, complexcsrSortedValA, complexcsrSortedValB, complexcsrSortedValC, complexcsrSortedValD, complextol, complexbsrSortedVal, complexbscVal, complexcscSortedVal, complexds, complexdl, complexd, complexdu, complexdw, complexx, complex_boost_val;
cuComplex complex, complexA, complexAlpha, complexB, complexBeta, complexC, complexF, complexX, complexY, complexbsrValA, complexbsrSortedValC, complexcsrSortedValA, complexcsrSortedValB, complexcsrSortedValC, complexcsrSortedValD, complextol, complexbsrSortedVal, complexbscVal, complexcscSortedVal, complexds, complexdl, complexd, complexdu, complexdw, complexx, complex_boost_val;

#if CUDA_VERSION >= 8000
// TODO: [#899] There should be rocsparse_datatype instead of hipDataType
cudaDataType_t dataType_t;
cudaDataType dataType;
#endif

#if (CUDA_VERSION >= 10010 && CUDA_VERSION < 11000 && !defined(_WIN32)) || CUDA_VERSION >= 11000
// CHECK: rocsparse_spmat_descr spMatDescr_t, spmatA, spmatB, spmatC;
cusparseSpMatDescr_t spMatDescr_t, spmatA, spmatB, spmatC;

// CHECK: rocsparse_dnmat_descr dnMatDescr_t, dnmatA, dnmatB, dnmatC;
cusparseDnMatDescr_t dnMatDescr_t, dnmatA, dnmatB, dnmatC;

// CHECK: rocsparse_spmm_alg spMMAlg_t;
cusparseSpMMAlg_t spMMAlg_t;
#endif

#if CUDA_VERSION >= 11010 && CUSPARSE_VERSION >= 11300
Expand Down Expand Up @@ -126,6 +137,11 @@ int main() {
// ROC: ROCSPARSE_EXPORT rocsparse_status rocsparse_dense_to_sparse(rocsparse_handle handle, rocsparse_const_dnmat_descr mat_A, rocsparse_spmat_descr mat_B, rocsparse_dense_to_sparse_alg alg, size_t* buffer_size, void* temp_buffer);
// CHECK: status_t = rocsparse_dense_to_sparse(handle_t, constDnMatDescr, spmatB, denseToSparseAlg_t, nullptr, tempBuffer);
status_t = cusparseDenseToSparse_analysis(handle_t, constDnMatDescr, spmatB, denseToSparseAlg_t, tempBuffer);

// CUDA: cusparseStatus_t CUSPARSEAPI cusparseSpMM_bufferSize(cusparseHandle_t handle, cusparseOperation_t opA, cusparseOperation_t opB, const void* alpha, cusparseConstSpMatDescr_t matA, cusparseConstDnMatDescr_t matB, const void* beta, cusparseDnMatDescr_t matC, cudaDataType computeType, cusparseSpMMAlg_t alg, size_t* bufferSize);
// ROC: ROCSPARSE_EXPORT rocsparse_status rocsparse_spmm(rocsparse_handle handle, rocsparse_operation trans_A, rocsparse_operation trans_B, const void* alpha, rocsparse_const_spmat_descr mat_A, rocsparse_const_dnmat_descr mat_B, const void* beta, const rocsparse_dnmat_descr mat_C, rocsparse_datatype compute_type, rocsparse_spmm_alg alg, rocsparse_spmm_stage stage, size_t* buffer_size, void* temp_buffer);
// CHECK: status_t = rocsparse_spmm(handle_t, opA, opB, alpha, constSpMatDescr, constDnMatDescr, beta, dnmatC, dataType, spMMAlg_t, rocsparse_spmm_stage_compute, &bufferSize, nullptr);
status_t = cusparseSpMM_bufferSize(handle_t, opA, opB, alpha, constSpMatDescr, constDnMatDescr, beta, dnmatC, dataType, spMMAlg_t, &bufferSize);
#endif

return 0;
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -75,8 +75,7 @@ int main() {
cusparseSolvePolicy_t SOLVE_POLICY_NO_LEVEL = CUSPARSE_SOLVE_POLICY_NO_LEVEL;
cusparseSolvePolicy_t SOLVE_POLICY_USE_LEVEL = CUSPARSE_SOLVE_POLICY_USE_LEVEL;

#if CUDA_VERSION >= 9020
#if CUDA_VERSION < 12000
#if CUDA_VERSION >= 9020 && CUDA_VERSION < 12000
// CHECK: rocsparse_mat_info csrsm2_info;
csrsm2Info_t csrsm2_info;

Expand Down Expand Up @@ -155,7 +154,6 @@ int main() {
// ROC: ROCSPARSE_EXPORT rocsparse_status rocsparse_destroy_mat_info(rocsparse_mat_info info);
// CHECK: status_t = rocsparse_destroy_mat_info(csrsm2_info);
status_t = cusparseDestroyCsrsm2Info(csrsm2_info);
#endif
#endif

return 0;
Expand Down

0 comments on commit e07a612

Please sign in to comment.