[HIPIFY][#1439] Add reinterpret_cast to args of some functions

+ Affected functions: hipFuncSetCacheConfig, hipFuncGetAttributes
+ Add a corresponding Matcher cudaReinterpretCastArgFuncCall
+ Add reinterpret_cast.cu test

TODO: Do the same for hipify-perl

hipamd/hipify-clang/src/HipifyAction.cpp

@@ -42,6 +42,8 @@ const std::string sCudaGetSymbolSize = "cudaGetSymbolSize";
 const std::string sCudaGetSymbolAddress = "cudaGetSymbolAddress";
 const std::string sCudaMemcpyFromSymbol = "cudaMemcpyFromSymbol";
 const std::string sCudaMemcpyFromSymbolAsync = "cudaMemcpyFromSymbolAsync";
+const std::string sCudaFuncSetCacheConfig = "cudaFuncSetCacheConfig";
+const std::string sCudaFuncGetAttributes = "cudaFuncGetAttributes";
 
 std::set<std::string> DeviceSymbolFunctions0 {
   {sCudaMemcpyToSymbol},
@@ -413,29 +415,31 @@ bool HipifyAction::cudaSharedIncompleteArrayVar(const clang::ast_matchers::Match
     insertReplacement(Rep, fullSL);
     hipCounter counter = {"HIP_DYNAMIC_SHARED", "", ConvTypes::CONV_MEMORY, ApiTypes::API_RUNTIME};
     Statistics::current().incrementCounter(counter, refName.str());
+    return true;
   }
-  return true;
+  return false;
 }
 
 bool HipifyAction::cudaDeviceFuncCall(const clang::ast_matchers::MatchFinder::MatchResult& Result) {
   if (const clang::CallExpr *call = Result.Nodes.getNodeAs<clang::CallExpr>("cudaDeviceFuncCall")) {
     const clang::FunctionDecl *funcDcl = call->getDirectCallee();
     if (!funcDcl) {
-      return true;
+      return false;
     }
     FindAndReplace(funcDcl->getDeclName().getAsString(), llcompat::getBeginLoc(call), CUDA_DEVICE_FUNC_MAP, false);
+    return true;
   }
-  return true;
+  return false;
 }
 
 bool HipifyAction::cudaSymbolFuncCall(const clang::ast_matchers::MatchFinder::MatchResult& Result) {
   if (const clang::CallExpr * call = Result.Nodes.getNodeAs<clang::CallExpr>("cudaSymbolFuncCall")) {
     if (!call->getNumArgs()) {
-      return true;
+      return false;
     }
     const clang::FunctionDecl* funcDcl = call->getDirectCallee();
     if (!funcDcl) {
-      return true;
+      return false;
     }
     std::string sName = funcDcl->getDeclName().getAsString();
     unsigned int argNum = 0;
@@ -444,7 +448,7 @@ bool HipifyAction::cudaSymbolFuncCall(const clang::ast_matchers::MatchFinder::Ma
     } else if (call->getNumArgs() > 1 && DeviceSymbolFunctions1.find(sName) != DeviceSymbolFunctions1.end()) {
       argNum = 1;
     } else {
-      return true;
+      return false;
     }
     clang::SmallString<40> XStr;
     llvm::raw_svector_ostream OS(XStr);
@@ -460,8 +464,46 @@ bool HipifyAction::cudaSymbolFuncCall(const clang::ast_matchers::MatchFinder::Ma
     ct::Replacement Rep(*SM, s, length, OS.str());
     clang::FullSourceLoc fullSL(s, *SM);
     insertReplacement(Rep, fullSL);
+    return true;
   }
-  return true;
+  return false;
+}
+
+bool HipifyAction::cudaReinterpretCastArgFuncCall(const clang::ast_matchers::MatchFinder::MatchResult& Result) {
+  if (const clang::CallExpr * call = Result.Nodes.getNodeAs<clang::CallExpr>("cudaReinterpretCastArgFuncCall")) {
+    if (!call->getNumArgs()) {
+      return false;
+    }
+    const clang::FunctionDecl* funcDcl = call->getDirectCallee();
+    if (!funcDcl) {
+      return false;
+    }
+    std::string sName = funcDcl->getDeclName().getAsString();
+    unsigned int argNum = 0;
+    if (sCudaFuncSetCacheConfig == sName) {
+      argNum = 0;
+    } else if (call->getNumArgs() > 1 && sCudaFuncGetAttributes == sName) {
+      argNum = 1;
+    } else {
+      return false;
+    }
+    clang::SmallString<40> XStr;
+    llvm::raw_svector_ostream OS(XStr);
+    clang::SourceRange sr = call->getArg(argNum)->getSourceRange();
+    clang::SourceManager* SM = Result.SourceManager;
+    const std::string sCast = "reinterpret_cast<const void*>";
+    OS << sCast << "(" << readSourceText(*SM, sr) << ")";
+    clang::SourceRange replacementRange = getWriteRange(*SM, { sr.getBegin(), sr.getEnd() });
+    clang::SourceLocation s = replacementRange.getBegin();
+    clang::SourceLocation e = replacementRange.getEnd();
+    clang::LangOptions DefaultLangOptions;
+    size_t length = SM->getCharacterData(clang::Lexer::getLocForEndOfToken(e, 0, *SM, DefaultLangOptions)) - SM->getCharacterData(s);
+    ct::Replacement Rep(*SM, s, length, OS.str());
+    clang::FullSourceLoc fullSL(s, *SM);
+    insertReplacement(Rep, fullSL);
+    return true;
+  }
+  return false;
 }
 
 void HipifyAction::insertReplacement(const ct::Replacement& rep, const clang::FullSourceLoc& fullSL) {
@@ -505,6 +547,20 @@ std::unique_ptr<clang::ASTConsumer> HipifyAction::CreateASTConsumer(clang::Compi
     ).bind("cudaSymbolFuncCall"),
     this
   );
+  Finder->addMatcher(
+    mat::callExpr(
+      mat::isExpansionInMainFile(),
+      mat::callee(
+        mat::functionDecl(
+          mat::hasAnyName(
+            sCudaFuncSetCacheConfig,
+            sCudaFuncGetAttributes
+          )
+        )
+      )
+    ).bind("cudaReinterpretCastArgFuncCall"),
+    this
+  );
   Finder->addMatcher(
     mat::callExpr(
       mat::isExpansionInMainFile(),
@@ -643,5 +699,6 @@ void HipifyAction::run(const clang::ast_matchers::MatchFinder::MatchResult& Resu
   if (cudaLaunchKernel(Result)) return;
   if (cudaSharedIncompleteArrayVar(Result)) return;
   if (cudaSymbolFuncCall(Result)) return;
+  if (cudaReinterpretCastArgFuncCall(Result)) return;
   if (cudaDeviceFuncCall(Result)) return;
 }

hipamd/hipify-clang/src/HipifyAction.h

@@ -72,6 +72,8 @@ public:
   bool cudaSharedIncompleteArrayVar(const clang::ast_matchers::MatchFinder::MatchResult& Result);
   bool cudaDeviceFuncCall(const clang::ast_matchers::MatchFinder::MatchResult& Result);
   bool cudaSymbolFuncCall(const clang::ast_matchers::MatchFinder::MatchResult& Result);
+  bool cudaReinterpretCastArgFuncCall(const clang::ast_matchers::MatchFinder::MatchResult& Result);
+
   // Called by the preprocessor for each include directive during the non-raw lexing pass.
   void InclusionDirective(clang::SourceLocation hash_loc,
                           const clang::Token &include_token,

hipamd/tests/hipify-clang/unit_tests/samples/reinterpret_cast.cu

@@ -0,0 +1,53 @@
+// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args
+
+/*
+Copyright (c) 2015-present Advanced Micro Devices, Inc. All rights reserved.
+
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+#include <stdio.h>
+// CHECK: #include <hip/hip_runtime.h>
+#include <cuda_runtime.h>
+
+__global__
+void fn(float* px, float* py) {
+  bool a[42];
+  __shared__ double b[69];
+  for (auto&& x : b) x = *py++;
+  for (auto&& x : a) x = *px++ > 0.0;
+  for (auto&& x : a) if (x)* --py = *--px;
+}
+
+int main() {
+  // CHECK: hipFuncCache_t cacheConfig;
+  cudaFuncCache cacheConfig;
+  void* func;
+  // CHECK: hipFuncSetCacheConfig(reinterpret_cast<const void*>(func), cacheConfig);
+  cudaFuncSetCacheConfig(func, cacheConfig);
+  // CHECK: hipFuncAttributes attr{};
+  cudaFuncAttributes attr{};
+  // CHECK: auto r = hipFuncGetAttributes(&attr, reinterpret_cast<const void*>(&fn));
+  auto r = cudaFuncGetAttributes(&attr, &fn);
+  // CHECK: if (r != hipSuccess || attr.maxThreadsPerBlock == 0) {
+  if (r != cudaSuccess || attr.maxThreadsPerBlock == 0) {
+    return 1;
+  }
+  return 0;
+}