Merge pull request #1161 from emankov/HIPIFY

[HIPIFY][#674][rocSPARSE][feature] rocSPARSE support - Step 64 - functions `rocsparse_(s|d|c|z)csrsv_solve`
ROCm · Nov 23, 2023 · 8de08bc · 8de08bc
2 parents 3b15203 + d0622a1
commit 8de08bc
Show file tree

Hide file tree

Showing 9 changed files with 184 additions and 21 deletions.
diff --git a/bin/hipify-perl b/bin/hipify-perl
@@ -1812,6 +1812,7 @@ sub rocSubstitutions {
     subst("cusparseCcsrsm2_analysis", "rocsparse_ccsrsm_analysis", "library");
     subst("cusparseCcsrsm2_bufferSizeExt", "rocsparse_ccsrsm_buffer_size", "library");
     subst("cusparseCcsrsm2_solve", "rocsparse_ccsrsm_solve", "library");
+    subst("cusparseCcsrsv2_solve", "rocsparse_ccsrsv_solve", "library");
     subst("cusparseCdense2csc", "rocsparse_cdense2csc", "library");
     subst("cusparseCdense2csr", "rocsparse_cdense2csr", "library");
     subst("cusparseCgebsr2csr", "rocsparse_cgebsr2csr", "library");
@@ -1889,6 +1890,7 @@ sub rocSubstitutions {
     subst("cusparseDcsrsm2_analysis", "rocsparse_dcsrsm_analysis", "library");
     subst("cusparseDcsrsm2_bufferSizeExt", "rocsparse_dcsrsm_buffer_size", "library");
     subst("cusparseDcsrsm2_solve", "rocsparse_dcsrsm_solve", "library");
+    subst("cusparseDcsrsv2_solve", "rocsparse_dcsrsv_solve", "library");
     subst("cusparseDdense2csc", "rocsparse_ddense2csc", "library");
     subst("cusparseDdense2csr", "rocsparse_ddense2csr", "library");
     subst("cusparseDestroy", "rocsparse_destroy_handle", "library");
@@ -1985,6 +1987,7 @@ sub rocSubstitutions {
     subst("cusparseScsrsm2_analysis", "rocsparse_scsrsm_analysis", "library");
     subst("cusparseScsrsm2_bufferSizeExt", "rocsparse_scsrsm_buffer_size", "library");
     subst("cusparseScsrsm2_solve", "rocsparse_scsrsm_solve", "library");
+    subst("cusparseScsrsv2_solve", "rocsparse_scsrsv_solve", "library");
     subst("cusparseSdense2csc", "rocsparse_sdense2csc", "library");
     subst("cusparseSdense2csr", "rocsparse_sdense2csr", "library");
     subst("cusparseSetMatDiagType", "rocsparse_set_mat_diag_type", "library");
@@ -2092,6 +2095,7 @@ sub rocSubstitutions {
     subst("cusparseZcsrsm2_analysis", "rocsparse_zcsrsm_analysis", "library");
     subst("cusparseZcsrsm2_bufferSizeExt", "rocsparse_zcsrsm_buffer_size", "library");
     subst("cusparseZcsrsm2_solve", "rocsparse_zcsrsm_solve", "library");
+    subst("cusparseZcsrsv2_solve", "rocsparse_zcsrsv_solve", "library");
     subst("cusparseZdense2csc", "rocsparse_zdense2csc", "library");
     subst("cusparseZdense2csr", "rocsparse_zdense2csr", "library");
     subst("cusparseZgebsr2csr", "rocsparse_zgebsr2csr", "library");
@@ -2132,6 +2136,8 @@ sub rocSubstitutions {
     subst("csrilu02Info_t", "rocsparse_mat_info", "type");
     subst("csrsm2Info", "_rocsparse_mat_info", "type");
     subst("csrsm2Info_t", "rocsparse_mat_info", "type");
+    subst("csrsv2Info", "_rocsparse_mat_descr", "type");
+    subst("csrsv2Info_t", "rocsparse_mat_descr", "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
@@ -124,8 +124,8 @@
 |`csrilu02Info_t`| |12.2| | |`csrilu02Info_t`|1.9.2| | | | |`rocsparse_mat_info`|1.9.0| | | | |
 |`csrsm2Info`|9.2| | |12.0| | | | | | |`_rocsparse_mat_info`|1.9.0| | | | |
 |`csrsm2Info_t`|9.2| | |12.0|`csrsm2Info_t`|3.1.0| | | | |`rocsparse_mat_info`|1.9.0| | | | |
-|`csrsv2Info`| | | |12.0| | | | | | | | | | | | |
-|`csrsv2Info_t`| | | |12.0|`csrsv2Info_t`|1.9.2| | | | | | | | | | |
+|`csrsv2Info`| | | |12.0| | | | | | |`_rocsparse_mat_descr`|1.9.0| | | | |
+|`csrsv2Info_t`| | | |12.0|`csrsv2Info_t`|1.9.2| | | | |`rocsparse_mat_descr`|1.9.0| | | | |
 |`csru2csrInfo`| |12.2| | |`csru2csrInfo`|4.2.0| | | | | | | | | | |
 |`csru2csrInfo_t`| |12.2| | |`csru2csrInfo_t`|4.2.0| | | | | | | | | | |
 |`cusparseAction_t`| | | | |`hipsparseAction_t`|1.9.2| | | | |`rocsparse_action`|1.9.0| | | | |
@@ -303,7 +303,7 @@
 |`cusparseCcsrsv2_analysis`| |11.3| |12.0|`hipsparseCcsrsv2_analysis`|3.1.0| | | | | | | | | | |
 |`cusparseCcsrsv2_bufferSize`| |11.3| |12.0|`hipsparseCcsrsv2_bufferSize`|3.1.0| | | | | | | | | | |
 |`cusparseCcsrsv2_bufferSizeExt`| |11.3| |12.0|`hipsparseCcsrsv2_bufferSizeExt`|3.1.0| | | | | | | | | | |
-|`cusparseCcsrsv2_solve`| |11.3| |12.0|`hipsparseCcsrsv2_solve`|3.1.0| | | | | | | | | | |
+|`cusparseCcsrsv2_solve`| |11.3| |12.0|`hipsparseCcsrsv2_solve`|3.1.0| | | | |`rocsparse_ccsrsv_solve`|2.10.0| | | | |
 |`cusparseCcsrsv_analysis`| |10.2| |11.0| | | | | | | | | | | | |
 |`cusparseCcsrsv_solve`| |10.2| |11.0| | | | | | | | | | | | |
 |`cusparseCgemvi`|7.5| | | |`hipsparseCgemvi`|4.3.0| | | | |`rocsparse_cgemvi`|4.3.0| | | | |
@@ -326,7 +326,7 @@
 |`cusparseDcsrsv2_analysis`| |11.3| |12.0|`hipsparseDcsrsv2_analysis`|1.9.2| | | | | | | | | | |
 |`cusparseDcsrsv2_bufferSize`| |11.3| |12.0|`hipsparseDcsrsv2_bufferSize`|1.9.2| | | | | | | | | | |
 |`cusparseDcsrsv2_bufferSizeExt`| |11.3| |12.0|`hipsparseDcsrsv2_bufferSizeExt`|1.9.2| | | | | | | | | | |
-|`cusparseDcsrsv2_solve`| |11.3| |12.0|`hipsparseDcsrsv2_solve`|1.9.2| | | | | | | | | | |
+|`cusparseDcsrsv2_solve`| |11.3| |12.0|`hipsparseDcsrsv2_solve`|1.9.2| | | | |`rocsparse_dcsrsv_solve`|1.9.0| | | | |
 |`cusparseDcsrsv_analysis`| |10.2| |11.0| | | | | | | | | | | | |
 |`cusparseDcsrsv_solve`| |10.2| |11.0| | | | | | | | | | | | |
 |`cusparseDgemvi`|7.5| | | |`hipsparseDgemvi`|4.3.0| | | | |`rocsparse_dgemvi`|4.3.0| | | | |
@@ -345,7 +345,7 @@
 |`cusparseScsrsv2_analysis`| |11.3| |12.0|`hipsparseScsrsv2_analysis`|1.9.2| | | | | | | | | | |
 |`cusparseScsrsv2_bufferSize`| |11.3| |12.0|`hipsparseScsrsv2_bufferSize`|1.9.2| | | | | | | | | | |
 |`cusparseScsrsv2_bufferSizeExt`| |11.3| |12.0|`hipsparseScsrsv2_bufferSizeExt`|1.9.2| | | | | | | | | | |
-|`cusparseScsrsv2_solve`| |11.3| |12.0|`hipsparseScsrsv2_solve`|1.9.2| | | | | | | | | | |
+|`cusparseScsrsv2_solve`| |11.3| |12.0|`hipsparseScsrsv2_solve`|1.9.2| | | | |`rocsparse_scsrsv_solve`|1.9.0| | | | |
 |`cusparseScsrsv_analysis`| |10.2| |11.0| | | | | | | | | | | | |
 |`cusparseScsrsv_solve`| |10.2| |11.0| | | | | | | | | | | | |
 |`cusparseSgemvi`|7.5| | | |`hipsparseSgemvi`|4.3.0| | | | |`rocsparse_sgemvi`|4.3.0| | | | |
@@ -366,7 +366,7 @@
 |`cusparseZcsrsv2_analysis`| |11.3| |12.0|`hipsparseZcsrsv2_analysis`|3.1.0| | | | | | | | | | |
 |`cusparseZcsrsv2_bufferSize`| |11.3| |12.0|`hipsparseZcsrsv2_bufferSize`|3.1.0| | | | | | | | | | |
 |`cusparseZcsrsv2_bufferSizeExt`| |11.3| |12.0|`hipsparseZcsrsv2_bufferSizeExt`|3.1.0| | | | | | | | | | |
-|`cusparseZcsrsv2_solve`| |11.3| |12.0|`hipsparseZcsrsv2_solve`|3.1.0| | | | | | | | | | |
+|`cusparseZcsrsv2_solve`| |11.3| |12.0|`hipsparseZcsrsv2_solve`|3.1.0| | | | |`rocsparse_zcsrsv_solve`|2.10.0| | | | |
 |`cusparseZcsrsv_analysis`| |10.2| |11.0| | | | | | | | | | | | |
 |`cusparseZcsrsv_solve`| |10.2| |11.0| | | | | | | | | | | | |
 |`cusparseZgemvi`|7.5| | | |`hipsparseZgemvi`|4.3.0| | | | |`rocsparse_zgemvi`|4.3.0| | | | |

diff --git a/docs/tables/CUSPARSE_API_supported_by_ROC.md b/docs/tables/CUSPARSE_API_supported_by_ROC.md
@@ -124,8 +124,8 @@
 |`csrilu02Info_t`| |12.2| | |`rocsparse_mat_info`|1.9.0| | | | |
 |`csrsm2Info`|9.2| | |12.0|`_rocsparse_mat_info`|1.9.0| | | | |
 |`csrsm2Info_t`|9.2| | |12.0|`rocsparse_mat_info`|1.9.0| | | | |
-|`csrsv2Info`| | | |12.0| | | | | | |
-|`csrsv2Info_t`| | | |12.0| | | | | | |
+|`csrsv2Info`| | | |12.0|`_rocsparse_mat_descr`|1.9.0| | | | |
+|`csrsv2Info_t`| | | |12.0|`rocsparse_mat_descr`|1.9.0| | | | |
 |`csru2csrInfo`| |12.2| | | | | | | | |
 |`csru2csrInfo_t`| |12.2| | | | | | | | |
 |`cusparseAction_t`| | | | |`rocsparse_action`|1.9.0| | | | |
@@ -303,7 +303,7 @@
 |`cusparseCcsrsv2_analysis`| |11.3| |12.0| | | | | | |
 |`cusparseCcsrsv2_bufferSize`| |11.3| |12.0| | | | | | |
 |`cusparseCcsrsv2_bufferSizeExt`| |11.3| |12.0| | | | | | |
-|`cusparseCcsrsv2_solve`| |11.3| |12.0| | | | | | |
+|`cusparseCcsrsv2_solve`| |11.3| |12.0|`rocsparse_ccsrsv_solve`|2.10.0| | | | |
 |`cusparseCcsrsv_analysis`| |10.2| |11.0| | | | | | |
 |`cusparseCcsrsv_solve`| |10.2| |11.0| | | | | | |
 |`cusparseCgemvi`|7.5| | | |`rocsparse_cgemvi`|4.3.0| | | | |
@@ -326,7 +326,7 @@
 |`cusparseDcsrsv2_analysis`| |11.3| |12.0| | | | | | |
 |`cusparseDcsrsv2_bufferSize`| |11.3| |12.0| | | | | | |
 |`cusparseDcsrsv2_bufferSizeExt`| |11.3| |12.0| | | | | | |
-|`cusparseDcsrsv2_solve`| |11.3| |12.0| | | | | | |
+|`cusparseDcsrsv2_solve`| |11.3| |12.0|`rocsparse_dcsrsv_solve`|1.9.0| | | | |
 |`cusparseDcsrsv_analysis`| |10.2| |11.0| | | | | | |
 |`cusparseDcsrsv_solve`| |10.2| |11.0| | | | | | |
 |`cusparseDgemvi`|7.5| | | |`rocsparse_dgemvi`|4.3.0| | | | |
@@ -345,7 +345,7 @@
 |`cusparseScsrsv2_analysis`| |11.3| |12.0| | | | | | |
 |`cusparseScsrsv2_bufferSize`| |11.3| |12.0| | | | | | |
 |`cusparseScsrsv2_bufferSizeExt`| |11.3| |12.0| | | | | | |
-|`cusparseScsrsv2_solve`| |11.3| |12.0| | | | | | |
+|`cusparseScsrsv2_solve`| |11.3| |12.0|`rocsparse_scsrsv_solve`|1.9.0| | | | |
 |`cusparseScsrsv_analysis`| |10.2| |11.0| | | | | | |
 |`cusparseScsrsv_solve`| |10.2| |11.0| | | | | | |
 |`cusparseSgemvi`|7.5| | | |`rocsparse_sgemvi`|4.3.0| | | | |
@@ -366,7 +366,7 @@
 |`cusparseZcsrsv2_analysis`| |11.3| |12.0| | | | | | |
 |`cusparseZcsrsv2_bufferSize`| |11.3| |12.0| | | | | | |
 |`cusparseZcsrsv2_bufferSizeExt`| |11.3| |12.0| | | | | | |
-|`cusparseZcsrsv2_solve`| |11.3| |12.0| | | | | | |
+|`cusparseZcsrsv2_solve`| |11.3| |12.0|`rocsparse_zcsrsv_solve`|2.10.0| | | | |
 |`cusparseZcsrsv_analysis`| |10.2| |11.0| | | | | | |
 |`cusparseZcsrsv_solve`| |10.2| |11.0| | | | | | |
 |`cusparseZgemvi`|7.5| | | |`rocsparse_zgemvi`|4.3.0| | | | |

diff --git a/src/CUDA2HIP_SPARSE_API_functions.cpp b/src/CUDA2HIP_SPARSE_API_functions.cpp
@@ -196,10 +196,10 @@ const std::map<llvm::StringRef, hipCounter> CUDA_SPARSE_FUNCTION_MAP {
   {"cusparseCcsrsv2_analysis",                          {"hipsparseCcsrsv2_analysis",                          "",                                                                 CONV_LIB_FUNC, API_SPARSE, 9, ROC_UNSUPPORTED | CUDA_DEPRECATED | CUDA_REMOVED}},
   {"cusparseZcsrsv2_analysis",                          {"hipsparseZcsrsv2_analysis",                          "",                                                                 CONV_LIB_FUNC, API_SPARSE, 9, ROC_UNSUPPORTED | CUDA_DEPRECATED | CUDA_REMOVED}},
 
-  {"cusparseScsrsv2_solve",                             {"hipsparseScsrsv2_solve",                             "",                                                                 CONV_LIB_FUNC, API_SPARSE, 9, ROC_UNSUPPORTED | CUDA_DEPRECATED | CUDA_REMOVED}},
-  {"cusparseDcsrsv2_solve",                             {"hipsparseDcsrsv2_solve",                             "",                                                                 CONV_LIB_FUNC, API_SPARSE, 9, ROC_UNSUPPORTED | CUDA_DEPRECATED | CUDA_REMOVED}},
-  {"cusparseCcsrsv2_solve",                             {"hipsparseCcsrsv2_solve",                             "",                                                                 CONV_LIB_FUNC, API_SPARSE, 9, ROC_UNSUPPORTED | CUDA_DEPRECATED | CUDA_REMOVED}},
-  {"cusparseZcsrsv2_solve",                             {"hipsparseZcsrsv2_solve",                             "",                                                                 CONV_LIB_FUNC, API_SPARSE, 9, ROC_UNSUPPORTED | CUDA_DEPRECATED | CUDA_REMOVED}},
+  {"cusparseScsrsv2_solve",                             {"hipsparseScsrsv2_solve",                             "rocsparse_scsrsv_solve",                                           CONV_LIB_FUNC, API_SPARSE, 9, CUDA_DEPRECATED | CUDA_REMOVED}},
+  {"cusparseDcsrsv2_solve",                             {"hipsparseDcsrsv2_solve",                             "rocsparse_dcsrsv_solve",                                           CONV_LIB_FUNC, API_SPARSE, 9, CUDA_DEPRECATED | CUDA_REMOVED}},
+  {"cusparseCcsrsv2_solve",                             {"hipsparseCcsrsv2_solve",                             "rocsparse_ccsrsv_solve",                                           CONV_LIB_FUNC, API_SPARSE, 9, CUDA_DEPRECATED | CUDA_REMOVED}},
+  {"cusparseZcsrsv2_solve",                             {"hipsparseZcsrsv2_solve",                             "rocsparse_zcsrsv_solve",                                           CONV_LIB_FUNC, API_SPARSE, 9, CUDA_DEPRECATED | CUDA_REMOVED}},
   {"cusparseXcsrsv2_zeroPivot",                         {"hipsparseXcsrsv2_zeroPivot",                         "",                                                                 CONV_LIB_FUNC, API_SPARSE, 9, ROC_UNSUPPORTED | CUDA_DEPRECATED | CUDA_REMOVED}},
 
   {"cusparseShybmv",                                    {"hipsparseShybmv",                                    "rocsparse_shybmv",                                                 CONV_LIB_FUNC, API_SPARSE, 9, CUDA_DEPRECATED | CUDA_REMOVED}},
@@ -2323,6 +2323,10 @@ const std::map<llvm::StringRef, hipAPIversions> HIP_SPARSE_FUNCTION_VER_MAP {
   {"rocsparse_chybmv",                                   {HIP_2100, HIP_0,    HIP_0   }},
   {"rocsparse_dhybmv",                                   {HIP_1090, HIP_0,    HIP_0   }},
   {"rocsparse_shybmv",                                   {HIP_1090, HIP_0,    HIP_0   }},
+  {"rocsparse_zcsrsv_solve",                             {HIP_2100, HIP_0,    HIP_0   }},
+  {"rocsparse_ccsrsv_solve",                             {HIP_2100, HIP_0,    HIP_0   }},
+  {"rocsparse_dcsrsv_solve",                             {HIP_1090, HIP_0,    HIP_0   }},
+  {"rocsparse_scsrsv_solve",                             {HIP_1090, HIP_0,    HIP_0   }},
 };
 
 const std::map<llvm::StringRef, cudaAPIChangedVersions> CUDA_SPARSE_FUNCTION_CHANGED_VER_MAP {

diff --git a/src/CUDA2HIP_SPARSE_API_types.cpp b/src/CUDA2HIP_SPARSE_API_types.cpp
@@ -38,8 +38,8 @@ const std::map<llvm::StringRef, hipCounter> CUDA_SPARSE_TYPE_NAME_MAP {
   {"cusparseSolveAnalysisInfo",                 {"hipsparseSolveAnalysisInfo",                 "",                                                   CONV_TYPE, API_SPARSE, 4, UNSUPPORTED | CUDA_DEPRECATED | CUDA_REMOVED}},
   {"cusparseSolveAnalysisInfo_t",               {"hipsparseSolveAnalysisInfo_t",               "",                                                   CONV_TYPE, API_SPARSE, 4, UNSUPPORTED | CUDA_DEPRECATED | CUDA_REMOVED}},
 
-  {"csrsv2Info",                                {"csrsv2Info",                                 "",                                                   CONV_TYPE, API_SPARSE, 4, UNSUPPORTED | CUDA_REMOVED}},
-  {"csrsv2Info_t",                              {"csrsv2Info_t",                               "",                                                   CONV_TYPE, API_SPARSE, 4, ROC_UNSUPPORTED | CUDA_REMOVED}},
+  {"csrsv2Info",                                {"csrsv2Info",                                 "_rocsparse_mat_descr",                               CONV_TYPE, API_SPARSE, 4, HIP_UNSUPPORTED | CUDA_REMOVED}},
+  {"csrsv2Info_t",                              {"csrsv2Info_t",                               "rocsparse_mat_descr",                                CONV_TYPE, API_SPARSE, 4, CUDA_REMOVED}},
 
   {"csrsm2Info",                                {"csrsm2Info",                                 "_rocsparse_mat_info",                                CONV_TYPE, API_SPARSE, 4, HIP_UNSUPPORTED | CUDA_REMOVED}},
   {"csrsm2Info_t",                              {"csrsm2Info_t",                               "rocsparse_mat_info",                                 CONV_TYPE, API_SPARSE, 4, CUDA_REMOVED}},

diff --git a/src/HipifyAction.cpp b/src/HipifyAction.cpp
@@ -155,6 +155,10 @@ const std::string sCusparseZgemvi_bufferSize = "cusparseZgemvi_bufferSize";
 const std::string sCusparseCgemvi_bufferSize = "cusparseCgemvi_bufferSize";
 const std::string sCusparseDgemvi_bufferSize = "cusparseDgemvi_bufferSize";
 const std::string sCusparseSgemvi_bufferSize = "cusparseSgemvi_bufferSize";
+const std::string sCusparseZcsrsv2_solve = "cusparseZcsrsv2_solve";
+const std::string sCusparseCcsrsv2_solve = "cusparseCcsrsv2_solve";
+const std::string sCusparseDcsrsv2_solve = "cusparseDcsrsv2_solve";
+const std::string sCusparseScsrsv2_solve = "cusparseScsrsv2_solve";
 
 // CUDA_OVERLOADED
 const std::string sCudaEventCreate = "cudaEventCreate";
@@ -1127,6 +1131,42 @@ std::map<std::string, ArgCastStruct> FuncArgCasts {
       false
     }
   },
+  {sCusparseZcsrsv2_solve,
+    {
+      {
+        {12, {e_replace_argument_with_const, cw_None, "rocsparse_solve_policy_auto"}}
+      },
+      true,
+      false
+    }
+  },
+  {sCusparseCcsrsv2_solve,
+    {
+      {
+        {12, {e_replace_argument_with_const, cw_None, "rocsparse_solve_policy_auto"}}
+      },
+      true,
+      false
+    }
+  },
+  {sCusparseDcsrsv2_solve,
+    {
+      {
+        {12, {e_replace_argument_with_const, cw_None, "rocsparse_solve_policy_auto"}}
+      },
+      true,
+      false
+    }
+  },
+  {sCusparseScsrsv2_solve,
+    {
+      {
+        {12, {e_replace_argument_with_const, cw_None, "rocsparse_solve_policy_auto"}}
+      },
+      true,
+      false
+    }
+  },
 };
 
 void HipifyAction::RewriteString(StringRef s, clang::SourceLocation start) {
@@ -1910,7 +1950,11 @@ std::unique_ptr<clang::ASTConsumer> HipifyAction::CreateASTConsumer(clang::Compi
             sCusparseZgemvi_bufferSize,
             sCusparseCgemvi_bufferSize,
             sCusparseDgemvi_bufferSize,
-            sCusparseSgemvi_bufferSize
+            sCusparseSgemvi_bufferSize,
+            sCusparseZcsrsv2_solve,
+            sCusparseCcsrsv2_solve,
+            sCusparseDcsrsv2_solve,
+            sCusparseScsrsv2_solve
           )
         )
       )

diff --git a/tests/lit.cfg b/tests/lit.cfg
@@ -104,6 +104,7 @@ if config.cuda_version_major >= 12:
     config.excludes.append('cub_02.cu')
     config.excludes.append('cub_03.cu')
     config.excludes.append('cusparse2rocsparse_9200_12000.cu')
+    config.excludes.append('cusparse2rocsparse_before_12000.cu')
 
 if config.llvm_version_major < 8:
     config.excludes.append('cd_intro.cu')