[HIPIFY][ROCm#674][rocSPARSE][feature] rocSPARSE support - Step 100 -…

… `cusparseDenseToSparse_bufferSize` -> `rocsparse_dense_to_sparse`

+ [IMP] `rocsparse_dense_to_sparse` 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

@@ -2438,6 +2438,7 @@ sub rocSubstitutions {
     subst("cusparseDdense2csc", "rocsparse_ddense2csc", "library");
     subst("cusparseDdense2csr", "rocsparse_ddense2csr", "library");
     subst("cusparseDdoti", "rocsparse_ddoti", "library");
+    subst("cusparseDenseToSparse_bufferSize", "rocsparse_dense_to_sparse", "library");
     subst("cusparseDestroy", "rocsparse_destroy_handle", "library");
     subst("cusparseDestroyBsric02Info", "rocsparse_destroy_mat_info", "library");
     subst("cusparseDestroyBsrilu02Info", "rocsparse_destroy_mat_info", "library");

docs/tables/CUSPARSE_API_supported_by_HIP_and_ROC.md

@@ -843,7 +843,7 @@
 |`cusparseCsrSetPointers`|11.0| | | |`hipsparseCsrSetPointers`|4.1.0| | | | |`rocsparse_csr_set_pointers`|4.1.0| | | | |
 |`cusparseCsrSetStridedBatch`|11.0| | | |`hipsparseCsrSetStridedBatch`|5.2.0| | | | |`rocsparse_csr_set_strided_batch`|5.2.0| | | | |
 |`cusparseDenseToSparse_analysis`|11.1| |12.0| |`hipsparseDenseToSparse_analysis`|4.2.0| |6.0.0| | | | | | | | |
-|`cusparseDenseToSparse_bufferSize`|11.1| |12.0| |`hipsparseDenseToSparse_bufferSize`|4.2.0| |6.0.0| | | | | | | | |
+|`cusparseDenseToSparse_bufferSize`|11.1| |12.0| |`hipsparseDenseToSparse_bufferSize`|4.2.0| |6.0.0| | |`rocsparse_dense_to_sparse`|4.1.0| |6.0.0| | |
 |`cusparseDenseToSparse_convert`|11.1| |12.0| |`hipsparseDenseToSparse_convert`|4.2.0| |6.0.0| | | | | | | | |
 |`cusparseDestroyDnMat`|10.1| |12.0| |`hipsparseDestroyDnMat`|4.2.0| |6.0.0| | |`rocsparse_destroy_dnmat_descr`|4.1.0| |6.0.0| | |
 |`cusparseDestroyDnVec`|10.2| |12.0| |`hipsparseDestroyDnVec`|4.1.0| |6.0.0| | |`rocsparse_destroy_dnvec_descr`|4.1.0| |6.0.0| | |

docs/tables/CUSPARSE_API_supported_by_ROC.md

@@ -843,7 +843,7 @@
 |`cusparseCsrSetPointers`|11.0| | | |`rocsparse_csr_set_pointers`|4.1.0| | | | |
 |`cusparseCsrSetStridedBatch`|11.0| | | |`rocsparse_csr_set_strided_batch`|5.2.0| | | | |
 |`cusparseDenseToSparse_analysis`|11.1| |12.0| | | | | | | |
-|`cusparseDenseToSparse_bufferSize`|11.1| |12.0| | | | | | | |
+|`cusparseDenseToSparse_bufferSize`|11.1| |12.0| |`rocsparse_dense_to_sparse`|4.1.0| |6.0.0| | |
 |`cusparseDenseToSparse_convert`|11.1| |12.0| | | | | | | |
 |`cusparseDestroyDnMat`|10.1| |12.0| |`rocsparse_destroy_dnmat_descr`|4.1.0| |6.0.0| | |
 |`cusparseDestroyDnVec`|10.2| |12.0| |`rocsparse_destroy_dnvec_descr`|4.1.0| |6.0.0| | |

src/CUDA2HIP_SPARSE_API_functions.cpp

@@ -865,7 +865,7 @@ const std::map<llvm::StringRef, hipCounter> CUDA_SPARSE_FUNCTION_MAP {
 
   {"cusparseSparseToDense",                             {"hipsparseSparseToDense",                             "rocsparse_sparse_to_dense",                                        CONV_LIB_FUNC, API_SPARSE, 15}},
   {"cusparseSparseToDense_bufferSize",                  {"hipsparseSparseToDense_bufferSize",                  "rocsparse_sparse_to_dense",                                        CONV_LIB_FUNC, API_SPARSE, 15}},
-  {"cusparseDenseToSparse_bufferSize",                  {"hipsparseDenseToSparse_bufferSize",                  "",                                                                 CONV_LIB_FUNC, API_SPARSE, 15, ROC_UNSUPPORTED}},
+  {"cusparseDenseToSparse_bufferSize",                  {"hipsparseDenseToSparse_bufferSize",                  "rocsparse_dense_to_sparse",                                        CONV_LIB_FUNC, API_SPARSE, 15}},
   {"cusparseDenseToSparse_analysis",                    {"hipsparseDenseToSparse_analysis",                    "",                                                                 CONV_LIB_FUNC, API_SPARSE, 15, ROC_UNSUPPORTED}},
   {"cusparseDenseToSparse_convert",                     {"hipsparseDenseToSparse_convert",                     "",                                                                 CONV_LIB_FUNC, API_SPARSE, 15, ROC_UNSUPPORTED}},
 
@@ -1224,7 +1224,7 @@ const std::map<llvm::StringRef, cudaAPIversions> CUDA_SPARSE_FUNCTION_VER_MAP {
   {"cusparseCooSetPointers",                            {CUDA_111, CUDA_0,   CUDA_0  }},
   {"cusparseSparseToDense_bufferSize",                  {CUDA_111, CUDA_0,   CUDA_0  }}, // A: CUSPARSE_VERSION 11300 C: CUSPARSE_VERSION 12000
   {"cusparseSparseToDense",                             {CUDA_111, CUDA_0,   CUDA_0  }}, // A: CUSPARSE_VERSION 11300 C: CUSPARSE_VERSION 12000
-  {"cusparseDenseToSparse_bufferSize",                  {CUDA_111, CUDA_0,   CUDA_0  }},
+  {"cusparseDenseToSparse_bufferSize",                  {CUDA_111, CUDA_0,   CUDA_0  }}, // A: CUSPARSE_VERSION 11300 C: CUSPARSE_VERSION 12000
   {"cusparseDenseToSparse_analysis",                    {CUDA_111, CUDA_0,   CUDA_0  }},
   {"cusparseDenseToSparse_convert",                     {CUDA_111, CUDA_0,   CUDA_0  }},
   {"cusparseCreateCsrsv2Info",                          {CUDA_0,   CUDA_113, CUDA_120}}, // D: CUSPARSE_VERSION 11500 R: CUSPARSE_VERSION 12000
@@ -2424,6 +2424,7 @@ const std::map<llvm::StringRef, hipAPIversions> HIP_SPARSE_FUNCTION_VER_MAP {
   {"rocsparse_zbsrilu0_buffer_size",                     {HIP_3080, HIP_0,    HIP_0   }},
   {"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   }},
 };
 
 const std::map<llvm::StringRef, cudaAPIChangedVersions> CUDA_SPARSE_FUNCTION_CHANGED_VER_MAP {
@@ -2527,6 +2528,7 @@ const std::map<llvm::StringRef, hipAPIChangedVersions> HIP_SPARSE_FUNCTION_CHANG
   {"rocsparse_destroy_dnmat_descr",                      {HIP_6000}},
   {"rocsparse_dnmat_get_strided_batch",                  {HIP_6000}},
   {"rocsparse_sparse_to_dense",                          {HIP_6000}},
+  {"rocsparse_dense_to_sparse",                          {HIP_6000}},
 };
 
 const std::map<unsigned int, llvm::StringRef> CUDA_SPARSE_API_SECTION_MAP {

src/HipifyAction.cpp

@@ -206,6 +206,7 @@ const std::string sCusparseSbsrilu02_bufferSize = "cusparseSbsrilu02_bufferSize"
 const std::string sCusparseCsr2cscEx2_bufferSize = "cusparseCsr2cscEx2_bufferSize";
 const std::string sCusparseSparseToDense = "cusparseSparseToDense";
 const std::string sCusparseSparseToDense_bufferSize = "cusparseSparseToDense_bufferSize";
+const std::string sCusparseDenseToSparse_bufferSize = "cusparseDenseToSparse_bufferSize";
 
 // CUDA_OVERLOADED
 const std::string sCudaEventCreate = "cudaEventCreate";
@@ -1651,6 +1652,15 @@ std::map<std::string, ArgCastStruct> FuncArgCasts {
       false
     }
   },
+  {sCusparseDenseToSparse_bufferSize,
+    {
+      {
+        {5, {e_add_const_argument, cw_None, "nullptr"}}
+      },
+      true,
+      false
+    }
+  },
 };
 
 void HipifyAction::RewriteString(StringRef s, clang::SourceLocation start) {
@@ -2497,7 +2507,8 @@ std::unique_ptr<clang::ASTConsumer> HipifyAction::CreateASTConsumer(clang::Compi
             sCusparseSbsrilu02_bufferSize,
             sCusparseCsr2cscEx2_bufferSize,
             sCusparseSparseToDense,
-            sCusparseSparseToDense_bufferSize
+            sCusparseSparseToDense_bufferSize,
+            sCusparseDenseToSparse_bufferSize
           )
         )
       )

tests/unit_tests/synthetic/libraries/cusparse2hipsparse.cu

@@ -2585,7 +2585,6 @@ int main() {
   // CHECK: status_t = hipsparseSparseToDense(handle_t, spmatA, dnmatB, sparseToDenseAlg_t, tempBuffer);
   status_t = cusparseSparseToDense(handle_t, spmatA, dnmatB, sparseToDenseAlg_t, tempBuffer);
 
-  // TODO: Mark as C-Changed in 12.0.0
   // CUDA: cusparseStatus_t CUSPARSEAPI cusparseDenseToSparse_bufferSize(cusparseHandle_t handle, cusparseDnMatDescr_t matA, cusparseSpMatDescr_t matB, cusparseDenseToSparseAlg_t alg, size_t* bufferSize);
   // HIP: HIPSPARSE_EXPORT hipsparseStatus_t hipsparseDenseToSparse_bufferSize(hipsparseHandle_t handle, hipsparseDnMatDescr_t matA, hipsparseSpMatDescr_t matB, hipsparseDenseToSparseAlg_t alg, size_t* bufferSize);
   // CHECK: status_t = hipsparseDenseToSparse_bufferSize(handle_t, dnmatA, spmatB, denseToSparseAlg_t, &bufferSize);

tests/unit_tests/synthetic/libraries/cusparse2rocsparse_12000.cu

@@ -80,11 +80,6 @@ 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 >= 11010 && CUSPARSE_VERSION >= 11300
-  // CHECK: rocsparse_sparse_to_dense_alg sparseToDenseAlg_t;
-  cusparseSparseToDenseAlg_t sparseToDenseAlg_t;
-#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;
@@ -93,6 +88,14 @@ int main() {
   cusparseDnMatDescr_t dnMatDescr_t, dnmatA, dnmatB, dnmatC;
 #endif
 
+#if CUDA_VERSION >= 11010 && CUSPARSE_VERSION >= 11300
+  // CHECK: rocsparse_sparse_to_dense_alg sparseToDenseAlg_t;
+  cusparseSparseToDenseAlg_t sparseToDenseAlg_t;
+
+  // CHECK: rocsparse_dense_to_sparse_alg denseToSparseAlg_t;
+  cusparseDenseToSparseAlg_t denseToSparseAlg_t;
+#endif
+
 #if CUDA_VERSION >= 12000
   // CHECK: rocsparse_const_spmat_descr constSpMatDescr = nullptr;
   // CHECK-NEXT: rocsparse_const_spmat_descr constSpMatDescrB = nullptr;
@@ -108,6 +111,11 @@ int main() {
   // ROC: ROCSPARSE_EXPORT rocsparse_status rocsparse_sparse_to_dense(rocsparse_handle handle, rocsparse_const_spmat_descr mat_A, rocsparse_dnmat_descr mat_B, rocsparse_sparse_to_dense_alg alg, size_t* buffer_size, void* temp_buffer);
   // CHECK: status_t = rocsparse_sparse_to_dense(handle_t, constSpMatDescr, dnmatB, sparseToDenseAlg_t, &bufferSize, nullptr);
   status_t = cusparseSparseToDense_bufferSize(handle_t, constSpMatDescr, dnmatB, sparseToDenseAlg_t, &bufferSize);
+
+  // CUDA: cusparseStatus_t CUSPARSEAPI cusparseDenseToSparse_bufferSize(cusparseHandle_t handle, cusparseConstDnMatDescr_t matA, cusparseSpMatDescr_t matB, cusparseDenseToSparseAlg_t alg, size_t* bufferSize);
+  // 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, dnmatB, spMatDescr_t, denseToSparseAlg_t, &bufferSize, nullptr);
+  status_t = cusparseDenseToSparse_bufferSize(handle_t, dnmatB, spMatDescr_t, denseToSparseAlg_t, &bufferSize);
 #endif
 
   return 0;

tests/unit_tests/synthetic/libraries/cusparse2rocsparse_before_12000.cu

@@ -80,17 +80,20 @@ 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 >= 11010 && CUSPARSE_VERSION >= 11300
-  // CHECK: rocsparse_sparse_to_dense_alg sparseToDenseAlg_t;
-  cusparseSparseToDenseAlg_t sparseToDenseAlg_t;
-#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 >= 11010 && CUSPARSE_VERSION >= 11300
+  // CHECK: rocsparse_sparse_to_dense_alg sparseToDenseAlg_t;
+  cusparseSparseToDenseAlg_t sparseToDenseAlg_t;
+
+  // CHECK: rocsparse_dense_to_sparse_alg denseToSparseAlg_t;
+  cusparseDenseToSparseAlg_t denseToSparseAlg_t;
 
 #if CUDA_VERSION < 12000
   // CUDA: cusparseStatus_t CUSPARSEAPI cusparseSparseToDense(cusparseHandle_t handle, cusparseSpMatDescr_t matA, cusparseDnMatDescr_t matB, cusparseSparseToDenseAlg_t alg, void* buffer);
@@ -102,6 +105,11 @@ int main() {
   // ROC: ROCSPARSE_EXPORT rocsparse_status rocsparse_sparse_to_dense(rocsparse_handle handle, const rocsparse_spmat_descr mat_A, rocsparse_dnmat_descr mat_B, rocsparse_sparse_to_dense_alg alg, size_t* buffer_size, void* temp_buffer);
   // CHECK: status_t = rocsparse_sparse_to_dense(handle_t, spmatA, dnmatB, sparseToDenseAlg_t, &bufferSize, nullptr);
   status_t = cusparseSparseToDense_bufferSize(handle_t, spmatA, dnmatB, sparseToDenseAlg_t, &bufferSize);
+
+  // CUDA: cusparseStatus_t CUSPARSEAPI cusparseDenseToSparse_bufferSize(cusparseHandle_t handle, cusparseDnMatDescr_t matA, cusparseSpMatDescr_t matB, cusparseDenseToSparseAlg_t alg, size_t* bufferSize);
+  // ROC: ROCSPARSE_EXPORT rocsparse_status rocsparse_dense_to_sparse(rocsparse_handle handle, const rocsparse_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, dnmatA, spmatB, denseToSparseAlg_t, &bufferSize, nullptr);
+  status_t = cusparseDenseToSparse_bufferSize(handle_t, dnmatA, spmatB, denseToSparseAlg_t, &bufferSize);
 #endif
 #endif