From e18e3675e225e00f1b92b5cf0901f6eb7d410c71 Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Wed, 24 Jan 2024 15:27:30 +0000 Subject: [PATCH] [HIPIFY][#674][rocSPARSE][feature] rocSPARSE support - Step 103 - `cusparseSpSM_analysis` -> `rocsparse_spsm` + [IMP] `rocsparse_spsm` 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 --- bin/hipify-perl | 1 + .../CUSPARSE_API_supported_by_HIP_and_ROC.md | 2 +- docs/tables/CUSPARSE_API_supported_by_ROC.md | 2 +- src/CUDA2HIP_SPARSE_API_functions.cpp | 4 +- src/HipifyAction.cpp | 14 ++- tests/lit.cfg | 5 + .../synthetic/libraries/cusparse2hipsparse.cu | 1 - .../cusparse2rocsparse_11030_12000.cu | 114 ++++++++++++++++++ .../libraries/cusparse2rocsparse_12000.cu | 15 +++ 9 files changed, 153 insertions(+), 5 deletions(-) create mode 100644 tests/unit_tests/synthetic/libraries/cusparse2rocsparse_11030_12000.cu diff --git a/bin/hipify-perl b/bin/hipify-perl index ce65876b..0540fb03 100755 --- a/bin/hipify-perl +++ b/bin/hipify-perl @@ -2609,6 +2609,7 @@ sub rocSubstitutions { subst("cusparseSpMatSetAttribute", "rocsparse_spmat_set_attribute", "library"); subst("cusparseSpMatSetStridedBatch", "rocsparse_spmat_set_strided_batch", "library"); subst("cusparseSpMatSetValues", "rocsparse_spmat_set_values", "library"); + subst("cusparseSpSM_analysis", "rocsparse_spsm", "library"); subst("cusparseSpVecGet", "rocsparse_spvec_get", "library"); subst("cusparseSpVecGetIndexBase", "rocsparse_spvec_get_index_base", "library"); subst("cusparseSpVecGetValues", "rocsparse_spvec_get_values", "library"); diff --git a/docs/tables/CUSPARSE_API_supported_by_HIP_and_ROC.md b/docs/tables/CUSPARSE_API_supported_by_HIP_and_ROC.md index c2f1c3f6..db9e930e 100644 --- a/docs/tables/CUSPARSE_API_supported_by_HIP_and_ROC.md +++ b/docs/tables/CUSPARSE_API_supported_by_HIP_and_ROC.md @@ -893,7 +893,7 @@ |`cusparseSpMatSetNumBatches`|10.1| | |10.2| | | | | | | | | | | | | |`cusparseSpMatSetStridedBatch`|10.2| | |12.0|`hipsparseSpMatSetStridedBatch`|5.2.0| | | | |`rocsparse_spmat_set_strided_batch`|5.2.0| | | | | |`cusparseSpMatSetValues`|10.2| | | |`hipsparseSpMatSetValues`|4.1.0| | | | |`rocsparse_spmat_set_values`|4.1.0| | | | | -|`cusparseSpSM_analysis`|11.3| |12.0| |`hipsparseSpSM_analysis`|4.5.0| |6.0.0| | | | | | | | | +|`cusparseSpSM_analysis`|11.3| |12.0| |`hipsparseSpSM_analysis`|4.5.0| |6.0.0| | |`rocsparse_spsm`|4.5.0| |6.0.0| | | |`cusparseSpSM_bufferSize`|11.3| |12.0| |`hipsparseSpSM_bufferSize`|4.5.0| |6.0.0| | | | | | | | | |`cusparseSpSM_createDescr`|11.3| | | |`hipsparseSpSM_createDescr`|4.5.0| | | | | | | | | | | |`cusparseSpSM_destroyDescr`|11.3| | | |`hipsparseSpSM_destroyDescr`|4.5.0| | | | | | | | | | | diff --git a/docs/tables/CUSPARSE_API_supported_by_ROC.md b/docs/tables/CUSPARSE_API_supported_by_ROC.md index f45190eb..7f9f2221 100644 --- a/docs/tables/CUSPARSE_API_supported_by_ROC.md +++ b/docs/tables/CUSPARSE_API_supported_by_ROC.md @@ -893,7 +893,7 @@ |`cusparseSpMatSetNumBatches`|10.1| | |10.2| | | | | | | |`cusparseSpMatSetStridedBatch`|10.2| | |12.0|`rocsparse_spmat_set_strided_batch`|5.2.0| | | | | |`cusparseSpMatSetValues`|10.2| | | |`rocsparse_spmat_set_values`|4.1.0| | | | | -|`cusparseSpSM_analysis`|11.3| |12.0| | | | | | | | +|`cusparseSpSM_analysis`|11.3| |12.0| |`rocsparse_spsm`|4.5.0| |6.0.0| | | |`cusparseSpSM_bufferSize`|11.3| |12.0| | | | | | | | |`cusparseSpSM_createDescr`|11.3| | | | | | | | | | |`cusparseSpSM_destroyDescr`|11.3| | | | | | | | | | diff --git a/src/CUDA2HIP_SPARSE_API_functions.cpp b/src/CUDA2HIP_SPARSE_API_functions.cpp index 533e8da9..42f7e975 100644 --- a/src/CUDA2HIP_SPARSE_API_functions.cpp +++ b/src/CUDA2HIP_SPARSE_API_functions.cpp @@ -839,7 +839,7 @@ const std::map CUDA_SPARSE_FUNCTION_MAP { {"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_analysis", {"hipsparseSpSM_analysis", "rocsparse_spsm", CONV_LIB_FUNC, API_SPARSE, 15}}, {"cusparseSpSM_solve", {"hipsparseSpSM_solve", "", CONV_LIB_FUNC, API_SPARSE, 15, ROC_UNSUPPORTED}}, // Sparse Matrix Multiplication (SpGEMM) Structure Reuse @@ -2430,6 +2430,7 @@ const std::map HIP_SPARSE_FUNCTION_VER_MAP { {"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 }}, + {"rocsparse_spsm", {HIP_4050, HIP_0, HIP_0 }}, }; const std::map CUDA_SPARSE_FUNCTION_CHANGED_VER_MAP { @@ -2535,6 +2536,7 @@ const std::map HIP_SPARSE_FUNCTION_CHANG {"rocsparse_sparse_to_dense", {HIP_6000}}, {"rocsparse_dense_to_sparse", {HIP_6000}}, {"rocsparse_spmm", {HIP_6000}}, + {"rocsparse_spsm", {HIP_6000}}, }; const std::map CUDA_SPARSE_API_SECTION_MAP { diff --git a/src/HipifyAction.cpp b/src/HipifyAction.cpp index 2d0927f5..acfd3da4 100644 --- a/src/HipifyAction.cpp +++ b/src/HipifyAction.cpp @@ -209,6 +209,7 @@ const std::string sCusparseSparseToDense_bufferSize = "cusparseSparseToDense_buf const std::string sCusparseDenseToSparse_bufferSize = "cusparseDenseToSparse_bufferSize"; const std::string sCusparseDenseToSparse_analysis = "cusparseDenseToSparse_analysis"; const std::string sCusparseSpMM_bufferSize = "cusparseSpMM_bufferSize"; +const std::string sCusparseSpSM_analysis = "cusparseSpSM_analysis"; // CUDA_OVERLOADED const std::string sCudaEventCreate = "cudaEventCreate"; @@ -1682,6 +1683,16 @@ std::map FuncArgCasts { false } }, + {sCusparseSpSM_analysis, + { + { + {9, {e_replace_argument_with_const, cw_None, "rocsparse_spsm_stage_compute"}}, + {10, {e_add_const_argument, cw_None, "nullptr"}} + }, + true, + false + } + }, }; void HipifyAction::RewriteString(StringRef s, clang::SourceLocation start) { @@ -2531,7 +2542,8 @@ std::unique_ptr HipifyAction::CreateASTConsumer(clang::Compi sCusparseSparseToDense_bufferSize, sCusparseDenseToSparse_bufferSize, sCusparseDenseToSparse_analysis, - sCusparseSpMM_bufferSize + sCusparseSpMM_bufferSize, + sCusparseSpSM_analysis ) ) ) diff --git a/tests/lit.cfg b/tests/lit.cfg index 45f0cd32..9cbcf3c4 100644 --- a/tests/lit.cfg +++ b/tests/lit.cfg @@ -91,6 +91,11 @@ if config.cuda_version_major < 11 and sys.platform in ['win32']: if config.cuda_version_major < 11 or (config.cuda_version_major == 11 and config.cuda_version_minor <= 1) or config.cuda_version_major >= 12: config.excludes.append('cusparse2rocsparse_11010_12000.cu') +# [NOTE] cmake doesn't load lib VERSIONs, including CUSPARSE_VERSION, which equals 11500 for CUDA 11.3.0 and 11600 for CUDA 11.3.1 +# [NOTE] Both CUDA 11.3.0 CUDA 11.3.1 have the same CUDA_VERSION 11030, and we can't distinguish them, thus exclude the below tests from both 11.3.0 and 11.3.1 +if config.cuda_version_major < 11 or (config.cuda_version_major == 11 and config.cuda_version_minor <= 3) or config.cuda_version_major >= 12: + config.excludes.append('cusparse2rocsparse_11030_12000.cu') + if config.cuda_version_major <= 10: config.excludes.append('headers_test_12_SOLVER_10010.cu') diff --git a/tests/unit_tests/synthetic/libraries/cusparse2hipsparse.cu b/tests/unit_tests/synthetic/libraries/cusparse2hipsparse.cu index 98adb182..fc799751 100644 --- a/tests/unit_tests/synthetic/libraries/cusparse2hipsparse.cu +++ b/tests/unit_tests/synthetic/libraries/cusparse2hipsparse.cu @@ -2756,7 +2756,6 @@ int main() { // CHECK: status_t = hipsparseSpSM_bufferSize(handle_t, opA, opB, alpha, spmatA, dnmatB, dnmatC, dataType, spSMAlg_t, spSMDescr, &bufferSize); status_t = cusparseSpSM_bufferSize(handle_t, opA, opB, alpha, spmatA, dnmatB, dnmatC, dataType, spSMAlg_t, spSMDescr, &bufferSize); - // TODO: Mark as C-Changed in 12.0.0 // CUDA: cusparseStatus_t CUSPARSEAPI cusparseSpSM_analysis(cusparseHandle_t handle, cusparseOperation_t opA, cusparseOperation_t opB, const void* alpha, cusparseSpMatDescr_t matA, cusparseDnMatDescr_t matB, cusparseDnMatDescr_t matC, cudaDataType computeType, cusparseSpSMAlg_t alg, cusparseSpSMDescr_t spsmDescr, void* externalBuffer); // HIP: HIPSPARSE_EXPORT hipsparseStatus_t hipsparseSpSM_analysis(hipsparseHandle_t handle, hipsparseOperation_t opA, hipsparseOperation_t opB, const void* alpha, const hipsparseSpMatDescr_t matA, const hipsparseDnMatDescr_t matB, const hipsparseDnMatDescr_t matC, hipDataType computeType, hipsparseSpSMAlg_t alg, hipsparseSpSMDescr_t spsmDescr, void* externalBuffer); // CHECK: status_t = hipsparseSpSM_analysis(handle_t, opA, opB, alpha, spmatA, dnmatB, dnmatC, dataType, spSMAlg_t, spSMDescr, tempBuffer); diff --git a/tests/unit_tests/synthetic/libraries/cusparse2rocsparse_11030_12000.cu b/tests/unit_tests/synthetic/libraries/cusparse2rocsparse_11030_12000.cu new file mode 100644 index 00000000..41c4edb5 --- /dev/null +++ b/tests/unit_tests/synthetic/libraries/cusparse2rocsparse_11030_12000.cu @@ -0,0 +1,114 @@ +// 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 +#include +// CHECK: #include "hip/hip_complex.h" +#include "cuComplex.h" +#include +// 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_status status_t; + cusparseStatus_t status_t; + + // CHECK: _rocsparse_handle *handle = nullptr; + // CHECK-NEXT: rocsparse_handle handle_t; + cusparseContext *handle = nullptr; + cusparseHandle_t handle_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_operation opA, opB, opX; + cusparseOperation_t opA, opB, opX; + + // CHECK: rocsparse_solve_policy solvePolicy_t; + cusparseSolvePolicy_t solvePolicy_t; + + int m = 0; + int n = 0; + int k = 0; + int innz = 0; + int nnza = 0; + int nnzb = 0; + int nnzc = 0; + int nnzd = 0; + int csrRowPtrA = 0; + int csrRowPtrB = 0; + int csrRowPtrC = 0; + int csrRowPtrD = 0; + int csrColIndA = 0; + int csrColIndB = 0; + int csrColIndC = 0; + int csrColIndD = 0; + int bufferSizeInBytes = 0; + size_t bufferSize = 0; + double dA = 0.f; + double dB = 0.f; + double dAlpha = 0.f; + double dF = 0.f; + double dX = 0.f; + double dcsrSortedValA = 0.f; + double dcsrSortedValB = 0.f; + double dcsrSortedValC = 0.f; + double dcsrSortedValD = 0.f; + float fAlpha = 0.f; + float fA = 0.f; + float fB = 0.f; + float fF = 0.f; + float fX = 0.f; + float csrSortedValA = 0.f; + float csrSortedValB = 0.f; + float csrSortedValC = 0.f; + float csrSortedValD = 0.f; + void *alpha = nullptr; + void *pBuffer = nullptr; + void *tempBuffer = nullptr; + + // TODO: should be rocsparse_double_complex + // TODO: add to TypeOverloads cuDoubleComplex -> rocsparse_double_complex under a new option --sparse + // CHECK: rocblas_double_complex dcomplex, dcomplexA, dcomplexAlpha, dcomplexB, dcomplexBeta, dcomplexC, dcomplexF, dcomplexX, dcomplexY, dComplexbsrSortedValA, dComplexbsrSortedValC, dComplexcsrSortedValA, dComplexcsrSortedValB, dComplexcsrSortedValC, dComplexcsrSortedValD, dcomplextol, dComplexbsrSortedVal, dComplexbscVal, dComplexcscSortedVal, dcomplexds, dcomplexdl, dcomplexd, dcomplexdu, dcomplexdw, dcomplexx, dcomplex_boost_val; + cuDoubleComplex dcomplex, dcomplexA, dcomplexAlpha, dcomplexB, dcomplexBeta, dcomplexC, dcomplexF, dcomplexX, dcomplexY, dComplexbsrSortedValA, dComplexbsrSortedValC, dComplexcsrSortedValA, dComplexcsrSortedValB, dComplexcsrSortedValC, dComplexcsrSortedValD, dcomplextol, dComplexbsrSortedVal, dComplexbscVal, dComplexcscSortedVal, dcomplexds, dcomplexdl, dcomplexd, dcomplexdu, dcomplexdw, dcomplexx, dcomplex_boost_val; + + // TODO: should be rocsparse_double_complex + // TODO: add to TypeOverloads cuComplex -> rocsparse_float_complex under a new option --sparse + // 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; +#endif + +#if CUDA_VERSION >= 11030 && CUSPARSE_VERSION >= 11600 + // CHECK: rocsparse_spsm_alg spSMAlg_t; + cusparseSpSMAlg_t spSMAlg_t; + + // NOTE:cusparseSpSMDescr_t doesn't have a correspondence in rocSPARSE, the corresponding function argument is removed in the hipified call of the rocsparse_spsm function + cusparseSpSMDescr_t spSMDescr; + +#if CUDA_VERSION < 12000 + // CUDA: cusparseStatus_t CUSPARSEAPI cusparseSpSM_analysis(cusparseHandle_t handle, cusparseOperation_t opA, cusparseOperation_t opB, const void* alpha, cusparseSpMatDescr_t matA, cusparseDnMatDescr_t matB, cusparseDnMatDescr_t matC, cudaDataType computeType, cusparseSpSMAlg_t alg, cusparseSpSMDescr_t spsmDescr, void* externalBuffer); + // ROC: ROCSPARSE_EXPORT rocsparse_status rocsparse_spsm(rocsparse_handle handle, rocsparse_operation trans_A, rocsparse_operation trans_B, const void* alpha, const rocsparse_spmat_descr matA, const rocsparse_dnmat_descr matB, const rocsparse_dnmat_descr matC, rocsparse_datatype compute_type, rocsparse_spsm_alg alg, rocsparse_spsm_stage stage, size_t* buffer_size, void* temp_buffer); + // CHECK: status_t = rocsparse_spsm(handle_t, opA, opB, alpha, spmatA, dnmatB, dnmatC, dataType, spSMAlg_t, rocsparse_spsm_stage_compute, nullptr, tempBuffer); + status_t = cusparseSpSM_analysis(handle_t, opA, opB, alpha, spmatA, dnmatB, dnmatC, dataType, spSMAlg_t, spSMDescr, tempBuffer); +#endif +#endif + + return 0; +} diff --git a/tests/unit_tests/synthetic/libraries/cusparse2rocsparse_12000.cu b/tests/unit_tests/synthetic/libraries/cusparse2rocsparse_12000.cu index 79ea9c21..5039329c 100644 --- a/tests/unit_tests/synthetic/libraries/cusparse2rocsparse_12000.cu +++ b/tests/unit_tests/synthetic/libraries/cusparse2rocsparse_12000.cu @@ -107,6 +107,16 @@ int main() { cusparseDenseToSparseAlg_t denseToSparseAlg_t; #endif +#if CUDA_VERSION >= 11030 && CUSPARSE_VERSION >= 11600 + // CHECK: rocsparse_spsm_alg spSMAlg_t; + // CHECK-NEXT: rocsparse_spsm_alg SPSM_ALG_DEFAULT = rocsparse_spsm_alg_default; + cusparseSpSMAlg_t spSMAlg_t; + cusparseSpSMAlg_t SPSM_ALG_DEFAULT = CUSPARSE_SPSM_ALG_DEFAULT; + + // NOTE:cusparseSpSMDescr_t doesn't have a correspondence in rocSPARSE, the corresponding function argument is removed in the hipified call of the rocsparse_spsm function + cusparseSpSMDescr_t spSMDescr; +#endif + #if CUDA_VERSION >= 12000 // CHECK: rocsparse_const_spmat_descr constSpMatDescr = nullptr; // CHECK-NEXT: rocsparse_const_spmat_descr constSpMatDescrB = nullptr; @@ -142,6 +152,11 @@ int main() { // 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); + + // CUDA: cusparseStatus_t CUSPARSEAPI cusparseSpSM_analysis(cusparseHandle_t handle, cusparseOperation_t opA, cusparseOperation_t opB, const void* alpha, cusparseConstSpMatDescr_t matA, cusparseConstDnMatDescr_t matB, cusparseDnMatDescr_t matC, cudaDataType computeType, cusparseSpSMAlg_t alg, cusparseSpSMDescr_t spsmDescr, void* externalBuffer); + // ROC: ROCSPARSE_EXPORT rocsparse_status rocsparse_spsm(rocsparse_handle handle, rocsparse_operation trans_A, rocsparse_operation trans_B, const void* alpha, rocsparse_const_spmat_descr matA, rocsparse_const_dnmat_descr matB, const rocsparse_dnmat_descr matC, rocsparse_datatype compute_type, rocsparse_spsm_alg alg, rocsparse_spsm_stage stage, size_t* buffer_size, void* temp_buffer); + // CHECK: status_t = rocsparse_spsm(handle_t, opA, opB, alpha, constSpMatDescr, constDnMatDescrB, dnmatC, dataType, spSMAlg_t, rocsparse_spsm_stage_compute, nullptr, tempBuffer); + status_t = cusparseSpSM_analysis(handle_t, opA, opB, alpha, constSpMatDescr, constDnMatDescrB, dnmatC, dataType, spSMAlg_t, spSMDescr, tempBuffer); #endif return 0;