[HIPIFY][ROCm#782][ROCm#783][feature] Initial hipification support fo…

…r CUDA overloaded functions

[IMP]
+ Introduced 1-to-N hipification for overloaded APIs with different number of arguments (see the `cudaEventCreate` example in ROCm#783)
+ Introduced a new marker `CUDA_OVERLOADED` and the corresponding Matcher `cudaOverloadedHostFuncCall`
+ Currently, `CUDA_OVERLOADED` APIs are hipified twice: first, by rewriting string blindly, and second, by overload matcher (correcting)
+ Added a synthetic test for cudaEventCreate to `hipEventCreateWithFlags` hipification in `runtime_functions.cu`

[ToDo]
+ Take into account other markers, firstly `HIP_UNSUPPORTED`, in overloaded hipification
+ Take into account `CUDA_VERSION` for overloaded APIs: currently, versioning is provided only for a single instance of the overloaded API
+ Find all the existing overloaded CUDA APIs, update HIPIFY correspondingly for those APIs, and provide tests
+ [sub-task] Implement overloading with the same number of arguments and different argument types
+ [optionally] Think about implementing the same in hipify-perl

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

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);

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;

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);

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:

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.

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);