From bc9b970f823a373bb7d49d97a36e2d06fce5e4aa Mon Sep 17 00:00:00 2001
From: Evgeny Mankov <Evgeniy.Mankov@amd.com>
Date: Wed, 24 May 2017 18:25:40 +0300
Subject: [PATCH] [FIX] [HIPIFY] Add matchers for function return types.

https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/issues/73

Examples (https://github.com/thrust/thrust/blob/master/thrust/system/cuda/detail/trivial_copy.inl):

template<typename System1,
         typename System2>
cudaStream_t cuda_memcpy_stream(const thrust::cpp::execution_policy<System1> &,
                                const thrust::cuda::execution_policy<System2> &exec)

template<typename System1,
         typename System2>
cudaMemcpyKind cuda_memcpy_kind(const thrust::cuda::execution_policy<System1> &,
                                const thrust::cpp::execution_policy<System2> &)


[ROCm/hip commit: c8632156114577404bd1279aa3903c7a022f4409]
---
 projects/hip/hipify-clang/src/Cuda2Hip.cpp | 38 ++++++++++++++++++++++
 1 file changed, 38 insertions(+)

diff --git a/projects/hip/hipify-clang/src/Cuda2Hip.cpp b/projects/hip/hipify-clang/src/Cuda2Hip.cpp
index dcb9c3d216..f17c3e2646 100644
--- a/projects/hip/hipify-clang/src/Cuda2Hip.cpp
+++ b/projects/hip/hipify-clang/src/Cuda2Hip.cpp
@@ -3123,6 +3123,33 @@ private:
     return false;
   }
 
+  bool cudaFunctionReturn(const MatchFinder::MatchResult &Result) {
+    if (const auto *ret = Result.Nodes.getNodeAs<FunctionDecl>("cudaFunctionReturn")) {
+      QualType QT = ret->getReturnType();
+      SourceManager *SM = Result.SourceManager;
+      SourceRange sr = ret->getReturnTypeSourceRange();
+      SourceLocation sl = sr.getBegin();
+      std::string name = QT.getAsString();
+      if (QT.getTypePtr()->isEnumeralType()) {
+        name = QT.getTypePtr()->getAs<EnumType>()->getDecl()->getNameAsString();
+      }
+      const auto found = N.cuda2hipRename.find(name);
+      if (found != N.cuda2hipRename.end()) {
+        updateCounters(found->second, name);
+        if (!found->second.unsupported) {
+          StringRef repName = found->second.hipName;
+          Replacement Rep(*SM, sl, name.size(), repName);
+          FullSourceLoc fullSL(sl, *SM);
+          insertReplacement(Rep, fullSL);
+        }
+      }
+      else {
+        std::string msg = "the following reference is not handled: '" + name + "' [function return].";
+        printHipifyMessage(*SM, sl, msg);
+      }
+    }
+    return false;
+  }
 
   bool cudaSharedIncompleteArrayVar(const MatchFinder::MatchResult &Result) {
     StringRef refName = "cudaSharedIncompleteArrayVar";
@@ -3269,6 +3296,7 @@ public:
       if (cudaParamDeclPtr(Result)) break;
       if (cudaLaunchKernel(Result)) break;
       if (cudaNewOperatorDecl(Result)) break;
+      if (cudaFunctionReturn(Result)) break;
       if (cudaSharedIncompleteArrayVar(Result)) break;
       if (stringLiteral(Result)) break;
       if (unresolvedTemplateName(Result)) break;
@@ -3373,6 +3401,16 @@ void addAllMatchers(ast_matchers::MatchFinder &Finder, Cuda2HipCallback *Callbac
                                hasType(pointsTo(namedDecl(matchesName("cu.*|CU.*")))))
                               .bind("cudaNewOperatorDecl"),
                                Callback);
+  // Examples:
+  // 1.
+  // cudaStream_t cuda_memcpy_stream(...)
+  // 2.
+  // template<typename System1, typename System2> cudaMemcpyKind cuda_memcpy_kind(...)
+  Finder.addMatcher(functionDecl(isExpansionInMainFile(),
+                                 returns(hasDeclaration(namedDecl(matchesName("cu.*|CU.*")))))
+                                .bind("cudaFunctionReturn"),
+                                 Callback);
+
 }
 
 int64_t printStats(const std::string &csvFile, const std::string &srcFile,