[HIPIFY][ROCm#674][ROCm#1014][rocSPARSE][feature] rocSPARSE support -…

… Step 39 - functions + Implemented a new function call transformation type "replace argument with a const" (`e_replace_argument_with_const`) + The new transformation is tested on rocSPARSE functions `rocsparse_(s|d|c|z)csrilu0`, where: - the penultimate argument `cusparseSolvePolicy_t policy` should always be the const value `rocsparse_solve_policy_auto`; - it is how `hipsparse(S|D|C|Z)csrilu02` calls `rocsparse_(s|d|c|z)csrilu0` in its implementation; + Updated synthetic tests and the regenerated hipify-perl and SPARSE docs + Added the missing `csrilu02Info_t` -> `rocsparse_mat_info` transformation
emankov · Sep 10, 2023 · 472483d · 472483d
1 parent 880b76c
commit 472483d
Show file tree

Hide file tree

Showing 9 changed files with 130 additions and 22 deletions.
diff --git a/bin/hipify-perl b/bin/hipify-perl
@@ -1682,6 +1682,7 @@ sub rocSubstitutions {
     subst("cusparseCcsr2gebsr_bufferSize", "rocsparse_ccsr2gebsr_buffer_size", "library");
     subst("cusparseCcsr2hyb", "rocsparse_ccsr2hyb", "library");
     subst("cusparseCcsrcolor", "rocsparse_ccsrcolor", "library");
+    subst("cusparseCcsrilu02", "rocsparse_ccsrilu0", "library");
     subst("cusparseCdense2csc", "rocsparse_cdense2csc", "library");
     subst("cusparseCdense2csr", "rocsparse_cdense2csr", "library");
     subst("cusparseCgebsr2csr", "rocsparse_cgebsr2csr", "library");
@@ -1732,6 +1733,7 @@ sub rocSubstitutions {
     subst("cusparseDcsr2gebsr_bufferSize", "rocsparse_dcsr2gebsr_buffer_size", "library");
     subst("cusparseDcsr2hyb", "rocsparse_dcsr2hyb", "library");
     subst("cusparseDcsrcolor", "rocsparse_dcsrcolor", "library");
+    subst("cusparseDcsrilu02", "rocsparse_dcsrilu0", "library");
     subst("cusparseDdense2csc", "rocsparse_ddense2csc", "library");
     subst("cusparseDdense2csr", "rocsparse_ddense2csr", "library");
     subst("cusparseDestroy", "rocsparse_destroy_handle", "library");
@@ -1801,6 +1803,7 @@ sub rocSubstitutions {
     subst("cusparseScsr2gebsr_bufferSize", "rocsparse_scsr2gebsr_buffer_size", "library");
     subst("cusparseScsr2hyb", "rocsparse_scsr2hyb", "library");
     subst("cusparseScsrcolor", "rocsparse_scsrcolor", "library");
+    subst("cusparseScsrilu02", "rocsparse_scsrilu0", "library");
     subst("cusparseSdense2csc", "rocsparse_sdense2csc", "library");
     subst("cusparseSdense2csr", "rocsparse_sdense2csr", "library");
     subst("cusparseSetMatDiagType", "rocsparse_set_mat_diag_type", "library");
@@ -1873,6 +1876,7 @@ sub rocSubstitutions {
     subst("cusparseZcsr2gebsr_bufferSize", "rocsparse_zcsr2gebsr_buffer_size", "library");
     subst("cusparseZcsr2hyb", "rocsparse_zcsr2hyb", "library");
     subst("cusparseZcsrcolor", "rocsparse_zcsrcolor", "library");
+    subst("cusparseZcsrilu02", "rocsparse_zcsrilu0", "library");
     subst("cusparseZdense2csc", "rocsparse_zdense2csc", "library");
     subst("cusparseZdense2csr", "rocsparse_zdense2csr", "library");
     subst("cusparseZgebsr2csr", "rocsparse_zgebsr2csr", "library");
@@ -1894,6 +1898,8 @@ sub rocSubstitutions {
     subst("cusparseZnnz_compress", "rocsparse_znnz_compress", "library");
     subst("cublas.h", "rocblas.h", "include_cuda_main_header");
     subst("cublas_v2.h", "rocblas.h", "include_cuda_main_header_v2");
+    subst("csrilu02Info", "_rocsparse_mat_info", "type");
+    subst("csrilu02Info_t", "rocsparse_mat_info", "type");
     subst("cuComplex", "rocblas_float_complex", "type");
     subst("cuDoubleComplex", "rocblas_double_complex", "type");
     subst("cuFloatComplex", "rocblas_float_complex", "type");

diff --git a/docs/tables/CUSPARSE_API_supported_by_HIP_and_ROC.md b/docs/tables/CUSPARSE_API_supported_by_HIP_and_ROC.md
@@ -120,8 +120,8 @@
 |`csrgemm2Info_t`| | |12.0|`csrgemm2Info_t`|2.8.0| | | | | | | | |
 |`csric02Info`| |12.2| |`csric02Info`|3.1.0| | | | | | | | |
 |`csric02Info_t`| |12.2| |`csric02Info_t`|3.1.0| | | | | | | | |
-|`csrilu02Info`| |12.2| |`csrilu02Info`|1.9.2| | | | | | | | |
-|`csrilu02Info_t`| |12.2| |`csrilu02Info_t`|1.9.2| | | | | | | | |
+|`csrilu02Info`| |12.2| |`csrilu02Info`|1.9.2| | | |`_rocsparse_mat_info`|1.9.0| | | |
+|`csrilu02Info_t`| |12.2| |`csrilu02Info_t`|1.9.2| | | |`rocsparse_mat_info`|1.9.0| | | |
 |`csrsm2Info`|9.2| |12.0| | | | | | | | | | |
 |`csrsm2Info_t`|9.2| |12.0|`csrsm2Info_t`|3.1.0| | | | | | | | |
 |`csrsv2Info`| | |12.0| | | | | | | | | | |
@@ -484,7 +484,7 @@
 |`cusparseCcsric02_bufferSize`| |12.2| |`hipsparseCcsric02_bufferSize`|3.1.0| | | | | | | | |
 |`cusparseCcsric02_bufferSizeExt`| |12.2| |`hipsparseCcsric02_bufferSizeExt`|3.1.0| | | | | | | | |
 |`cusparseCcsrilu0`| |10.2|11.0| | | | | | | | | | |
-|`cusparseCcsrilu02`| |12.2| |`hipsparseCcsrilu02`|3.1.0| | | | | | | | |
+|`cusparseCcsrilu02`| |12.2| |`hipsparseCcsrilu02`|3.1.0| | | |`rocsparse_ccsrilu0`|2.10.0| | | |
 |`cusparseCcsrilu02_analysis`| |12.2| |`hipsparseCcsrilu02_analysis`|3.1.0| | | | | | | | |
 |`cusparseCcsrilu02_bufferSize`| |12.2| |`hipsparseCcsrilu02_bufferSize`|3.1.0| | | | | | | | |
 |`cusparseCcsrilu02_bufferSizeExt`| |12.2| |`hipsparseCcsrilu02_bufferSizeExt`|3.1.0| | | | | | | | |
@@ -518,7 +518,7 @@
 |`cusparseDcsric02_bufferSize`| |12.2| |`hipsparseDcsric02_bufferSize`|3.1.0| | | | | | | | |
 |`cusparseDcsric02_bufferSizeExt`| |12.2| |`hipsparseDcsric02_bufferSizeExt`|3.1.0| | | | | | | | |
 |`cusparseDcsrilu0`| |10.2|11.0| | | | | | | | | | |
-|`cusparseDcsrilu02`| |12.2| |`hipsparseDcsrilu02`|1.9.2| | | | | | | | |
+|`cusparseDcsrilu02`| |12.2| |`hipsparseDcsrilu02`|1.9.2| | | |`rocsparse_dcsrilu0`|1.9.0| | | |
 |`cusparseDcsrilu02_analysis`| |12.2| |`hipsparseDcsrilu02_analysis`|1.9.2| | | | | | | | |
 |`cusparseDcsrilu02_bufferSize`| |12.2| |`hipsparseDcsrilu02_bufferSize`|1.9.2| | | | | | | | |
 |`cusparseDcsrilu02_bufferSizeExt`| |12.2| |`hipsparseDcsrilu02_bufferSizeExt`|1.9.2| | | | | | | | |
@@ -551,7 +551,7 @@
 |`cusparseScsric02_bufferSize`| |12.2| |`hipsparseScsric02_bufferSize`|3.1.0| | | | | | | | |
 |`cusparseScsric02_bufferSizeExt`| |12.2| |`hipsparseScsric02_bufferSizeExt`|3.1.0| | | | | | | | |
 |`cusparseScsrilu0`| |10.2|11.0| | | | | | | | | | |
-|`cusparseScsrilu02`| |12.2| |`hipsparseScsrilu02`|1.9.2| | | | | | | | |
+|`cusparseScsrilu02`| |12.2| |`hipsparseScsrilu02`|1.9.2| | | |`rocsparse_scsrilu0`|1.9.0| | | |
 |`cusparseScsrilu02_analysis`| |12.2| |`hipsparseScsrilu02_analysis`|1.9.2| | | | | | | | |
 |`cusparseScsrilu02_bufferSize`| |12.2| |`hipsparseScsrilu02_bufferSize`|1.9.2| | | | | | | | |
 |`cusparseScsrilu02_bufferSizeExt`| |12.2| |`hipsparseScsrilu02_bufferSizeExt`|1.9.2| | | | | | | | |
@@ -588,7 +588,7 @@
 |`cusparseZcsric02_bufferSize`| |12.2| |`hipsparseZcsric02_bufferSize`|3.1.0| | | | | | | | |
 |`cusparseZcsric02_bufferSizeExt`| |12.2| |`hipsparseZcsric02_bufferSizeExt`|3.1.0| | | | | | | | |
 |`cusparseZcsrilu0`| |10.2|11.0| | | | | | | | | | |
-|`cusparseZcsrilu02`| |12.2| |`hipsparseZcsrilu02`|3.1.0| | | | | | | | |
+|`cusparseZcsrilu02`| |12.2| |`hipsparseZcsrilu02`|3.1.0| | | |`rocsparse_zcsrilu0`|2.10.0| | | |
 |`cusparseZcsrilu02_analysis`| |12.2| |`hipsparseZcsrilu02_analysis`|3.1.0| | | | | | | | |
 |`cusparseZcsrilu02_bufferSize`| |12.2| |`hipsparseZcsrilu02_bufferSize`|3.1.0| | | | | | | | |
 |`cusparseZcsrilu02_bufferSizeExt`| |12.2| |`hipsparseZcsrilu02_bufferSizeExt`|3.1.0| | | | | | | | |

diff --git a/docs/tables/CUSPARSE_API_supported_by_ROC.md b/docs/tables/CUSPARSE_API_supported_by_ROC.md
@@ -120,8 +120,8 @@
 |`csrgemm2Info_t`| | |12.0| | | | | |
 |`csric02Info`| |12.2| | | | | | |
 |`csric02Info_t`| |12.2| | | | | | |
-|`csrilu02Info`| |12.2| | | | | | |
-|`csrilu02Info_t`| |12.2| | | | | | |
+|`csrilu02Info`| |12.2| |`_rocsparse_mat_info`|1.9.0| | | |
+|`csrilu02Info_t`| |12.2| |`rocsparse_mat_info`|1.9.0| | | |
 |`csrsm2Info`|9.2| |12.0| | | | | |
 |`csrsm2Info_t`|9.2| |12.0| | | | | |
 |`csrsv2Info`| | |12.0| | | | | |
@@ -484,7 +484,7 @@
 |`cusparseCcsric02_bufferSize`| |12.2| | | | | | |
 |`cusparseCcsric02_bufferSizeExt`| |12.2| | | | | | |
 |`cusparseCcsrilu0`| |10.2|11.0| | | | | |
-|`cusparseCcsrilu02`| |12.2| | | | | | |
+|`cusparseCcsrilu02`| |12.2| |`rocsparse_ccsrilu0`|2.10.0| | | |
 |`cusparseCcsrilu02_analysis`| |12.2| | | | | | |
 |`cusparseCcsrilu02_bufferSize`| |12.2| | | | | | |
 |`cusparseCcsrilu02_bufferSizeExt`| |12.2| | | | | | |
@@ -518,7 +518,7 @@
 |`cusparseDcsric02_bufferSize`| |12.2| | | | | | |
 |`cusparseDcsric02_bufferSizeExt`| |12.2| | | | | | |
 |`cusparseDcsrilu0`| |10.2|11.0| | | | | |
-|`cusparseDcsrilu02`| |12.2| | | | | | |
+|`cusparseDcsrilu02`| |12.2| |`rocsparse_dcsrilu0`|1.9.0| | | |
 |`cusparseDcsrilu02_analysis`| |12.2| | | | | | |
 |`cusparseDcsrilu02_bufferSize`| |12.2| | | | | | |
 |`cusparseDcsrilu02_bufferSizeExt`| |12.2| | | | | | |
@@ -551,7 +551,7 @@
 |`cusparseScsric02_bufferSize`| |12.2| | | | | | |
 |`cusparseScsric02_bufferSizeExt`| |12.2| | | | | | |
 |`cusparseScsrilu0`| |10.2|11.0| | | | | |
-|`cusparseScsrilu02`| |12.2| | | | | | |
+|`cusparseScsrilu02`| |12.2| |`rocsparse_scsrilu0`|1.9.0| | | |
 |`cusparseScsrilu02_analysis`| |12.2| | | | | | |
 |`cusparseScsrilu02_bufferSize`| |12.2| | | | | | |
 |`cusparseScsrilu02_bufferSizeExt`| |12.2| | | | | | |
@@ -588,7 +588,7 @@
 |`cusparseZcsric02_bufferSize`| |12.2| | | | | | |
 |`cusparseZcsric02_bufferSizeExt`| |12.2| | | | | | |
 |`cusparseZcsrilu0`| |10.2|11.0| | | | | |
-|`cusparseZcsrilu02`| |12.2| | | | | | |
+|`cusparseZcsrilu02`| |12.2| |`rocsparse_zcsrilu0`|2.10.0| | | |
 |`cusparseZcsrilu02_analysis`| |12.2| | | | | | |
 |`cusparseZcsrilu02_bufferSize`| |12.2| | | | | | |
 |`cusparseZcsrilu02_bufferSizeExt`| |12.2| | | | | | |

diff --git a/src/CUDA2HIP_SPARSE_API_functions.cpp b/src/CUDA2HIP_SPARSE_API_functions.cpp
@@ -391,10 +391,10 @@ const std::map<llvm::StringRef, hipCounter> CUDA_SPARSE_FUNCTION_MAP {
   {"cusparseCcsrilu02_analysis",                        {"hipsparseCcsrilu02_analysis",                        "",                                                                 CONV_LIB_FUNC, API_SPARSE, 12, ROC_UNSUPPORTED | CUDA_DEPRECATED}},
   {"cusparseZcsrilu02_analysis",                        {"hipsparseZcsrilu02_analysis",                        "",                                                                 CONV_LIB_FUNC, API_SPARSE, 12, ROC_UNSUPPORTED | CUDA_DEPRECATED}},
 
-  {"cusparseScsrilu02",                                 {"hipsparseScsrilu02",                                 "",                                                                 CONV_LIB_FUNC, API_SPARSE, 12, ROC_UNSUPPORTED | CUDA_DEPRECATED}},
-  {"cusparseDcsrilu02",                                 {"hipsparseDcsrilu02",                                 "",                                                                 CONV_LIB_FUNC, API_SPARSE, 12, ROC_UNSUPPORTED | CUDA_DEPRECATED}},
-  {"cusparseCcsrilu02",                                 {"hipsparseCcsrilu02",                                 "",                                                                 CONV_LIB_FUNC, API_SPARSE, 12, ROC_UNSUPPORTED | CUDA_DEPRECATED}},
-  {"cusparseZcsrilu02",                                 {"hipsparseZcsrilu02",                                 "",                                                                 CONV_LIB_FUNC, API_SPARSE, 12, ROC_UNSUPPORTED | CUDA_DEPRECATED}},
+  {"cusparseScsrilu02",                                 {"hipsparseScsrilu02",                                 "rocsparse_scsrilu0",                                               CONV_LIB_FUNC, API_SPARSE, 12, CUDA_DEPRECATED}},
+  {"cusparseDcsrilu02",                                 {"hipsparseDcsrilu02",                                 "rocsparse_dcsrilu0",                                               CONV_LIB_FUNC, API_SPARSE, 12, CUDA_DEPRECATED}},
+  {"cusparseCcsrilu02",                                 {"hipsparseCcsrilu02",                                 "rocsparse_ccsrilu0",                                               CONV_LIB_FUNC, API_SPARSE, 12, CUDA_DEPRECATED}},
+  {"cusparseZcsrilu02",                                 {"hipsparseZcsrilu02",                                 "rocsparse_zcsrilu0",                                               CONV_LIB_FUNC, API_SPARSE, 12, CUDA_DEPRECATED}},
   {"cusparseXcsrilu02_zeroPivot",                       {"hipsparseXcsrilu02_zeroPivot",                       "",                                                                 CONV_LIB_FUNC, API_SPARSE, 12, ROC_UNSUPPORTED | CUDA_DEPRECATED}},
 
   {"cusparseSbsrilu02_numericBoost",                    {"hipsparseSbsrilu02_numericBoost",                    "",                                                                 CONV_LIB_FUNC, API_SPARSE, 12, ROC_UNSUPPORTED | CUDA_DEPRECATED}},
@@ -2175,6 +2175,10 @@ const std::map<llvm::StringRef, hipAPIversions> HIP_SPARSE_FUNCTION_VER_MAP {
   {"rocsparse_cgtsv_buffer_size",                        {HIP_4030, HIP_0,    HIP_0   }},
   {"rocsparse_dgtsv_buffer_size",                        {HIP_4030, HIP_0,    HIP_0   }},
   {"rocsparse_sgtsv_buffer_size",                        {HIP_4030, HIP_0,    HIP_0   }},
+  {"rocsparse_zcsrilu0",                                 {HIP_2100, HIP_0,    HIP_0   }},
+  {"rocsparse_ccsrilu0",                                 {HIP_2100, HIP_0,    HIP_0   }},
+  {"rocsparse_dcsrilu0",                                 {HIP_1090, HIP_0,    HIP_0   }},
+  {"rocsparse_scsrilu0",                                 {HIP_1090, HIP_0,    HIP_0   }},
 };
 
 const std::map<unsigned int, llvm::StringRef> CUDA_SPARSE_API_SECTION_MAP {

diff --git a/src/CUDA2HIP_SPARSE_API_types.cpp b/src/CUDA2HIP_SPARSE_API_types.cpp
@@ -53,8 +53,8 @@ const std::map<llvm::StringRef, hipCounter> CUDA_SPARSE_TYPE_NAME_MAP {
   {"bsric02Info",                               {"bsric02Info",                                "",                                                   CONV_TYPE, API_SPARSE, 4, ROC_UNSUPPORTED | CUDA_DEPRECATED}},
   {"bsric02Info_t",                             {"bsric02Info_t",                              "",                                                   CONV_TYPE, API_SPARSE, 4, ROC_UNSUPPORTED | CUDA_DEPRECATED}},
 
-  {"csrilu02Info",                              {"csrilu02Info",                               "",                                                   CONV_TYPE, API_SPARSE, 4, ROC_UNSUPPORTED | CUDA_DEPRECATED}},
-  {"csrilu02Info_t",                            {"csrilu02Info_t",                             "",                                                   CONV_TYPE, API_SPARSE, 4, ROC_UNSUPPORTED | CUDA_DEPRECATED}},
+  {"csrilu02Info",                              {"csrilu02Info",                               "_rocsparse_mat_info",                                CONV_TYPE, API_SPARSE, 4, CUDA_DEPRECATED}},
+  {"csrilu02Info_t",                            {"csrilu02Info_t",                             "rocsparse_mat_info",                                 CONV_TYPE, API_SPARSE, 4, CUDA_DEPRECATED}},
 
   {"bsrilu02Info",                              {"bsrilu02Info",                               "",                                                   CONV_TYPE, API_SPARSE, 4, ROC_UNSUPPORTED | CUDA_DEPRECATED}},
   {"bsrilu02Info_t",                            {"bsrilu02Info_t",                             "",                                                   CONV_TYPE, API_SPARSE, 4, ROC_UNSUPPORTED | CUDA_DEPRECATED}},

diff --git a/src/CUDA2HIP_Scripting.h b/src/CUDA2HIP_Scripting.h
@@ -33,6 +33,7 @@ namespace hipify {
     e_add_const_argument,
     e_add_var_argument,
     e_move_argument,
+    e_replace_argument_with_const,
   };
 
   enum OverloadTypes {
@@ -51,7 +52,7 @@ namespace hipify {
   struct CastInfo {
     CastTypes castType;
     CastWarning castWarn;
-    std::string constValToAdd = "";
+    std::string constValToAddOrReplace = "";
     unsigned moveOrCopyTo = 0;
     unsigned numberToMoveOrCopy = 1;
   };

diff --git a/src/HipifyAction.cpp b/src/HipifyAction.cpp
@@ -98,6 +98,10 @@ const std::string sCusparseZgtsvInterleavedBatch_bufferSizeExt = "cusparseZgtsvI
 const std::string sCusparseCgtsvInterleavedBatch_bufferSizeExt = "cusparseCgtsvInterleavedBatch_bufferSizeExt";
 const std::string sCusparseDgtsvInterleavedBatch_bufferSizeExt = "cusparseDgtsvInterleavedBatch_bufferSizeExt";
 const std::string sCusparseSgtsvInterleavedBatch_bufferSizeExt = "cusparseSgtsvInterleavedBatch_bufferSizeExt";
+const std::string sCusparseZcsrilu02 = "cusparseZcsrilu02";
+const std::string sCusparseCcsrilu02 = "cusparseCcsrilu02";
+const std::string sCusparseDcsrilu02 = "cusparseDcsrilu02";
+const std::string sCusparseScsrilu02 = "cusparseScsrilu02";
 // CUDA_OVERLOADED
 const std::string sCudaEventCreate = "cudaEventCreate";
 const std::string sCudaGraphInstantiate = "cudaGraphInstantiate";
@@ -122,6 +126,7 @@ std::string getCastType(hipify::CastTypes c) {
     case e_add_const_argument: return "";
     case e_add_var_argument: return "";
     case e_move_argument: return "";
+    case e_replace_argument_with_const: return "";
     default: return "";
   }
 }
@@ -542,6 +547,42 @@ std::map<std::string, ArgCastStruct> FuncArgCasts {
       false
     }
   },
+  {sCusparseZcsrilu02,
+    {
+      {
+        {8, {e_replace_argument_with_const, cw_None, "rocsparse_solve_policy_auto"}}
+      },
+      true,
+      false
+    }
+  },
+  {sCusparseCcsrilu02,
+    {
+      {
+        {8, {e_replace_argument_with_const, cw_None, "rocsparse_solve_policy_auto"}}
+      },
+      true,
+      false
+    }
+  },
+  {sCusparseDcsrilu02,
+    {
+      {
+        {8, {e_replace_argument_with_const, cw_None, "rocsparse_solve_policy_auto"}}
+      },
+      true,
+      false
+    }
+  },
+  {sCusparseScsrilu02,
+    {
+      {
+        {8, {e_replace_argument_with_const, cw_None, "rocsparse_solve_policy_auto"}}
+      },
+      true,
+      false
+    }
+  },
 };
 
 void HipifyAction::RewriteString(StringRef s, clang::SourceLocation start) {
@@ -1048,9 +1089,9 @@ bool HipifyAction::cudaHostFuncCall(const mat::MatchFinder::MatchResult &Result)
         case e_add_const_argument:
         {
           if (argNum < call->getNumArgs())
-            OS << c.second.constValToAdd << ", ";
+            OS << c.second.constValToAddOrReplace << ", ";
           else
-            OS << ", " << c.second.constValToAdd;
+            OS << ", " << c.second.constValToAddOrReplace;
           break;
         }
         case e_add_var_argument:
@@ -1070,6 +1111,14 @@ bool HipifyAction::cudaHostFuncCall(const mat::MatchFinder::MatchResult &Result)
           }
           break;
         }
+        case e_replace_argument_with_const:
+        {
+          if (argNum >= call->getNumArgs())
+            break;
+          OS << c.second.constValToAddOrReplace;
+          length = SM->getCharacterData(clang::Lexer::getLocForEndOfToken(e, 0, *SM, DefaultLangOptions)) - SM->getCharacterData(s);
+          break;
+        }
         default:
           OS << getCastType(c.second.castType) << "(" << readSourceText(*SM, sr) << ")";
           length = SM->getCharacterData(clang::Lexer::getLocForEndOfToken(e, 0, *SM, DefaultLangOptions)) - SM->getCharacterData(s);
@@ -1259,7 +1308,11 @@ std::unique_ptr<clang::ASTConsumer> HipifyAction::CreateASTConsumer(clang::Compi
             sCusparseZgtsvInterleavedBatch_bufferSizeExt,
             sCusparseCgtsvInterleavedBatch_bufferSizeExt,
             sCusparseDgtsvInterleavedBatch_bufferSizeExt,
-            sCusparseSgtsvInterleavedBatch_bufferSizeExt
+            sCusparseSgtsvInterleavedBatch_bufferSizeExt,
+            sCusparseZcsrilu02,
+            sCusparseCcsrilu02,
+            sCusparseDcsrilu02,
+            sCusparseScsrilu02
           )
         )
       )