Merge pull request #1838 from emankov/HIPIFY

[HIPIFY][SWDEV-430220][LRT][Workaround] Use `const_cast<const char**>` in `hiprtcCreateProgram` and `hiprtcCompileProgram`
ROCm · Jan 21, 2025 · e939890 · e939890
2 parents 8004179 + 18c4c7a
commit e939890
Show file tree

Hide file tree

Showing 5 changed files with 68 additions and 3 deletions.
diff --git a/bin/hipify-perl b/bin/hipify-perl
@@ -6843,7 +6843,7 @@ sub simpleSubstitutions {
     subst("cusolver_common.h", "hipsolver.h", "include_cuda_main_header");
     subst("cusparse.h", "hipsparse.h", "include_cuda_main_header");
     subst("cutensor.h", "hiptensor.h", "include_cuda_main_header");
-    subst("nvrtc.h", "hiprtc.h", "include_cuda_main_header");
+    subst("nvrtc.h", "hip\/hiprtc.h", "include_cuda_main_header");
     subst("cublas_v2.h", "hipblas.h", "include_cuda_main_header_v2");
     subst("cusparse_v2.h", "hipsparse.h", "include_cuda_main_header_v2");
     subst("CUDAContext", "HIPContext", "type");

diff --git a/src/CUDA2HIP.cpp b/src/CUDA2HIP.cpp
@@ -98,7 +98,7 @@ const std::map <llvm::StringRef, hipCounter> CUDA_INCLUDE_MAP {
   {"caffe2/operators/gather_op.cuh",                        {"caffe2/operators/math/gather_op.cuh",                   "",                                                               CONV_INCLUDE,                API_CAFFE2, 0, UNSUPPORTED}},
   {"caffe2/core/common_cudnn.h",                            {"caffe2/core/hip/common_miopen.h",                       "",                                                               CONV_INCLUDE,                API_CAFFE2, 0}},
   // RTC includes
-  {"nvrtc.h",                                               {"hiprtc.h",                                              "", CONV_INCLUDE_CUDA_MAIN_H, API_RTC, 0}},
+  {"nvrtc.h",                                               {"hip/hiprtc.h",                                          "",                                                               CONV_INCLUDE_CUDA_MAIN_H, API_RTC, 0}},
 };
 
 const std::map<llvm::StringRef, hipCounter> &CUDA_RENAMES_MAP() {

diff --git a/src/CUDA2HIP_Scripting.h b/src/CUDA2HIP_Scripting.h
@@ -32,6 +32,7 @@ namespace hipify {
     e_HIP_SYMBOL,
     e_reinterpret_cast,
     e_reinterpret_cast_size_t,
+    e_const_cast_char_pp,
     e_static_cast_uint64_t,
     e_int32_t,
     e_int64_t,

diff --git a/src/HipifyAction.cpp b/src/HipifyAction.cpp
@@ -47,6 +47,7 @@ std::string sHIP_SYMBOL = "HIP_SYMBOL";
 std::string s_reinterpret_cast = "reinterpret_cast<const void*>";
 std::string s_reinterpret_cast_size_t = "reinterpret_cast<size_t*>";
 std::string s_static_cast_uint64_t = "static_cast<uint64_t>";
+std::string s_const_cast_char_pp = "const_cast<const char**>";
 std::string s_int32_t = "int32_t";
 std::string s_int64_t = "int64_t";
 const std::string sHipLaunchKernelGGL = "hipLaunchKernelGGL";
@@ -233,6 +234,8 @@ const std::string sCusparseSpMV_bufferSize = "cusparseSpMV_bufferSize";
 const std::string sCusparseSpMM_preprocess = "cusparseSpMM_preprocess";
 const std::string sCusparseSpSV_bufferSize = "cusparseSpSV_bufferSize";
 const std::string sCudaGetDriverEntryPoint = "cudaGetDriverEntryPoint";
+const std::string sNvrtcCompileProgram = "nvrtcCompileProgram";
+const std::string sNvrtcCreateProgram = "nvrtcCreateProgram";
 
 // CUDA_OVERLOADED
 const std::string sCudaEventCreate = "cudaEventCreate";
@@ -253,6 +256,7 @@ std::string getCastType(hipify::CastTypes c) {
     case e_HIP_SYMBOL: return sHIP_SYMBOL;
     case e_reinterpret_cast: return s_reinterpret_cast;
     case e_reinterpret_cast_size_t: return s_reinterpret_cast_size_t;
+    case e_const_cast_char_pp: return s_const_cast_char_pp;
     case e_static_cast_uint64_t: return s_static_cast_uint64_t;
     case e_int32_t: return s_int32_t;
     case e_int64_t: return s_int64_t;
@@ -291,6 +295,25 @@ std::map<std::string, hipify::FuncOverloadsStruct> FuncOverloads {
 };
 
 std::map<std::string, std::vector<ArgCastStruct>> FuncArgCasts {
+  {sNvrtcCompileProgram,
+    {
+      {
+        {
+          {2, {e_const_cast_char_pp, cw_None}}
+        }
+      }
+    }
+  },
+  {sNvrtcCreateProgram,
+    {
+      {
+        {
+          {4, {e_const_cast_char_pp, cw_None}},
+          {5, {e_const_cast_char_pp, cw_None}}
+        }
+      }
+    }
+  },
   {sCudaGetDriverEntryPoint,
     {
       {
@@ -3112,7 +3135,9 @@ std::unique_ptr<clang::ASTConsumer> HipifyAction::CreateASTConsumer(clang::Compi
             sCusparseSpMV_bufferSize,
             sCusparseSpMM_preprocess,
             sCusparseSpSV_bufferSize,
-            sCudaGetDriverEntryPoint
+            sCudaGetDriverEntryPoint,
+            sNvrtcCompileProgram,
+            sNvrtcCreateProgram
           )
         )
       )

diff --git a/tests/unit_tests/synthetic/nvrtc2hiprtc.cu b/tests/unit_tests/synthetic/nvrtc2hiprtc.cu
@@ -0,0 +1,39 @@
+// RUN: %run_test hipify "%s" "%t" %hipify_args 3 --amap --skip-excluded-preprocessor-conditional-blocks --experimental %clang_args -ferror-limit=500
+
+// CHECK: #include <hip/hip_runtime.h>
+#include <cuda_runtime.h>
+#include <stdio.h>
+// CHECK: #include "hip/hiprtc.h"
+#include "nvrtc.h"
+// CHECK-NOT: #include "hip/hiprtc.h"
+// CHECK-NOT: #include "nvrtc.h"
+
+int main() {
+  printf("25. CUDA nvRTC API to HIP hipRTC API synthetic test\n");
+
+  int numOptions = 0;
+  int numHeaders = 0;
+  const char* pOptions = nullptr;
+  const char* pHeadedrs = nullptr;
+  const char* pIncludeNames = nullptr;
+  char* pchSrc = nullptr;
+  char* pchName = nullptr;
+
+  // CHECK: hiprtcProgram rtcProgram = nullptr;
+  nvrtcProgram rtcProgram = nullptr;
+
+  // CHECK: hiprtcResult rtcResult;
+  nvrtcResult rtcResult;
+
+  // CUDA: nvrtcResult nvrtcCompileProgram(nvrtcProgram prog, int numOptions, const char* const* options);
+  // HIP: hiprtcResult hiprtcCompileProgram(hiprtcProgram prog, int numOptions, const char** options);
+  // CHECK: rtcResult = hiprtcCompileProgram(rtcProgram, numOptions, const_cast<const char**>(&pOptions));
+  rtcResult = nvrtcCompileProgram(rtcProgram, numOptions, &pOptions);
+
+  // CUDA: nvrtcResult nvrtcCreateProgram(nvrtcProgram *prog, const char* src, const char* name, int numHeaders, const char* const* headers, const char* const* includeNames);
+  // HIP: hiprtcResult hiprtcCreateProgram(hiprtcProgram* prog, const char* src, const char* name, int numHeaders, const char** headers, const char** includeNames);
+  // CHECK: rtcResult = hiprtcCreateProgram(&rtcProgram, pchSrc, pchName, numHeaders, const_cast<const char**>(&pHeadedrs), const_cast<const char**>(&pIncludeNames));
+  rtcResult = nvrtcCreateProgram(&rtcProgram, pchSrc, pchName, numHeaders, &pHeadedrs, &pIncludeNames);
+
+  return 0;
+}