Merge pull request #853 from emankov/HIPIFY

[HIPIFY][#782][#783][feature] Initial hipification support for overloaded CUDA functions
ROCm · May 2, 2023 · c0494a4 · c0494a4
2 parents b276023 + 51db06a
commit c0494a4
Show file tree

Hide file tree

Showing 7 changed files with 104 additions and 2 deletions.
diff --git a/src/CUDA2HIP_Runtime_API_functions.cpp b/src/CUDA2HIP_Runtime_API_functions.cpp
@@ -170,7 +170,7 @@ const std::map<llvm::StringRef, hipCounter> CUDA_RUNTIME_FUNCTION_MAP {
   // 5. Event Management
   // no analogue
   // NOTE: Not equal to cuEventCreate due to different signatures
-  {"cudaEventCreate",                                         {"hipEventCreate",                                         "", CONV_EVENT, API_RUNTIME, SEC::EVENT}},
+  {"cudaEventCreate",                                         {"hipEventCreate",                                         "", CONV_EVENT, API_RUNTIME, SEC::EVENT, CUDA_OVERLOADED}},
   // cuEventCreate
   {"cudaEventCreateWithFlags",                                {"hipEventCreateWithFlags",                                "", CONV_EVENT, API_RUNTIME, SEC::EVENT}},
   // cuEventDestroy

diff --git a/src/CUDA2HIP_Scripting.h b/src/CUDA2HIP_Scripting.h
@@ -34,11 +34,19 @@ namespace hipify {
     e_move_argument,
   };
 
+  enum OverloadTypes {
+    ot_arguments_number,
+  };
+
   enum CastWarning {
     cw_None,
     cw_DataLoss,
   };
 
+  enum OverloadWarning {
+    ow_None,
+  };
+
   struct CastInfo {
     CastTypes castType;
     CastWarning castWarn;
@@ -54,11 +62,27 @@ namespace hipify {
     bool isToRoc = false;
     bool isToMIOpen = false;
   };
+
+  struct OverloadInfo {
+    hipCounter counter;
+    OverloadTypes overloadType;
+    OverloadWarning overloadWarn;
+  };
+
+  typedef std::map<unsigned, OverloadInfo> OverloadMap;
+
+  struct FuncOverloadsStruct {
+    OverloadMap overloadMap;
+    bool isToRoc = false;
+    bool isToMIOpen = false;
+  };
 }
 
 extern std::string getCastType(hipify::CastTypes c);
 extern std::map<std::string, hipify::ArgCastStruct> FuncArgCasts;
 
+extern std::map<std::string, hipify::FuncOverloadsStruct> FuncOverloads;
+
 namespace perl {
 
   bool generate(bool Generate = true);

diff --git a/src/HipifyAction.cpp b/src/HipifyAction.cpp
@@ -82,9 +82,12 @@ const std::string sCudnnSoftmaxBackward = "cudnnSoftmaxBackward";
 const std::string sCudnnConvolutionForward = "cudnnConvolutionForward";
 const std::string sCudnnConvolutionBackwardData = "cudnnConvolutionBackwardData";
 const std::string sCudnnRNNBackwardWeights = "cudnnRNNBackwardWeights";
+// CUDA_OVERLOADED
+const std::string sCudaEventCreate = "cudaEventCreate";
 // Matchers' names
 const StringRef sCudaLaunchKernel = "cudaLaunchKernel";
 const StringRef sCudaHostFuncCall = "cudaHostFuncCall";
+const StringRef sCudaOverloadedHostFuncCall = "cudaOverloadedHostFuncCall";
 const StringRef sCudaDeviceFuncCall = "cudaDeviceFuncCall";
 const StringRef sCubNamespacePrefix = "cubNamespacePrefix";
 const StringRef sCubFunctionTemplateDecl = "cubFunctionTemplateDecl";
@@ -104,6 +107,17 @@ std::string getCastType(hipify::CastTypes c) {
   }
 }
 
+std::map<std::string, hipify::FuncOverloadsStruct> FuncOverloads {
+  {sCudaEventCreate,
+    {
+      {
+        {1, {{"hipEventCreate", "", CONV_EVENT, API_RUNTIME, runtime::CUDA_RUNTIME_API_SECTIONS::EVENT}, ot_arguments_number, ow_None}},
+        {2, {{"hipEventCreateWithFlags", "", CONV_EVENT, API_RUNTIME, runtime::CUDA_RUNTIME_API_SECTIONS::EVENT}, ot_arguments_number, ow_None}}
+      }
+    }
+  },
+};
+
 std::map<std::string, ArgCastStruct> FuncArgCasts {
   {sCudaMemcpyToSymbol,
     {
@@ -882,6 +896,43 @@ bool HipifyAction::cudaHostFuncCall(const mat::MatchFinder::MatchResult &Result)
   return false;
 }
 
+bool HipifyAction::cudaOverloadedHostFuncCall(const mat::MatchFinder::MatchResult& Result) {
+  if (auto* call = Result.Nodes.getNodeAs<clang::CallExpr>(sCudaOverloadedHostFuncCall)) {
+    if (!call->getNumArgs()) return false;
+    auto* funcDcl = call->getDirectCallee();
+    if (!funcDcl) return false;
+    std::string name = funcDcl->getDeclName().getAsString();
+    const auto found = CUDA_RENAMES_MAP().find(name);
+    if (found == CUDA_RENAMES_MAP().end()) return false;
+    if (!Statistics::isCudaOverloaded(found->second)) return false;
+    auto it = FuncOverloads.find(name);
+    if (it == FuncOverloads.end()) return false;
+    auto FuncOverloadsStruct = it->second;
+    if (FuncOverloadsStruct.isToMIOpen != TranslateToMIOpen || FuncOverloadsStruct.isToRoc != TranslateToRoc) return false;
+    unsigned numArgs = call->getNumArgs();
+    auto itNumArgs = FuncOverloadsStruct.overloadMap.find(numArgs);
+    if (itNumArgs == FuncOverloadsStruct.overloadMap.end()) return false;
+    auto overrideInfo = itNumArgs->second;
+    auto counter = overrideInfo.counter;
+    // check if SUPPORTED
+    auto* SM = Result.SourceManager;
+    clang::SourceLocation s;
+    switch (overrideInfo.overloadType) {
+      case ot_arguments_number:
+      default:
+      {
+        s = call->getBeginLoc();
+        ct::Replacement Rep(*SM, s, name.size(), counter.hipName.str());
+        clang::FullSourceLoc fullSL(s, *SM);
+        insertReplacement(Rep, fullSL);
+        break;
+      }
+    }
+    return true;
+  }
+  return false;
+}
+
 bool HipifyAction::half2Member(const mat::MatchFinder::MatchResult &Result) {
   if (auto *expr = Result.Nodes.getNodeAs<clang::MemberExpr>(sHalf2Member)) {
     auto *baseExpr = expr->getBase();
@@ -982,6 +1033,19 @@ std::unique_ptr<clang::ASTConsumer> HipifyAction::CreateASTConsumer(clang::Compi
     ).bind(sCudaHostFuncCall),
     this
   );
+  Finder->addMatcher(
+    mat::callExpr(
+      mat::isExpansionInMainFile(),
+      mat::callee(
+        mat::functionDecl(
+          mat::hasAnyName(
+            sCudaEventCreate
+          )
+        )
+      )
+    ).bind(sCudaOverloadedHostFuncCall),
+    this
+  );
   Finder->addMatcher(
     mat::callExpr(
       mat::isExpansionInMainFile(),
@@ -1148,6 +1212,7 @@ void HipifyAction::ExecuteAction() {
 void HipifyAction::run(const mat::MatchFinder::MatchResult &Result) {
   if (cudaLaunchKernel(Result)) return;
   if (cudaHostFuncCall(Result)) return;
+  if (cudaOverloadedHostFuncCall(Result)) return;
   if (cudaDeviceFuncCall(Result)) return;
   if (cubNamespacePrefix(Result)) return;
   if (cubFunctionTemplateDecl(Result)) return;

diff --git a/src/HipifyAction.h b/src/HipifyAction.h
@@ -74,6 +74,7 @@ class HipifyAction : public clang::ASTFrontendAction,
   bool cudaLaunchKernel(const mat::MatchFinder::MatchResult &Result);
   bool cudaDeviceFuncCall(const mat::MatchFinder::MatchResult &Result);
   bool cudaHostFuncCall(const mat::MatchFinder::MatchResult &Result);
+  bool cudaOverloadedHostFuncCall(const mat::MatchFinder::MatchResult &Result);
   bool cubNamespacePrefix(const mat::MatchFinder::MatchResult &Result);
   bool cubFunctionTemplateDecl(const mat::MatchFinder::MatchResult &Result);
   bool cubUsingNamespaceDecl(const mat::MatchFinder::MatchResult &Result);

diff --git a/src/Statistics.cpp b/src/Statistics.cpp
@@ -425,6 +425,10 @@ bool Statistics::isRocMiopenOnly(const hipCounter& counter) {
   return ROC_MIOPEN_ONLY == (counter.supportDegree & ROC_MIOPEN_ONLY);
 }
 
+bool Statistics::isCudaOverloaded(const hipCounter& counter) {
+  return CUDA_OVERLOADED == (counter.supportDegree & CUDA_OVERLOADED);
+}
+
 std::string Statistics::getCudaVersion(const cudaVersions& ver) {
   switch (ver) {
     case CUDA_0:

diff --git a/src/Statistics.h b/src/Statistics.h
@@ -193,7 +193,8 @@ enum SupportDegree {
   REMOVED = 0x100,
   HIP_EXPERIMENTAL = 0x200,
   HIP_SUPPORTED_V2_ONLY = 0x400,
-  ROC_MIOPEN_ONLY = 0x800
+  ROC_MIOPEN_ONLY = 0x800,
+  CUDA_OVERLOADED = 0x1000
 };
 
 enum cudaVersions {
@@ -468,6 +469,8 @@ class Statistics {
   static bool isHipSupportedV2Only(const hipCounter& counter);
   // Check whether the counter is ROC_MIOPEN_ONLY or not.
   static bool isRocMiopenOnly(const hipCounter& counter);
+  // Check whether the counter is CUDA_OVERLOADED or not.
+  static bool isCudaOverloaded(const hipCounter& counter);
   // Get string CUDA version.
   static std::string getCudaVersion(const cudaVersions &ver);
   // Get string HIP version.

diff --git a/tests/unit_tests/synthetic/runtime_functions.cu b/tests/unit_tests/synthetic/runtime_functions.cu
@@ -1049,6 +1049,11 @@ int main() {
   // CHECK: result = hipEventCreate(&Event_t);
   result = cudaEventCreate(&Event_t);
 
+  // CUDA: static __inline__ __host__ cudaError_t cudaEventCreate(cudaEvent_t* event, unsigned int flags);
+  // HIP: hipError_t hipEventCreateWithFlags(hipEvent_t* event, unsigned flags);
+  // CHECK: result = hipEventCreateWithFlags(&Event_t, flags);
+  result = cudaEventCreate(&Event_t, flags);
+
   // CUDA: extern __host__ __cudart_builtin__ cudaError_t CUDARTAPI cudaEventCreateWithFlags(cudaEvent_t *event, unsigned int flags);
   // HIP: hipError_t hipEventCreateWithFlags(hipEvent_t* event, unsigned flags);
   // CHECK: result = hipEventCreateWithFlags(&Event_t, flags);