From 434c070ea7ed2bf6c5da134584e31dcb0ac51fa1 Mon Sep 17 00:00:00 2001 From: atimofee Date: Wed, 17 Feb 2016 18:34:58 +0300 Subject: [PATCH] Hipify tool in it current state --- src/Cuda2Hip.cpp | 177 ++++++++++++++++++++++++++++------------------- 1 file changed, 105 insertions(+), 72 deletions(-) diff --git a/src/Cuda2Hip.cpp b/src/Cuda2Hip.cpp index c4c8f58810..fe3f360dce 100644 --- a/src/Cuda2Hip.cpp +++ b/src/Cuda2Hip.cpp @@ -1,29 +1,39 @@ -/* -Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved. +//===---- tools/extra/ToolTemplate.cpp - Template for refactoring tool ----===// +// +// The LLVM Compiler Infrastructure +// +// This file is distributed under the University of Illinois Open Source +// License. See LICENSE.TXT for details. +// +//===----------------------------------------------------------------------===// +// +// This file implements an empty refactoring tool using the clang tooling. +// The goal is to lower the "barrier to entry" for writing refactoring tools. +// +// Usage: +// tool-template ... +// +// Where is a CMake build directory in which a file named +// compile_commands.json exists (enable -DCMAKE_EXPORT_COMPILE_COMMANDS in +// CMake to get this output). +// +// ... specify the paths of files in the CMake source tree. This path +// is looked up in the compile command database. If the path of a file is +// absolute, it needs to point into CMake's source tree. If the path is +// relative, the current working directory needs to be in the CMake source +// tree and the file must be in a subdirectory of the current working +// directory. "./" prefixes in the relative files will be automatically +// removed, but the rest of a relative path must be a suffix of a path in +// the compile command line database. +// +// For example, to use tool-template on all files in a subtree of the +// source tree, use: +// +// /path/in/subtree $ find . -name '*.cpp'| +// xargs tool-template /path/to/build +// +//===----------------------------------------------------------------------===// -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. -*/ -/** - * @file Cuda2Hip.cpp - * - * This file is compiled and linked into clang based hipify tool. - */ #include "clang/ASTMatchers/ASTMatchers.h" #include "clang/ASTMatchers/ASTMatchFinder.h" #include "clang/Basic/SourceManager.h" @@ -65,7 +75,7 @@ namespace { // Error codes and return types: cuda2hipRename["cudaError_t"] = "hipError_t"; - cuda2hipRename["cudaError"] = "hipError_t"; + cuda2hipRename["cudaError"] = "hipError"; cuda2hipRename["cudaSuccess"] = "hipSuccess"; cuda2hipRename["cudaErrorUnknown"] = "hipErrorUnknown"; @@ -138,7 +148,7 @@ namespace { cuda2hipRename["warpSize"] = "hipWarpSize"; - // Events + // Events cuda2hipRename["cudaEvent_t"] = "hipEvent_t"; cuda2hipRename["cudaEventCreate"] = "hipEventCreate"; cuda2hipRename["cudaEventCreateWithFlags"] = "hipEventCreateWithFlags"; @@ -157,7 +167,7 @@ namespace { cuda2hipRename["cudaStreamDefault"] = "hipStreamDefault"; cuda2hipRename["cudaStreamNonBlocking"] = "hipStreamNonBlocking"; - // Other synchronization + // Other synchronization cuda2hipRename["cudaDeviceSynchronize"] = "hipDeviceSynchronize"; cuda2hipRename["cudaThreadSynchronize"] = "hipDeviceSynchronize"; // translate deprecated cudaThreadSynchronize cuda2hipRename["cudaDeviceReset"] = "hipDeviceReset"; @@ -165,7 +175,7 @@ namespace { cuda2hipRename["cudaSetDevice"] = "hipSetDevice"; cuda2hipRename["cudaGetDevice"] = "hipGetDevice"; - // Device + // Device cuda2hipRename["cudaDeviceProp"] = "hipDeviceProp_t"; cuda2hipRename["cudaGetDeviceProperties"] = "hipDeviceGetProperties"; @@ -219,22 +229,22 @@ namespace { DenseMap cuda2hipRename; }; - struct HipifyPPCallbacks : public PPCallbacks, public SourceFileCallbacks { - HipifyPPCallbacks(Replacements * R) - : SeenEnd(false), _sm(nullptr), Replace(R) - { - } - - virtual bool handleBeginSource(CompilerInstance &CI, StringRef Filename) override - { - Preprocessor &PP = CI.getPreprocessor(); - SourceManager & SM = CI.getSourceManager(); - setSourceManager(&SM); - PP.addPPCallbacks(std::unique_ptr(this)); - PP.Retain(); - return true; - } - + struct HipifyPPCallbacks : public PPCallbacks, public SourceFileCallbacks { + HipifyPPCallbacks(Replacements * R) + : SeenEnd(false), _sm(nullptr), Replace(R) + { + } + + virtual bool handleBeginSource(CompilerInstance &CI, StringRef Filename) override + { + Preprocessor &PP = CI.getPreprocessor(); + SourceManager & SM = CI.getSourceManager(); + setSourceManager(&SM); + PP.addPPCallbacks(std::unique_ptr(this)); + PP.Retain(); + return true; + } + virtual void InclusionDirective( SourceLocation hash_loc, const Token &include_token, @@ -255,8 +265,8 @@ namespace { llvm::errs() << "\nWill be replaced with " << repName << "\n"; SourceLocation sl = filename_range.getBegin(); SourceLocation sle = filename_range.getEnd(); - const char* B = _sm->getCharacterData(sl); - const char* E = _sm->getCharacterData(sle); + const char* B = _sm->getCharacterData(sl); + const char* E = _sm->getCharacterData(sle); repName = "<" + repName + ">"; Replacement Rep(*_sm, sl, E - B, repName); Replace->insert(Rep); @@ -290,22 +300,22 @@ namespace { } } } - - void EndOfMainFile() override - { - - } - - bool SeenEnd; - void setSourceManager(SourceManager * sm) { _sm = sm; } - - private: - - SourceManager * _sm; - Replacements * Replace; - struct hipName N; - }; - + + void EndOfMainFile() override + { + + } + + bool SeenEnd; + void setSourceManager(SourceManager * sm) { _sm = sm; } + + private: + + SourceManager * _sm; + Replacements * Replace; + struct hipName N; + }; + class Cuda2HipCallback : public MatchFinder::MatchCallback { public: Cuda2HipCallback(Replacements *Replace) : Replace(Replace) {} @@ -349,7 +359,7 @@ class Cuda2HipCallback : public MatchFinder::MatchCallback { std::string name = kernelDecl->getDeclName().getAsString(); std::string repName = "hipLaunchKernel(HIP_KERNEL_NAME(" + name + "), "; - + const CallExpr * config = launchKernel->getConfig(); llvm::outs() << "\nKernel config arguments:\n"; @@ -359,7 +369,7 @@ class Cuda2HipCallback : public MatchFinder::MatchCallback { if (!isa(arg)) { std::string typeCtor = ""; const ParmVarDecl * pvd = config->getDirectCallee()->getParamDecl(argno); - + SourceLocation sl(arg->getLocStart()); SourceLocation el(arg->getLocEnd()); SourceLocation stop = clang::Lexer::getLocForEndOfToken(el, 0, *SM, DefaultLangOptions); @@ -372,7 +382,7 @@ class Cuda2HipCallback : public MatchFinder::MatchCallback { } else repName += " 0,"; } - + for (unsigned argno = 0; argno < launchKernel->getNumArgs(); argno++) { const Expr * arg = launchKernel->getArg(argno); @@ -409,11 +419,21 @@ class Cuda2HipCallback : public MatchFinder::MatchCallback { } } - if (const DeclRefExpr * cudaEnumConstant = Result.Nodes.getNodeAs("cudaEnumConstant")) + if (const DeclRefExpr * cudaEnumConstantRef = Result.Nodes.getNodeAs("cudaEnumConstantRef")) { - std::string name = cudaEnumConstant->getDecl()->getNameAsString(); + std::string name = cudaEnumConstantRef->getDecl()->getNameAsString(); std::string repName = N.cuda2hipRename[name]; - SourceLocation sl = cudaEnumConstant->getLocStart(); + SourceLocation sl = cudaEnumConstantRef->getLocStart(); + Replacement Rep(*SM, SM->isMacroArgExpansion(sl) ? + SM->getImmediateSpellingLoc(sl) : sl, name.length(), repName); + Replace->insert(Rep); + } + + if (const VarDecl * cudaEnumConstantDecl = Result.Nodes.getNodeAs("cudaEnumConstantDecl")) + { + std::string name = cudaEnumConstantDecl->getType()->getAsTagDecl()->getNameAsString(); + std::string repName = N.cuda2hipRename[name]; + SourceLocation sl = cudaEnumConstantDecl->getLocStart(); Replacement Rep(*SM, SM->isMacroArgExpansion(sl) ? SM->getImmediateSpellingLoc(sl) : sl, name.length(), repName); Replace->insert(Rep); @@ -429,6 +449,17 @@ class Cuda2HipCallback : public MatchFinder::MatchCallback { SM->getImmediateSpellingLoc(sl) : sl, name.length(), repName); Replace->insert(Rep); } + + if (const ParmVarDecl * cudaParamDecl = Result.Nodes.getNodeAs("cudaParamDecl")) + { + QualType QT = cudaParamDecl->getOriginalType(); + std::string name = QT.getAsString(); + std::string repName = N.cuda2hipRename[name]; + SourceLocation sl = cudaParamDecl->getLocStart(); + Replacement Rep(*SM, SM->isMacroArgExpansion(sl) ? + SM->getImmediateSpellingLoc(sl) : sl, name.length(), repName); + Replace->insert(Rep); + } } private: @@ -447,7 +478,7 @@ int main(int argc, const char **argv) { llvm::sys::PrintStackTraceOnErrorSignal(); int Result; - + CommonOptionsParser OptionsParser(argc, argv, ToolTemplateCategory); std::vector savedSources; for (auto I : OptionsParser.getSourcePathList()) @@ -472,8 +503,10 @@ int main(int argc, const char **argv) { Finder.addMatcher(callExpr(isExpansionInMainFile(), callee(functionDecl(matchesName("cuda.*")))).bind("cudaCall"), &Callback); Finder.addMatcher(cudaKernelCallExpr().bind("cudaLaunchKernel"), &Callback); Finder.addMatcher(memberExpr(isExpansionInMainFile(), hasObjectExpression(hasType(cxxRecordDecl(matchesName("__cuda_builtin_"))))).bind("cudaBuiltin"), &Callback); - Finder.addMatcher(declRefExpr(isExpansionInMainFile(), to(enumConstantDecl(matchesName("cuda.*")))).bind("cudaEnumConstant"), &Callback); + Finder.addMatcher(declRefExpr(isExpansionInMainFile(), to(enumConstantDecl(matchesName("cuda.*")))).bind("cudaEnumConstantRef"), &Callback); + Finder.addMatcher(varDecl(isExpansionInMainFile(), hasType(enumDecl(matchesName("cuda.*")))).bind("cudaEnumConstantDecl"), &Callback); Finder.addMatcher(varDecl(isExpansionInMainFile(), hasType(cxxRecordDecl(matchesName("cuda.*")))).bind("cudaStructVar"), &Callback); + Finder.addMatcher(parmVarDecl(isExpansionInMainFile(), hasType(namedDecl(matchesName("cuda.*")))).bind("cudaParamDecl"), &Callback); auto action = newFrontendActionFactory(&Finder, &PPCallbacks); @@ -491,7 +524,7 @@ int main(int argc, const char **argv) { Tool.clearArgumentsAdjusters(); } - + LangOptions DefaultLangOptions; IntrusiveRefCntPtr DiagOpts = new DiagnosticOptions(); TextDiagnosticPrinter DiagnosticPrinter(llvm::errs(), &*DiagOpts);