From e265ed41ed89fbaa26b10f8fda020670fad3af97 Mon Sep 17 00:00:00 2001 From: atimofee Date: Fri, 4 Mar 2016 22:07:41 +0300 Subject: [PATCH] Template kernels calls and declaration conversion added [ROCm/clr commit: 236c38719c04d10becb7b60e5508756f3f091bfa] --- projects/clr/hipamd/src/Cuda2Hip.cpp | 96 ++++++++++++++++++---------- 1 file changed, 63 insertions(+), 33 deletions(-) diff --git a/projects/clr/hipamd/src/Cuda2Hip.cpp b/projects/clr/hipamd/src/Cuda2Hip.cpp index f2a18262a9..434d74e584 100644 --- a/projects/clr/hipamd/src/Cuda2Hip.cpp +++ b/projects/clr/hipamd/src/Cuda2Hip.cpp @@ -324,7 +324,7 @@ namespace { const MacroDefinition &MD, SourceRange Range, const MacroArgs *Args) override { - if (_sm->isWrittenInMainFile(Range.getBegin())) + if (_sm->isWrittenInMainFile(MacroNameTok.getLocation())) { for (unsigned int i = 0; Args && i < MD.getMacroInfo()->getNumArgs(); i++) { @@ -386,12 +386,43 @@ namespace { }; class Cuda2HipCallback : public MatchFinder::MatchCallback { - public: - Cuda2HipCallback(Replacements *Replace) : Replace(Replace) {} + public: + Cuda2HipCallback(Replacements *Replace, ast_matchers::MatchFinder *parent) + : Replace(Replace), owner(parent) {} - void run(const MatchFinder::MatchResult &Result) override { + void convertKernelDecl(const FunctionDecl * kernelDecl, const MatchFinder::MatchResult &Result) + { + SourceManager * SM = Result.SourceManager; + LangOptions DefaultLangOptions; - SourceManager * SM = Result.SourceManager; + SmallString<40> XStr; + raw_svector_ostream OS(XStr); + StringRef initialParamList; + OS << "hipLaunchParm lp"; + size_t replacementLength = OS.str().size(); + SourceLocation sl = kernelDecl->getNameInfo().getEndLoc(); + SourceLocation kernelArgListStart = clang::Lexer::findLocationAfterToken(sl, clang::tok::l_paren, *SM, DefaultLangOptions, true); + kernelArgListStart.dump(*SM); + if (kernelDecl->getNumParams() > 0) { + const ParmVarDecl * pvdFirst = kernelDecl->getParamDecl(0); + const ParmVarDecl * pvdLast = kernelDecl->getParamDecl(kernelDecl->getNumParams() - 1); + SourceLocation kernelArgListStart(pvdFirst->getLocStart()); + SourceLocation kernelArgListEnd(pvdLast->getLocEnd()); + SourceLocation stop = clang::Lexer::getLocForEndOfToken(kernelArgListEnd, 0, *SM, DefaultLangOptions); + size_t replacementLength = SM->getCharacterData(stop) - SM->getCharacterData(kernelArgListStart); + initialParamList = StringRef(SM->getCharacterData(kernelArgListStart), replacementLength); + OS << ", " << initialParamList; + } + llvm::outs() << "initial paramlist: " << initialParamList << "\n"; + llvm::outs() << "new paramlist: " << OS.str() << "\n"; + Replacement Rep0(*(Result.SourceManager), kernelArgListStart, replacementLength, OS.str()); + Replace->insert(Rep0); + } + + void run(const MatchFinder::MatchResult &Result) override { + + SourceManager * SM = Result.SourceManager; + LangOptions DefaultLangOptions; if (const CallExpr * call = Result.Nodes.getNodeAs("cudaCall")) { @@ -406,35 +437,27 @@ class Cuda2HipCallback : public MatchFinder::MatchCallback { } } - if (const CUDAKernelCallExpr * launchKernel = Result.Nodes.getNodeAs("cudaLaunchKernel")) - { - LangOptions DefaultLangOptions; - StringRef initialParamList; - SmallString<40> XStr; - raw_svector_ostream OS(XStr); - OS << "hipLaunchParm lp"; - const FunctionDecl * kernelDecl = launchKernel->getDirectCallee(); - SourceLocation l1 = kernelDecl->getNameInfo().getLocStart(); - l1.dump(*SM);llvm::outs() << "\n"; - SourceLocation kernelArgListStart = clang::Lexer::findLocationAfterToken(l1, clang::tok::l_paren, *SM, DefaultLangOptions, true); - kernelArgListStart.dump(*SM);llvm::outs() << "\n"; - size_t replacementLength = 0; - if (kernelDecl->getNumParams() > 0) { - const ParmVarDecl * pvdLast = kernelDecl->getParamDecl(kernelDecl->getNumParams()-1); - SourceLocation kernelArgListEnd = SourceLocation(pvdLast->getLocEnd()); - SourceLocation stop = clang::Lexer::getLocForEndOfToken(kernelArgListEnd, 0, *SM, DefaultLangOptions); - replacementLength = SM->getCharacterData(stop) - SM->getCharacterData(kernelArgListStart); - initialParamList = StringRef(SM->getCharacterData(kernelArgListStart), replacementLength); - OS << ", " << initialParamList; + if (const CUDAKernelCallExpr * launchKernel = Result.Nodes.getNodeAs("cudaLaunchKernel")) + { + SmallString<40> XStr; + raw_svector_ostream OS(XStr); + StringRef calleeName; + const FunctionDecl * kernelDecl = launchKernel->getDirectCallee(); + if (kernelDecl) { + calleeName = kernelDecl->getName(); + convertKernelDecl(kernelDecl, Result); + } + else { + const Expr * e = launchKernel->getCallee(); + if (const UnresolvedLookupExpr * ule = dyn_cast(e)) { + calleeName = ule->getName().getAsIdentifierInfo()->getName(); + owner->addMatcher(functionTemplateDecl(hasName(calleeName)).bind("unresolvedTemplateName"), this); } + } - llvm::outs() << "initial paramlist: " << initialParamList << "\n"; - llvm::outs() << "new paramlist: " << OS.str() << "\n"; - Replacement Rep0(*(Result.SourceManager), kernelArgListStart, replacementLength, OS.str()); - Replace->insert(Rep0); XStr.clear(); - OS << "hipLaunchKernel(HIP_KERNEL_NAME(" << kernelDecl->getName() << "), "; + OS << "hipLaunchKernel(HIP_KERNEL_NAME(" << calleeName << "), "; const CallExpr * config = launchKernel->getConfig(); llvm::outs() << "\nKernel config arguments:\n"; @@ -475,6 +498,12 @@ class Cuda2HipCallback : public MatchFinder::MatchCallback { Replace->insert(Rep); } + if (const FunctionTemplateDecl * templateDecl = Result.Nodes.getNodeAs("unresolvedTemplateName")) + { + FunctionDecl * kernelDecl = templateDecl->getTemplatedDecl(); + convertKernelDecl(kernelDecl, Result); + } + if (const MemberExpr * threadIdx = Result.Nodes.getNodeAs("cudaBuiltin")) { if (const OpaqueValueExpr * refBase = dyn_cast(threadIdx->getBase())) { @@ -591,6 +620,7 @@ class Cuda2HipCallback : public MatchFinder::MatchCallback { private: Replacements *Replace; + ast_matchers::MatchFinder * owner; struct hipName N; }; @@ -631,7 +661,7 @@ int main(int argc, const char **argv) { RefactoringTool Tool(OptionsParser.getCompilations(), dst); ast_matchers::MatchFinder Finder; - Cuda2HipCallback Callback(&Tool.getReplacements()); + Cuda2HipCallback Callback(&Tool.getReplacements(), &Finder); HipifyPPCallbacks PPCallbacks(&Tool.getReplacements()); Finder.addMatcher(callExpr(isExpansionInMainFile(), callee(functionDecl(matchesName("cuda.*")))).bind("cudaCall"), &Callback); Finder.addMatcher(cudaKernelCallExpr().bind("cudaLaunchKernel"), &Callback); @@ -642,8 +672,8 @@ int main(int argc, const char **argv) { Finder.addMatcher(varDecl(isExpansionInMainFile(), hasType(pointsTo(cxxRecordDecl(matchesName("cuda.*"))))).bind("cudaStructVarPtr"), &Callback); Finder.addMatcher(parmVarDecl(isExpansionInMainFile(), hasType(namedDecl(matchesName("cuda.*")))).bind("cudaParamDecl"), &Callback); Finder.addMatcher(parmVarDecl(isExpansionInMainFile(), hasType(pointsTo(namedDecl(matchesName("cuda.*"))))).bind("cudaParamDeclPtr"), &Callback); - Finder.addMatcher(expr(sizeOfExpr(hasArgumentOfType(recordType(hasDeclaration(cxxRecordDecl(matchesName("cuda.*"))))))).bind("cudaStructSizeOf"), &Callback); - Finder.addMatcher(stringLiteral().bind("stringLiteral"), &Callback); + Finder.addMatcher(expr(isExpansionInMainFile(), sizeOfExpr(hasArgumentOfType(recordType(hasDeclaration(cxxRecordDecl(matchesName("cuda.*"))))))).bind("cudaStructSizeOf"), &Callback); + Finder.addMatcher(stringLiteral(isExpansionInMainFile()).bind("stringLiteral"), &Callback); auto action = newFrontendActionFactory(&Finder, &PPCallbacks);