diff --git a/projects/clr/hipamd/hipify-clang/src/ArgParse.cpp b/projects/clr/hipamd/hipify-clang/src/ArgParse.cpp new file mode 100644 index 0000000000..b27643710f --- /dev/null +++ b/projects/clr/hipamd/hipify-clang/src/ArgParse.cpp @@ -0,0 +1,40 @@ +#include "ArgParse.h" + +cl::OptionCategory ToolTemplateCategory("CUDA to HIP source translator options"); + +cl::opt OutputFilename("o", + cl::desc("Output filename"), + cl::value_desc("filename"), + cl::cat(ToolTemplateCategory)); + +cl::opt Inplace("inplace", + cl::desc("Modify input file inplace, replacing input with hipified output, save backup in .prehip file"), + cl::value_desc("inplace"), + cl::cat(ToolTemplateCategory)); + +cl::opt NoBackup("no-backup", + cl::desc("Don't create a backup file for the hipified source"), + cl::value_desc("no-backup"), + cl::cat(ToolTemplateCategory)); + +cl::opt NoOutput("no-output", + cl::desc("Don't write any translated output to stdout"), + cl::value_desc("no-output"), + cl::cat(ToolTemplateCategory)); + +cl::opt PrintStats("print-stats", + cl::desc("Print translation statistics"), + cl::value_desc("print-stats"), + cl::cat(ToolTemplateCategory)); + +cl::opt OutputStatsFilename("o-stats", + cl::desc("Output filename for statistics"), + cl::value_desc("filename"), + cl::cat(ToolTemplateCategory)); + +cl::opt Examine("examine", + cl::desc("Combines -no-output and -print-stats options"), + cl::value_desc("examine"), + cl::cat(ToolTemplateCategory)); + +cl::extrahelp CommonHelp(ct::CommonOptionsParser::HelpMessage); diff --git a/projects/clr/hipamd/hipify-clang/src/ArgParse.h b/projects/clr/hipamd/hipify-clang/src/ArgParse.h new file mode 100644 index 0000000000..b937a8dd15 --- /dev/null +++ b/projects/clr/hipamd/hipify-clang/src/ArgParse.h @@ -0,0 +1,19 @@ +#pragma once + +#include "clang/Tooling/CommonOptionsParser.h" +#include "llvm/Support/CommandLine.h" + +namespace cl = llvm::cl; +namespace ct = clang::tooling; + +extern cl::OptionCategory ToolTemplateCategory; + +extern cl::opt OutputFilename; +extern cl::opt Inplace; +extern cl::opt NoBackup; +extern cl::opt NoOutput; +extern cl::opt PrintStats; +extern cl::opt OutputStatsFilename; +extern cl::opt Examine; + +extern cl::extrahelp CommonHelp; diff --git a/projects/clr/hipamd/hipify-clang/src/CUDA2HipMap.cpp b/projects/clr/hipamd/hipify-clang/src/CUDA2HipMap.cpp index de6ddb2d74..650fac2079 100644 --- a/projects/clr/hipamd/hipify-clang/src/CUDA2HipMap.cpp +++ b/projects/clr/hipamd/hipify-clang/src/CUDA2HipMap.cpp @@ -2677,7 +2677,6 @@ const std::map& CUDA_RENAMES_MAP() { // First run, so compute the union map. ret = CUDA_IDENTIFIER_MAP; - ret.insert(CUDA_INCLUDE_MAP.begin(), CUDA_INCLUDE_MAP.end()); ret.insert(CUDA_TYPE_NAME_MAP.begin(), CUDA_TYPE_NAME_MAP.end()); return ret; diff --git a/projects/clr/hipamd/hipify-clang/src/Cuda2Hip.cpp b/projects/clr/hipamd/hipify-clang/src/Cuda2Hip.cpp deleted file mode 100644 index a1cf80fde9..0000000000 --- a/projects/clr/hipamd/hipify-clang/src/Cuda2Hip.cpp +++ /dev/null @@ -1,861 +0,0 @@ -/* -Copyright (c) 2015-2017 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. -*/ -/** - * @file Cuda2Hip.cpp - * - * This file is compiled and linked into clang based hipify tool. - */ -#include "clang/ASTMatchers/ASTMatchFinder.h" -#include "clang/ASTMatchers/ASTMatchers.h" -#include "clang/Basic/SourceManager.h" -#include "clang/Frontend/CompilerInstance.h" -#include "clang/Frontend/FrontendActions.h" -#include "clang/Frontend/TextDiagnosticPrinter.h" -#include "clang/Lex/Lexer.h" -#include "clang/Lex/MacroArgs.h" -#include "clang/Lex/MacroInfo.h" -#include "clang/Lex/PPCallbacks.h" -#include "clang/Lex/Preprocessor.h" -#include "clang/Rewrite/Core/Rewriter.h" -#include "clang/Tooling/CommonOptionsParser.h" -#include "clang/Tooling/Refactoring.h" -#include "clang/Tooling/Tooling.h" -#include "llvm/Support/CommandLine.h" -#include "llvm/Support/Debug.h" -#include "llvm/Support/MemoryBuffer.h" -#include "llvm/Support/Signals.h" - -#include -#include -#include -#include -#include -#include -#include - -#include "CUDA2HipMap.h" -#include "LLVMCompat.h" -#include "StringUtils.h" - -using namespace clang; -using namespace clang::ast_matchers; -using namespace clang::tooling; -using namespace llvm; - -#define DEBUG_TYPE "cuda2hip" - - -// Set up the command line options -static cl::OptionCategory ToolTemplateCategory("CUDA to HIP source translator options"); - -static cl::opt OutputFilename("o", - cl::desc("Output filename"), - cl::value_desc("filename"), - cl::cat(ToolTemplateCategory)); - -static cl::opt Inplace("inplace", - cl::desc("Modify input file inplace, replacing input with hipified output, save backup in .prehip file"), - cl::value_desc("inplace"), - cl::cat(ToolTemplateCategory)); - -static cl::opt NoBackup("no-backup", - cl::desc("Don't create a backup file for the hipified source"), - cl::value_desc("no-backup"), - cl::cat(ToolTemplateCategory)); - -static cl::opt NoOutput("no-output", - cl::desc("Don't write any translated output to stdout"), - cl::value_desc("no-output"), - cl::cat(ToolTemplateCategory)); - -static cl::opt PrintStats("print-stats", - cl::desc("Print translation statistics"), - cl::value_desc("print-stats"), - cl::cat(ToolTemplateCategory)); - -static cl::opt OutputStatsFilename("o-stats", - cl::desc("Output filename for statistics"), - cl::value_desc("filename"), - cl::cat(ToolTemplateCategory)); - -static cl::opt Examine("examine", - cl::desc("Combines -no-output and -print-stats options"), - cl::value_desc("examine"), - cl::cat(ToolTemplateCategory)); - -static cl::extrahelp CommonHelp(CommonOptionsParser::HelpMessage); - -class Cuda2Hip { -public: - Cuda2Hip(Replacements& R, const std::string &srcFileName) : - Replace(R), mainFileName(srcFileName) {} - - enum msgTypes { - HIPIFY_ERROR = 0, - HIPIFY_WARNING - }; - - std::string getMsgType(msgTypes type) { - switch (type) { - case HIPIFY_ERROR: return "error"; - default: - case HIPIFY_WARNING: return "warning"; - } - } - -protected: - Replacements& Replace; - std::string mainFileName; - - virtual void insertReplacement(const Replacement &rep, const FullSourceLoc &fullSL) { - llcompat::insertReplacement(Replace, rep); - if (PrintStats) { - rep.getLength(); - Statistics::current().lineTouched(fullSL.getExpansionLineNumber()); - Statistics::current().bytesChanged(rep.getLength()); - } - } - void insertHipHeaders(Cuda2Hip *owner, const SourceManager &SM) { - if (Replace.size() > 0) { - std::string repName = "#include "; - Statistics::current().incrementCounter({repName, ConvTypes::CONV_INCLUDE_CUDA_MAIN_H, ApiTypes::API_RUNTIME}, "#include "); - SourceLocation sl = SM.getLocForStartOfFile(SM.getMainFileID()); - FullSourceLoc fullSL(sl, SM); - Replacement Rep(SM, sl, 0, repName + "\n"); - insertReplacement(Rep, fullSL); - } - } - - void printHipifyMessage(const SourceManager &SM, const SourceLocation &sl, const std::string &message, msgTypes msgType = HIPIFY_WARNING) { - FullSourceLoc fullSL(sl, SM); - llvm::errs() << "[HIPIFY] " << getMsgType(msgType) << ": " << mainFileName << ":" << fullSL.getExpansionLineNumber() << ":" << fullSL.getExpansionColumnNumber() << ": " << message << "\n"; - } - - void processString(StringRef s, SourceManager &SM, SourceLocation start) { - size_t begin = 0; - while ((begin = s.find("cu", begin)) != StringRef::npos) { - const size_t end = s.find_first_of(" ", begin + 4); - StringRef name = s.slice(begin, end); - const auto found = CUDA_RENAMES_MAP().find(name); - if (found != CUDA_RENAMES_MAP().end()) { - StringRef repName = found->second.hipName; - hipCounter counter = {"[string literal]", ConvTypes::CONV_LITERAL, ApiTypes::API_RUNTIME, found->second.unsupported}; - Statistics::current().incrementCounter(counter, name.str()); - if (!counter.unsupported) { - SourceLocation sl = start.getLocWithOffset(begin + 1); - Replacement Rep(SM, sl, name.size(), repName); - FullSourceLoc fullSL(sl, SM); - insertReplacement(Rep, fullSL); - } - } - - if (end == StringRef::npos) { - break; - } - begin = end + 1; - } - } -}; - -class Cuda2HipCallback; - -class HipifyPPCallbacks : public PPCallbacks, public SourceFileCallbacks, public Cuda2Hip { -public: - HipifyPPCallbacks(Replacements& R, const std::string &mainFileName) - : Cuda2Hip(R, mainFileName) {} - - virtual bool handleBeginSource(CompilerInstance &CI -#if LLVM_VERSION_MAJOR <= 4 - , StringRef Filename -#endif - ) override { - Preprocessor &PP = CI.getPreprocessor(); - SourceManager &SM = CI.getSourceManager(); - setSourceManager(&SM); - PP.addPPCallbacks(std::unique_ptr(this)); - setPreprocessor(&PP); - return true; - } - - virtual void handleEndSource() override; - - virtual void InclusionDirective(SourceLocation hash_loc, - const Token &include_token, - StringRef file_name, bool is_angled, - CharSourceRange filename_range, - const FileEntry *file, StringRef search_path, - StringRef relative_path, - const clang::Module *imported) override { - if (!_sm->isWrittenInMainFile(hash_loc) || !is_angled) { - return; // We're looking to rewrite angle-includes in the main file to point to hip. - } - - const auto found = CUDA_INCLUDE_MAP.find(file_name); - if (found == CUDA_INCLUDE_MAP.end()) { - // Not a CUDA include - don't touch it. - return; - } - - Statistics::current().incrementCounter(found->second, file_name.str()); - if (found->second.unsupported) { - // An unsupported CUDA header? Oh dear. Print a warning. - printHipifyMessage(*_sm, hash_loc, "Unsupported CUDA header used: " + file_name.str()); - return; - } - - StringRef repName = found->second.hipName; - DEBUG(dbgs() << "Include file found: " << file_name << "\n" - << "SourceLocation: " - << filename_range.getBegin().printToString(*_sm) << "\n" - << "Will 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); - SmallString<128> tmpData; - Replacement Rep(*_sm, sl, E - B, Twine("<" + repName + ">").toStringRef(tmpData)); - FullSourceLoc fullSL(sl, *_sm); - insertReplacement(Rep, fullSL); - } - - /** - * Look at, and consider altering, a given token. - * - * If it's not a CUDA identifier, nothing happens. - * If it's an unsupported CUDA identifier, a warning is emitted. - * Otherwise, the source file is updated with the corresponding hipification. - */ - void RewriteToken(Token t) { - // String literals containing CUDA references need fixing... - if (t.is(tok::string_literal)) { - StringRef s(t.getLiteralData(), t.getLength()); - processString(unquoteStr(s), *_sm, t.getLocation()); - return; - } else if (!t.isAnyIdentifier()) { - // If it's neither a string nor an identifier, we don't care. - return; - } - - StringRef name = t.getIdentifierInfo()->getName(); - const auto found = CUDA_RENAMES_MAP().find(name); - if (found == CUDA_RENAMES_MAP().end()) { - // So it's an identifier, but not CUDA? Boring. - return; - } - Statistics::current().incrementCounter(found->second, name.str()); - - SourceLocation sl = t.getLocation(); - if (found->second.unsupported) { - // An unsupported identifier? Curses! Warn the user. - printHipifyMessage(*_sm, sl, "Unsupported CUDA identifier used: " + name.str()); - return; - } - - StringRef repName = found->second.hipName; - Replacement Rep(*_sm, sl, name.size(), repName); - FullSourceLoc fullSL(sl, *_sm); - insertReplacement(Rep, fullSL); - } - - virtual void MacroDefined(const Token &MacroNameTok, - const MacroDirective *MD) override { - if (!_sm->isWrittenInMainFile(MD->getLocation()) || - MD->getKind() != MacroDirective::MD_Define) { - return; - } - - for (auto T : MD->getMacroInfo()->tokens()) { - RewriteToken(T); - } - } - - virtual void MacroExpands(const Token &MacroNameTok, - const MacroDefinition &MD, SourceRange Range, - const MacroArgs *Args) override { - - if (!_sm->isWrittenInMainFile(MacroNameTok.getLocation())) { - return; // Macros in headers are not our concern. - } - - // Is the macro itself a CUDA identifier? If so, rewrite it - RewriteToken(MacroNameTok); - - // If it's a macro with arguments, rewrite all the arguments as hip, too. - for (unsigned int i = 0; Args && i < MD.getMacroInfo()->GET_NUM_ARGS(); i++) { - std::vector toks; - // Code below is a kind of stolen from 'MacroArgs::getPreExpArgument' - // to workaround the 'const' MacroArgs passed into this hook. - const Token *start = Args->getUnexpArgument(i); - size_t len = Args->getArgLength(start) + 1; - llcompat::EnterPreprocessorTokenStream(*_pp, start, len, false); - - do { - toks.push_back(Token()); - Token &tk = toks.back(); - _pp->Lex(tk); - } while (toks.back().isNot(tok::eof)); - - _pp->RemoveTopOfLexerStack(); - // end of stolen code - - for (auto tok : toks) { - RewriteToken(tok); - } - } - } - - void EndOfMainFile() override {} - - void setSourceManager(SourceManager *sm) { _sm = sm; } - void setPreprocessor(Preprocessor *pp) { _pp = pp; } - void setMatch(Cuda2HipCallback *match) { Match = match; } - -private: - SourceManager *_sm = nullptr; - Preprocessor *_pp = nullptr; - Cuda2HipCallback *Match = nullptr; -}; - -class Cuda2HipCallback : public MatchFinder::MatchCallback, public Cuda2Hip { -private: - bool cudaCall(const MatchFinder::MatchResult &Result) { - const CallExpr *call = Result.Nodes.getNodeAs("cudaCall"); - if (!call) { - return false; // Another handler will do it. - } - - const FunctionDecl *funcDcl = call->getDirectCallee(); - std::string name = funcDcl->getDeclName().getAsString(); - SourceManager *SM = Result.SourceManager; - SourceLocation sl = call->getLocStart(); - - // TODO: Make a lookup table just for functions to improve performance. - const auto found = CUDA_IDENTIFIER_MAP.find(name); - if (found == CUDA_IDENTIFIER_MAP.end()) { - std::string msg = "the following reference is not handled: '" + name + "' [function call]."; - printHipifyMessage(*SM, sl, msg); - return true; - } - - const hipCounter& hipCtr = found->second; - Statistics::current().incrementCounter(hipCtr, name); - - if (hipCtr.unsupported) { - return true; // Silently fail when you find an unsupported member. - // TODO: Print a warning with the diagnostics API? - } - - size_t length = name.size(); - Replacement Rep(*SM, sl, length, hipCtr.hipName); - FullSourceLoc fullSL(sl, *SM); - insertReplacement(Rep, fullSL); - - return true; - } - - SourceRange getReadRange(clang::SourceManager &SM, const SourceRange &exprRange) { - SourceLocation begin = exprRange.getBegin(); - SourceLocation end = exprRange.getEnd(); - - bool beginSafe = !SM.isMacroBodyExpansion(begin) || Lexer::isAtStartOfMacroExpansion(begin, SM, LangOptions{}); - bool endSafe = !SM.isMacroBodyExpansion(end) || Lexer::isAtEndOfMacroExpansion(end, SM, LangOptions{}); - - if (beginSafe && endSafe) { - return {SM.getFileLoc(begin), SM.getFileLoc(end)}; - } else { - return {SM.getSpellingLoc(begin), SM.getSpellingLoc(end)}; - } - } - - SourceRange getWriteRange(clang::SourceManager &SM, const SourceRange &exprRange) { - SourceLocation begin = exprRange.getBegin(); - SourceLocation end = exprRange.getEnd(); - - // If the range is contained within a macro, update the macro definition. - // Otherwise, use the file location and hope for the best. - if (!SM.isMacroBodyExpansion(begin) || !SM.isMacroBodyExpansion(end)) { - return {SM.getFileLoc(begin), SM.getFileLoc(end)}; - } - - return {SM.getSpellingLoc(begin), SM.getSpellingLoc(end)}; - } - - StringRef readSourceText(clang::SourceManager& SM, const SourceRange& exprRange) { - return Lexer::getSourceText(CharSourceRange::getTokenRange(getReadRange(SM, exprRange)), SM, LangOptions(), nullptr); - } - - /** - * Get a string representation of the expression `arg`, unless it's a defaulting function - * call argument, in which case get a 0. Used for building argument lists to kernel calls. - */ - std::string stringifyZeroDefaultedArg(SourceManager& SM, const Expr* arg) { - if (isa(arg)) { - return "0"; - } else { - return readSourceText(SM, arg->getSourceRange()); - } - } - - bool cudaLaunchKernel(const MatchFinder::MatchResult &Result) { - StringRef refName = "cudaLaunchKernel"; - if (const CUDAKernelCallExpr *launchKernel = Result.Nodes.getNodeAs(refName)) { - SmallString<40> XStr; - raw_svector_ostream OS(XStr); - - LangOptions DefaultLangOptions; - SourceManager *SM = Result.SourceManager; - - const Expr& calleeExpr = *(launchKernel->getCallee()); - OS << "hipLaunchKernelGGL(" << readSourceText(*SM, calleeExpr.getSourceRange()) << ", "; - - // Next up are the four kernel configuration parameters, the last two of which are optional and default to zero. - const CallExpr& config = *(launchKernel->getConfig()); - - // Copy the two dimensional arguments verbatim. - OS << "dim3(" << readSourceText(*SM, config.getArg(0)->getSourceRange()) << "), "; - OS << "dim3(" << readSourceText(*SM, config.getArg(1)->getSourceRange()) << "), "; - - // The stream/memory arguments default to zero if omitted. - OS << stringifyZeroDefaultedArg(*SM, config.getArg(2)) << ", "; - OS << stringifyZeroDefaultedArg(*SM, config.getArg(3)); - - // If there are ordinary arguments to the kernel, just copy them verbatim into our new call. - int numArgs = launchKernel->getNumArgs(); - if (numArgs > 0) { - OS << ", "; - - // Start of the first argument. - SourceLocation argStart = launchKernel->getArg(0)->getLocStart(); - - // End of the last argument. - SourceLocation argEnd = launchKernel->getArg(numArgs - 1)->getLocEnd(); - - OS << readSourceText(*SM, {argStart, argEnd}); - } - - OS << ")"; - - SourceRange replacementRange = getWriteRange(*SM, {launchKernel->getLocStart(), launchKernel->getLocEnd()}); - SourceLocation launchStart = replacementRange.getBegin(); - SourceLocation launchEnd = replacementRange.getEnd(); - - size_t length = SM->getCharacterData(Lexer::getLocForEndOfToken( - launchEnd, 0, *SM, DefaultLangOptions)) - - SM->getCharacterData(launchStart); - - Replacement Rep(*SM, launchStart, length, OS.str()); - FullSourceLoc fullSL(launchStart, *SM); - insertReplacement(Rep, fullSL); - hipCounter counter = {"hipLaunchKernelGGL", ConvTypes::CONV_KERN, ApiTypes::API_RUNTIME}; - Statistics::current().incrementCounter(counter, refName.str()); - return true; - } - return false; - } - - bool cudaBuiltin(const MatchFinder::MatchResult &Result) { - if (const MemberExpr *threadIdx = Result.Nodes.getNodeAs("cudaBuiltin")) { - if (const OpaqueValueExpr *refBase = - dyn_cast(threadIdx->getBase())) { - if (const DeclRefExpr *declRef = - dyn_cast(refBase->getSourceExpr())) { - SourceLocation sl = threadIdx->getLocStart(); - SourceManager *SM = Result.SourceManager; - StringRef name = declRef->getDecl()->getName(); - StringRef memberName = threadIdx->getMemberDecl()->getName(); - size_t pos = memberName.find_first_not_of("__fetch_builtin_"); - memberName = memberName.slice(pos, memberName.size()); - SmallString<128> tmpData; - name = Twine(name + "." + memberName).toStringRef(tmpData); - - // TODO: Make a lookup table just for builtins to improve performance. - const auto found = CUDA_IDENTIFIER_MAP.find(name); - if (found != CUDA_IDENTIFIER_MAP.end()) { - Statistics::current().incrementCounter(found->second, name.str()); - 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.str() + "' [builtin]."; - printHipifyMessage(*SM, sl, msg); - } - } - } - return true; - } - return false; - } - - bool cudaEnumConstantRef(const MatchFinder::MatchResult &Result) { - if (const DeclRefExpr *enumConstantRef = Result.Nodes.getNodeAs("cudaEnumConstantRef")) { - StringRef name = enumConstantRef->getDecl()->getName(); - SourceLocation sl = enumConstantRef->getLocStart(); - SourceManager *SM = Result.SourceManager; - - // TODO: Make a lookup table just for enum values to improve performance. - const auto found = CUDA_IDENTIFIER_MAP.find(name); - if (found != CUDA_IDENTIFIER_MAP.end()) { - Statistics::current().incrementCounter(found->second, name.str()); - 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.str() + "' [enum constant ref]."; - printHipifyMessage(*SM, sl, msg); - } - return true; - } - return false; - } - - bool cudaType(const MatchFinder::MatchResult& Result) { - const clang::TypeLoc* ret = Result.Nodes.getNodeAs("cudaType"); - if (!ret) { - return false; - } - - // Ignore qualifiers - they don't alter our decision to rename. - clang::UnqualTypeLoc tl = ret->getUnqualifiedLoc(); - const Type& typeObject = *(tl.getTypePtr()); - - std::string typeName = tl.getType().getAsString(); - - // Irritatingly, enum/struct types are identified as `enum/struct `, and unlike most compound - // types (such as pointers or references), there isn't another type node inside. So we have - // to make do with what we've got. There's probably a better way of doing this... - if (typeObject.isEnumeralType()) { - removePrefixIfPresent(typeName, "enum "); - } - if (typeObject.isStructureType()) { - removePrefixIfPresent(typeName, "struct "); - } - - // Do we have a replacement for this type? - const auto found = CUDA_TYPE_NAME_MAP.find(typeName); - if (found == CUDA_TYPE_NAME_MAP.end()) { - return false; - } - - SourceManager &SM = *(Result.SourceManager); - - // Start of the type expression to replace. - SourceLocation sl = tl.getBeginLoc(); - - const hipCounter& hipCtr = found->second; - if (hipCtr.unsupported) { - printHipifyMessage(SM, sl, "Unsupported CUDA '" + typeName); - return false; - } - - // Apply the rename! - Replacement Rep(SM, sl, typeName.size(), hipCtr.hipName); - FullSourceLoc fullSL(sl, SM); - insertReplacement(Rep, fullSL); - - return true; - } - - bool cudaSharedIncompleteArrayVar(const MatchFinder::MatchResult &Result) { - StringRef refName = "cudaSharedIncompleteArrayVar"; - if (const VarDecl *sharedVar = Result.Nodes.getNodeAs(refName)) { - // Example: extern __shared__ uint sRadix1[]; - if (sharedVar->hasExternalFormalLinkage()) { - QualType QT = sharedVar->getType(); - std::string typeName; - if (QT->isIncompleteArrayType()) { - const ArrayType *AT = QT.getTypePtr()->getAsArrayTypeUnsafe(); - QT = AT->getElementType(); - if (QT.getTypePtr()->isBuiltinType()) { - QT = QT.getCanonicalType(); - const BuiltinType *BT = dyn_cast(QT); - if (BT) { - LangOptions LO; - LO.CUDA = true; - PrintingPolicy policy(LO); - typeName = BT->getName(policy); - } - } else { - typeName = QT.getAsString(); - } - } - if (!typeName.empty()) { - SourceLocation slStart = sharedVar->getLocStart(); - SourceLocation slEnd = sharedVar->getLocEnd(); - SourceManager *SM = Result.SourceManager; - size_t repLength = SM->getCharacterData(slEnd) - SM->getCharacterData(slStart) + 1; - std::string varName = sharedVar->getNameAsString(); - std::string repName = "HIP_DYNAMIC_SHARED(" + typeName + ", " + varName + ")"; - Replacement Rep(*SM, slStart, repLength, repName); - FullSourceLoc fullSL(slStart, *SM); - insertReplacement(Rep, fullSL); - hipCounter counter = { "HIP_DYNAMIC_SHARED", ConvTypes::CONV_MEM, ApiTypes::API_RUNTIME }; - Statistics::current().incrementCounter(counter, refName.str()); - } - } - return true; - } - return false; - } - - bool stringLiteral(const MatchFinder::MatchResult &Result) { - if (const clang::StringLiteral *sLiteral = Result.Nodes.getNodeAs("stringLiteral")) { - if (sLiteral->getCharByteWidth() == 1) { - StringRef s = sLiteral->getString(); - SourceManager *SM = Result.SourceManager; - processString(s, *SM, sLiteral->getLocStart()); - } - return true; - } - return false; - } - -public: - Cuda2HipCallback(Replacements& Replace, ast_matchers::MatchFinder *parent, HipifyPPCallbacks *PPCallbacks, const std::string &mainFileName) - : Cuda2Hip(Replace, mainFileName), owner(parent), PP(PPCallbacks) { - PP->setMatch(this); - } - - void run(const MatchFinder::MatchResult &Result) override { - if (cudaType(Result)) return; - if (cudaCall(Result)) return; - if (cudaBuiltin(Result)) return; - if (cudaEnumConstantRef(Result)) return; - if (cudaLaunchKernel(Result)) return; - if (cudaSharedIncompleteArrayVar(Result)) return; - if (stringLiteral(Result)) return; - } - -private: - ast_matchers::MatchFinder *owner; - HipifyPPCallbacks *PP; -}; - -void HipifyPPCallbacks::handleEndSource() { - insertHipHeaders(Match, *_sm); -} - -void addAllMatchers(ast_matchers::MatchFinder &Finder, Cuda2HipCallback *Callback) { - // Rewrite CUDA api calls to hip ones. - Finder.addMatcher( - callExpr( - isExpansionInMainFile(), - callee( - functionDecl( - matchesName("cu.*"), - unless( - // Clang generates structs with functions on them to represent things like - // threadIdx.x. We have other logic to handle those builtins directly, so - // we need to suppress the call-handling. - // We can't handle those directly in the call-handler without special-casing - // it unpleasantly, since the names of the functions are unique only per-struct. - matchesName("__fetch_builtin.*") - ) - ) - ) - ).bind("cudaCall"), - Callback - ); - - // Rewrite all references to CUDA types to their corresponding hip types. - Finder.addMatcher( - typeLoc( - isExpansionInMainFile() - ).bind("cudaType"), - Callback - ); - - // Replace references to CUDA names in string literals with the equivalent hip names. - Finder.addMatcher(stringLiteral(isExpansionInMainFile()).bind("stringLiteral"), Callback); - - // Replace the <<<...>>> language extension with a hip kernel launch - Finder.addMatcher(cudaKernelCallExpr(isExpansionInMainFile()).bind("cudaLaunchKernel"), Callback); - - // Replace cuda builtins. - Finder.addMatcher( - memberExpr( - isExpansionInMainFile(), - hasObjectExpression( - hasType( - cxxRecordDecl( - matchesName("__cuda_builtin_") - ) - ) - ) - ).bind("cudaBuiltin"), - Callback - ); - - // Map CUDA enum _values_ to their hip equivalents. - Finder.addMatcher( - declRefExpr( - isExpansionInMainFile(), - to( - enumConstantDecl( - matchesName("cu.*|CU.*") - ) - ) - ).bind("cudaEnumConstantRef"), - Callback - ); - - Finder.addMatcher( - varDecl( - isExpansionInMainFile(), - allOf( - hasAttr(attr::CUDAShared), - hasType(incompleteArrayType()) - ) - ).bind("cudaSharedIncompleteArrayVar"), - Callback - ); -} - -void copyFile(const std::string& src, const std::string& dst) { - std::ifstream source(src, std::ios::binary); - std::ofstream dest(dst, std::ios::binary); - dest << source.rdbuf(); -} - -int main(int argc, const char **argv) { - llcompat::PrintStackTraceOnErrorSignal(); - - CommonOptionsParser OptionsParser(argc, argv, ToolTemplateCategory, llvm::cl::OneOrMore); - std::vector fileSources = OptionsParser.getSourcePathList(); - std::string dst = OutputFilename; - if (!dst.empty() && fileSources.size() > 1) { - llvm::errs() << "[HIPIFY] conflict: -o and multiple source files are specified.\n"; - return 1; - } - - if (NoOutput) { - if (Inplace) { - llvm::errs() << "[HIPIFY] conflict: both -no-output and -inplace options are specified.\n"; - return 1; - } - if (!dst.empty()) { - llvm::errs() << "[HIPIFY] conflict: both -no-output and -o options are specified.\n"; - return 1; - } - } - - if (Examine) { - NoOutput = PrintStats = true; - } - - int Result = 0; - - // Arguments for the Statistics print routines. - std::unique_ptr csv = nullptr; - llvm::raw_ostream* statPrint = nullptr; - if (!OutputStatsFilename.empty()) { - csv = std::unique_ptr(new std::ofstream(OutputStatsFilename, std::ios_base::trunc)); - } - if (PrintStats) { - statPrint = &llvm::errs(); - } - - for (const auto & src : fileSources) { - if (dst.empty()) { - if (Inplace) { - dst = src; - } else { - dst = src + ".hip"; - } - } else if (Inplace) { - llvm::errs() << "[HIPIFY] conflict: both -o and -inplace options are specified.\n"; - return 1; - } - - std::string tmpFile = src + ".hipify-tmp"; - - // Create a copy of the file to work on. When we're done, we'll move this onto the - // output (which may mean overwriting the input, if we're in-place). - // Should we fail for some reason, we'll just leak this file and not corrupt the input. - copyFile(src, tmpFile); - - // Initialise the statistics counters for this file. - Statistics::setActive(src); - - // RefactoringTool operates on the file in-place. Giving it the output path is no good, - // because that'll break relative includes, and we don't want to overwrite the input file. - // So what we do is operate on a copy, which we then move to the output. - RefactoringTool Tool(OptionsParser.getCompilations(), tmpFile); - ast_matchers::MatchFinder Finder; - - // The Replacements to apply to the file `src`. - Replacements& replacementsToUse = llcompat::getReplacements(Tool, tmpFile); - HipifyPPCallbacks* PPCallbacks = new HipifyPPCallbacks(replacementsToUse, tmpFile); - Cuda2HipCallback Callback(replacementsToUse, &Finder, PPCallbacks, tmpFile); - - addAllMatchers(Finder, &Callback); - - auto action = newFrontendActionFactory(&Finder, PPCallbacks); - - Tool.appendArgumentsAdjuster(getInsertArgumentAdjuster("--cuda-host-only", ArgumentInsertPosition::BEGIN)); - - // Ensure at least c++11 is used. - Tool.appendArgumentsAdjuster(getInsertArgumentAdjuster("-std=c++11", ArgumentInsertPosition::BEGIN)); -#if defined(HIPIFY_CLANG_RES) - Tool.appendArgumentsAdjuster(getInsertArgumentAdjuster("-resource-dir=" HIPIFY_CLANG_RES)); -#endif - Tool.appendArgumentsAdjuster(getClangSyntaxOnlyAdjuster()); - Result += Tool.run(action.get()); - Tool.clearArgumentsAdjusters(); - - LangOptions DefaultLangOptions; - IntrusiveRefCntPtr DiagOpts = new DiagnosticOptions(); - TextDiagnosticPrinter DiagnosticPrinter(llvm::errs(), &*DiagOpts); - DiagnosticsEngine Diagnostics(IntrusiveRefCntPtr(new DiagnosticIDs()), &*DiagOpts, &DiagnosticPrinter, false); - - SourceManager SM(Diagnostics, Tool.getFiles()); - - Rewriter Rewrite(SM, DefaultLangOptions); - if (!Tool.applyAllReplacements(Rewrite)) { - DEBUG(dbgs() << "Skipped some replacements.\n"); - } - - // Either move the tmpfile to the output, or remove it. - if (!NoOutput) { - Result += Rewrite.overwriteChangedFiles(); - rename(tmpFile.c_str(), dst.c_str()); - } else { - remove(tmpFile.c_str()); - } - - Statistics::current().markCompletion(); - Statistics::current().print(csv.get(), statPrint); - - dst.clear(); - } - - if (fileSources.size() > 1) { - Statistics::printAggregate(csv.get(), statPrint); - } - - return Result; -} diff --git a/projects/clr/hipamd/hipify-clang/src/HipifyAction.cpp b/projects/clr/hipamd/hipify-clang/src/HipifyAction.cpp new file mode 100644 index 0000000000..7608e4040b --- /dev/null +++ b/projects/clr/hipamd/hipify-clang/src/HipifyAction.cpp @@ -0,0 +1,457 @@ +#include "HipifyAction.h" + +#include + +#include "clang/Basic/SourceLocation.h" +#include "clang/Frontend/CompilerInstance.h" +#include "clang/ASTMatchers/ASTMatchFinder.h" +#include "clang/ASTMatchers/ASTMatchers.h" + +#include "LLVMCompat.h" +#include "CUDA2HipMap.h" +#include "StringUtils.h" +#include "ArgParse.h" + +namespace ct = clang::tooling; +namespace mat = clang::ast_matchers; + +void HipifyAction::RewriteString(StringRef s, clang::SourceLocation start) { + clang::SourceManager& SM = getCompilerInstance().getSourceManager(); + + size_t begin = 0; + while ((begin = s.find("cu", begin)) != StringRef::npos) { + const size_t end = s.find_first_of(" ", begin + 4); + StringRef name = s.slice(begin, end); + const auto found = CUDA_RENAMES_MAP().find(name); + if (found != CUDA_RENAMES_MAP().end()) { + StringRef repName = found->second.hipName; + hipCounter counter = {"[string literal]", ConvTypes::CONV_LITERAL, ApiTypes::API_RUNTIME, found->second.unsupported}; + Statistics::current().incrementCounter(counter, name.str()); + + if (!counter.unsupported) { + clang::SourceLocation sl = start.getLocWithOffset(begin + 1); + ct::Replacement Rep(SM, sl, name.size(), repName); + clang::FullSourceLoc fullSL(sl, SM); + insertReplacement(Rep, fullSL); + } + } + + if (end == StringRef::npos) { + break; + } + + begin = end + 1; + } +} + +/** + * Look at, and consider altering, a given token. + * + * If it's not a CUDA identifier, nothing happens. + * If it's an unsupported CUDA identifier, a warning is emitted. + * Otherwise, the source file is updated with the corresponding hipification. + */ +void HipifyAction::RewriteToken(const clang::Token& t) { + clang::SourceManager& SM = getCompilerInstance().getSourceManager(); + + // String literals containing CUDA references need fixing... + if (t.is(clang::tok::string_literal)) { + StringRef s(t.getLiteralData(), t.getLength()); + RewriteString(unquoteStr(s), t.getLocation()); + return; + } else if (!t.isAnyIdentifier()) { + // If it's neither a string nor an identifier, we don't care. + return; + } + + StringRef name = t.getRawIdentifier(); + const auto found = CUDA_RENAMES_MAP().find(name); + if (found == CUDA_RENAMES_MAP().end()) { + // So it's an identifier, but not CUDA? Boring. + return; + } + + Statistics::current().incrementCounter(found->second, name.str()); + + clang::SourceLocation sl = t.getLocation(); + if (found->second.unsupported) { + // An unsupported identifier? Curses! Warn the user. + llvm::errs() << "Unsupported CUDA identifier used: " + name.str() << "\n"; + return; + } + + StringRef repName = found->second.hipName; + ct::Replacement Rep(SM, sl, name.size(), repName); + clang::FullSourceLoc fullSL(sl, SM); + insertReplacement(Rep, fullSL); +} + +namespace { + +clang::SourceRange getReadRange(clang::SourceManager& SM, const clang::SourceRange& exprRange) { + clang::SourceLocation begin = exprRange.getBegin(); + clang::SourceLocation end = exprRange.getEnd(); + + bool beginSafe = !SM.isMacroBodyExpansion(begin) || clang::Lexer::isAtStartOfMacroExpansion(begin, SM, clang::LangOptions{}); + bool endSafe = !SM.isMacroBodyExpansion(end) || clang::Lexer::isAtEndOfMacroExpansion(end, SM, clang::LangOptions{}); + + if (beginSafe && endSafe) { + return {SM.getFileLoc(begin), SM.getFileLoc(end)}; + } else { + return {SM.getSpellingLoc(begin), SM.getSpellingLoc(end)}; + } +} + + +clang::SourceRange getWriteRange(clang::SourceManager& SM, const clang::SourceRange& exprRange) { + clang::SourceLocation begin = exprRange.getBegin(); + clang::SourceLocation end = exprRange.getEnd(); + + // If the range is contained within a macro, update the macro definition. + // Otherwise, use the file location and hope for the best. + if (!SM.isMacroBodyExpansion(begin) || !SM.isMacroBodyExpansion(end)) { + return {SM.getFileLoc(begin), SM.getFileLoc(end)}; + } + + return {SM.getSpellingLoc(begin), SM.getSpellingLoc(end)}; +} + + +StringRef readSourceText(clang::SourceManager& SM, const clang::SourceRange& exprRange) { + return clang::Lexer::getSourceText(clang::CharSourceRange::getTokenRange(getReadRange(SM, exprRange)), SM, clang::LangOptions(), nullptr); +} + +/** + * Get a string representation of the expression `arg`, unless it's a defaulting function + * call argument, in which case get a 0. Used for building argument lists to kernel calls. + */ +std::string stringifyZeroDefaultedArg(clang::SourceManager& SM, const clang::Expr* arg) { + if (clang::isa(arg)) { + return "0"; + } else { + return readSourceText(SM, arg->getSourceRange()); + } +} + +} // anonymous namespace + + +void HipifyAction::InclusionDirective(clang::SourceLocation hash_loc, + const clang::Token&, + StringRef file_name, + bool is_angled, + clang::CharSourceRange filename_range, + const clang::FileEntry*, StringRef, + StringRef, const clang::Module*) { + clang::SourceManager& SM = getCompilerInstance().getSourceManager(); + if (!SM.isWrittenInMainFile(hash_loc)) { + return; + } + + const auto found = CUDA_INCLUDE_MAP.find(file_name); + if (found == CUDA_INCLUDE_MAP.end()) { + // Not a CUDA include - don't touch it. + return; + } + + // Special-casing to avoid duplication of the hip_runtime include. + if (found->second.hipName == "hip/hip_runtime.h") { + if (insertedRuntimeHeader) { + return; + } + + insertedRuntimeHeader = true; + } + + Statistics::current().incrementCounter(found->second, file_name.str()); + + clang::SourceLocation sl = filename_range.getBegin(); + if (found->second.unsupported) { + // An unsupported CUDA header? Oh dear. Print a warning. + clang::DiagnosticsEngine& DE = getCompilerInstance().getDiagnostics(); + DE.Report(sl, DE.getCustomDiagID(clang::DiagnosticsEngine::Warning, "Unsupported CUDA header")); + return; + } + + const char *B = SM.getCharacterData(sl); + const char *E = SM.getCharacterData(filename_range.getEnd()); + clang::SmallString<128> includeBuffer; + clang::StringRef newInclude; + + // Keep the same include type that the user gave. + if (is_angled) { + newInclude = llvm::Twine("<" + found->second.hipName + ">").toStringRef(includeBuffer); + } else { + newInclude = llvm::Twine("\"" + found->second.hipName + "\"").toStringRef(includeBuffer); + } + + ct::Replacement Rep(SM, sl, E - B, newInclude); + insertReplacement(Rep, clang::FullSourceLoc{sl, SM}); +} + + +bool HipifyAction::cudaLaunchKernel(const clang::ast_matchers::MatchFinder::MatchResult& Result) { + StringRef refName = "cudaLaunchKernel"; + + const auto* launchKernel = Result.Nodes.getNodeAs(refName); + if (!launchKernel) { + return false; + } + + clang::SmallString<40> XStr; + llvm::raw_svector_ostream OS(XStr); + + clang::LangOptions DefaultLangOptions; + clang::SourceManager* SM = Result.SourceManager; + + const clang::Expr& calleeExpr = *(launchKernel->getCallee()); + OS << "hipLaunchKernelGGL(" << readSourceText(*SM, calleeExpr.getSourceRange()) << ", "; + + // Next up are the four kernel configuration parameters, the last two of which are optional and default to zero. + const clang::CallExpr& config = *(launchKernel->getConfig()); + + // Copy the two dimensional arguments verbatim. + OS << "dim3(" << readSourceText(*SM, config.getArg(0)->getSourceRange()) << "), "; + OS << "dim3(" << readSourceText(*SM, config.getArg(1)->getSourceRange()) << "), "; + + // The stream/memory arguments default to zero if omitted. + OS << stringifyZeroDefaultedArg(*SM, config.getArg(2)) << ", "; + OS << stringifyZeroDefaultedArg(*SM, config.getArg(3)); + + // If there are ordinary arguments to the kernel, just copy them verbatim into our new call. + int numArgs = launchKernel->getNumArgs(); + if (numArgs > 0) { + OS << ", "; + + // Start of the first argument. + clang::SourceLocation argStart = launchKernel->getArg(0)->getLocStart(); + + // End of the last argument. + clang::SourceLocation argEnd = launchKernel->getArg(numArgs - 1)->getLocEnd(); + + OS << readSourceText(*SM, {argStart, argEnd}); + } + + OS << ")"; + + clang::SourceRange replacementRange = getWriteRange(*SM, {launchKernel->getLocStart(), launchKernel->getLocEnd()}); + clang::SourceLocation launchStart = replacementRange.getBegin(); + clang::SourceLocation launchEnd = replacementRange.getEnd(); + + size_t length = SM->getCharacterData(clang::Lexer::getLocForEndOfToken(launchEnd, 0, *SM, DefaultLangOptions)) - SM->getCharacterData(launchStart); + + ct::Replacement Rep(*SM, launchStart, length, OS.str()); + clang::FullSourceLoc fullSL(launchStart, *SM); + insertReplacement(Rep, fullSL); + hipCounter counter = {"hipLaunchKernelGGL", ConvTypes::CONV_KERN, ApiTypes::API_RUNTIME}; + Statistics::current().incrementCounter(counter, refName.str()); + + return true; +} + +bool HipifyAction::cudaBuiltin(const clang::ast_matchers::MatchFinder::MatchResult& Result) { + const clang::MemberExpr* threadIdx = Result.Nodes.getNodeAs("cudaBuiltin"); + if (!threadIdx) { + return false; + } + + const clang::OpaqueValueExpr* refBase = clang::dyn_cast(threadIdx->getBase()); + if (!refBase) { + return false; + } + + const clang::DeclRefExpr* declRef = clang::dyn_cast(refBase->getSourceExpr()); + if (!declRef) { + return false; + } + + clang::SourceLocation sl = threadIdx->getLocStart(); + clang::SourceManager* SM = Result.SourceManager; + StringRef name = declRef->getDecl()->getName(); + StringRef memberName = threadIdx->getMemberDecl()->getName(); + size_t pos = memberName.find_first_not_of("__fetch_builtin_"); + memberName = memberName.slice(pos, memberName.size()); + clang::SmallString<128> tmpData; + name = clang::Twine(name + "." + memberName).toStringRef(tmpData); + + const auto found = CUDA_IDENTIFIER_MAP.find(name); + if (found != CUDA_IDENTIFIER_MAP.end()) { + Statistics::current().incrementCounter(found->second, name.str()); + if (!found->second.unsupported) { + StringRef repName = found->second.hipName; + ct::Replacement Rep(*SM, sl, name.size(), repName); + clang::FullSourceLoc fullSL(sl, *SM); + insertReplacement(Rep, fullSL); + } + } else { + std::string msg = "the following reference is not handled: '" + name.str() + "' [builtin]."; + llvm::errs() << msg << "\n"; + } + + return true; +} + +bool HipifyAction::cudaSharedIncompleteArrayVar(const clang::ast_matchers::MatchFinder::MatchResult& Result) { + StringRef refName = "cudaSharedIncompleteArrayVar"; + auto* sharedVar = Result.Nodes.getNodeAs(refName); + if (!sharedVar) { + return false; + } + + // Example: extern __shared__ uint sRadix1[]; + if (!sharedVar->hasExternalFormalLinkage()) { + return false; + } + + clang::QualType QT = sharedVar->getType(); + std::string typeName; + if (QT->isIncompleteArrayType()) { + const clang::ArrayType* AT = QT.getTypePtr()->getAsArrayTypeUnsafe(); + QT = AT->getElementType(); + if (QT.getTypePtr()->isBuiltinType()) { + QT = QT.getCanonicalType(); + const auto* BT = clang::dyn_cast(QT); + if (BT) { + clang::LangOptions LO; + LO.CUDA = true; + clang::PrintingPolicy policy(LO); + typeName = BT->getName(policy); + } + } else { + typeName = QT.getAsString(); + } + } + + if (!typeName.empty()) { + clang::SourceLocation slStart = sharedVar->getLocStart(); + clang::SourceLocation slEnd = sharedVar->getLocEnd(); + clang::SourceManager* SM = Result.SourceManager; + size_t repLength = SM->getCharacterData(slEnd) - SM->getCharacterData(slStart) + 1; + std::string varName = sharedVar->getNameAsString(); + std::string repName = "HIP_DYNAMIC_SHARED(" + typeName + ", " + varName + ")"; + ct::Replacement Rep(*SM, slStart, repLength, repName); + clang::FullSourceLoc fullSL(slStart, *SM); + insertReplacement(Rep, fullSL); + hipCounter counter = {"HIP_DYNAMIC_SHARED", ConvTypes::CONV_MEM, ApiTypes::API_RUNTIME}; + Statistics::current().incrementCounter(counter, refName.str()); + } + + return true; +} + +void HipifyAction::insertReplacement(const ct::Replacement& rep, const clang::FullSourceLoc& fullSL) { + llcompat::insertReplacement(*replacements, rep); + if (PrintStats) { + rep.getLength(); + Statistics::current().lineTouched(fullSL.getExpansionLineNumber()); + Statistics::current().bytesChanged(rep.getLength()); + } +} + +std::unique_ptr HipifyAction::CreateASTConsumer(clang::CompilerInstance& CI, llvm::StringRef) { + Finder.reset(new clang::ast_matchers::MatchFinder); + + // Replace the <<<...>>> language extension with a hip kernel launch + Finder->addMatcher(mat::cudaKernelCallExpr(mat::isExpansionInMainFile()).bind("cudaLaunchKernel"), this); + + // Replace cuda builtins. + Finder->addMatcher( + mat::memberExpr( + mat::isExpansionInMainFile(), + mat::hasObjectExpression( + mat::hasType( + mat::cxxRecordDecl( + mat::matchesName("__cuda_builtin_") + ) + ) + ) + ).bind("cudaBuiltin"), + this + ); + + Finder->addMatcher( + mat::varDecl( + mat::isExpansionInMainFile(), + mat::allOf( + mat::hasAttr(clang::attr::CUDAShared), + mat::hasType(mat::incompleteArrayType()) + ) + ).bind("cudaSharedIncompleteArrayVar"), + this + ); + + // Ownership is transferred to the caller... + return Finder->newASTConsumer(); +} + +void HipifyAction::EndSourceFileAction() { + // Insert the hip header, if we didn't already do it by accident during substitution. + if (!insertedRuntimeHeader) { + // It's not sufficient to just replace CUDA headers with hip ones, because numerous CUDA headers are + // implicitly included by the compiler. Instead, we _delete_ CUDA headers, and unconditionally insert + // one copy of the hip include into every file. + clang::SourceManager& SM = getCompilerInstance().getSourceManager(); + + clang::SourceLocation sl = SM.getLocForStartOfFile(SM.getMainFileID()); + clang::FullSourceLoc fullSL(sl, SM); + ct::Replacement Rep(SM, sl, 0, "#include \n"); + insertReplacement(Rep, fullSL); + } + + clang::ASTFrontendAction::EndSourceFileAction(); +} + + +namespace { + +/** + * A silly little class to proxy PPCallbacks back to the HipifyAction class. + */ +class PPCallbackProxy : public clang::PPCallbacks { + HipifyAction& hipifyAction; + +public: + explicit PPCallbackProxy(HipifyAction& action): hipifyAction(action) {} + + void InclusionDirective(clang::SourceLocation hash_loc, const clang::Token& include_token, + StringRef file_name, bool is_angled, clang::CharSourceRange filename_range, + const clang::FileEntry* file, StringRef search_path, StringRef relative_path, + const clang::Module* imported) override { + hipifyAction.InclusionDirective(hash_loc, include_token, file_name, is_angled, filename_range, file, search_path, relative_path, imported); + } +}; + +} + +void HipifyAction::ExecuteAction() { + clang::Preprocessor& PP = getCompilerInstance().getPreprocessor(); + clang::SourceManager& SM = getCompilerInstance().getSourceManager(); + + // Start lexing the specified input file. + const llvm::MemoryBuffer* FromFile = SM.getBuffer(SM.getMainFileID()); + clang::Lexer RawLex(SM.getMainFileID(), FromFile, SM, PP.getLangOpts()); + RawLex.SetKeepWhitespaceMode(true); + + // Perform a token-level rewrite of CUDA identifiers to hip ones. The raw-mode lexer gives us enough + // information to tell the difference between identifiers, string literals, and "other stuff". It also + // ignores preprocessor directives, so this transformation will operate inside preprocessor-deleted + // code. + clang::Token RawTok; + RawLex.LexFromRawLexer(RawTok); + while (RawTok.isNot(clang::tok::eof)) { + RewriteToken(RawTok); + RawLex.LexFromRawLexer(RawTok); + } + + // Register yourself as the preprocessor callback, by proxy. + PP.addPPCallbacks(std::unique_ptr(new PPCallbackProxy(*this))); + + // Now we're done futzing with the lexer, have the subclass proceeed with Sema and AST matching. + clang::ASTFrontendAction::ExecuteAction(); +} + +void HipifyAction::run(const clang::ast_matchers::MatchFinder::MatchResult& Result) { + if (cudaBuiltin(Result)) return; + if (cudaLaunchKernel(Result)) return; + if (cudaSharedIncompleteArrayVar(Result)) return; +} diff --git a/projects/clr/hipamd/hipify-clang/src/HipifyAction.h b/projects/clr/hipamd/hipify-clang/src/HipifyAction.h new file mode 100644 index 0000000000..03d34601f3 --- /dev/null +++ b/projects/clr/hipamd/hipify-clang/src/HipifyAction.h @@ -0,0 +1,83 @@ +#pragma once + +#include "clang/Lex/PPCallbacks.h" +#include "clang/Tooling/Tooling.h" +#include "clang/Frontend/FrontendAction.h" +#include "clang/Tooling/Core/Replacement.h" +#include "clang/ASTMatchers/ASTMatchFinder.h" +#include "ReplacementsFrontendActionFactory.h" + +namespace ct = clang::tooling; + +/** + * A FrontendAction that hipifies CUDA programs. + */ +class HipifyAction : public clang::ASTFrontendAction, + public clang::ast_matchers::MatchFinder::MatchCallback { +private: + ct::Replacements* replacements; + std::unique_ptr Finder; + + /// CUDA implicitly adds its runtime header. We rewrite explicitly-provided CUDA includes with equivalent + // ones, and track - using this flag - if the result led to us including the hip runtime header. If it did + // not, we insert it at the top of the file when we finish processing it. + // This approach means we do the best it's possible to do w.r.t preserving the user's include order. + bool insertedRuntimeHeader = false; + + /** + * Rewrite a string literal to refer to hip, not CUDA. + */ + void RewriteString(StringRef s, clang::SourceLocation start); + + /** + * Replace a CUDA identifier with the corresponding hip identifier, if applicable. + */ + void RewriteToken(const clang::Token &t); + +public: + explicit HipifyAction(ct::Replacements *replacements): + clang::ASTFrontendAction(), + replacements(replacements) {} + + // MatchCallback listeners + bool cudaBuiltin(const clang::ast_matchers::MatchFinder::MatchResult& Result); + bool cudaLaunchKernel(const clang::ast_matchers::MatchFinder::MatchResult& Result); + bool cudaSharedIncompleteArrayVar(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, + StringRef file_name, + bool is_angled, + clang::CharSourceRange filename_range, + const clang::FileEntry *file, + StringRef search_path, + StringRef relative_path, + const clang::Module *imported); + +protected: + /** + * Add a Replacement for the current file. These will all be applied after executing the FrontendAction. + */ + void insertReplacement(const ct::Replacement& rep, const clang::FullSourceLoc& fullSL); + + /** + * FrontendAction entry point. + */ + void ExecuteAction() override; + + /** + * Called at the start of each new file to process. + */ + void EndSourceFileAction() override; + + /** + * MatchCallback API entry point. Called by the AST visitor while searching the AST for things we registered an + * interest for. + */ + void run(const clang::ast_matchers::MatchFinder::MatchResult& Result) override; + + std::unique_ptr CreateASTConsumer(clang::CompilerInstance &CI, llvm::StringRef InFile) override; +}; diff --git a/projects/clr/hipamd/hipify-clang/src/ReplacementsFrontendActionFactory.h b/projects/clr/hipamd/hipify-clang/src/ReplacementsFrontendActionFactory.h new file mode 100644 index 0000000000..7896635ef6 --- /dev/null +++ b/projects/clr/hipamd/hipify-clang/src/ReplacementsFrontendActionFactory.h @@ -0,0 +1,28 @@ +#pragma once + +#include "clang/Tooling/Tooling.h" +#include "clang/Frontend/FrontendAction.h" +#include "clang/Tooling/Core/Replacement.h" + +namespace ct = clang::tooling; + + +/** + * A FrontendActionFactory that propagates a set of Replacements into the FrontendAction. + * This is necessary boilerplate for using a custom FrontendAction with a RefactoringTool. + * + * @tparam T The FrontendAction to create. + */ +template +class ReplacementsFrontendActionFactory : public ct::FrontendActionFactory { + ct::Replacements* replacements; + +public: + explicit ReplacementsFrontendActionFactory(ct::Replacements* r): + ct::FrontendActionFactory(), + replacements(r) {} + + clang::FrontendAction* create() override { + return new T(replacements); + } +}; diff --git a/projects/clr/hipamd/hipify-clang/src/main.cpp b/projects/clr/hipamd/hipify-clang/src/main.cpp new file mode 100644 index 0000000000..0cc3594466 --- /dev/null +++ b/projects/clr/hipamd/hipify-clang/src/main.cpp @@ -0,0 +1,155 @@ +/* +Copyright (c) 2015-2017 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. +*/ +/** + * @file Cuda2Hip.cpp + * + * This file is compiled and linked into clang based hipify tool. + */ +#include +#include +#include +#include +#include +#include +#include + +#include "CUDA2HipMap.h" +#include "LLVMCompat.h" +#include "HipifyAction.h" +#include "ArgParse.h" + +#define DEBUG_TYPE "cuda2hip" + +namespace ct = clang::tooling; + + +namespace { + +void copyFile(const std::string& src, const std::string& dst) { + std::ifstream source(src, std::ios::binary); + std::ofstream dest(dst, std::ios::binary); + dest << source.rdbuf(); +} + +} // anonymous namespace + +int main(int argc, const char **argv) { + llcompat::PrintStackTraceOnErrorSignal(); + + ct::CommonOptionsParser OptionsParser(argc, argv, ToolTemplateCategory, llvm::cl::OneOrMore); + std::vector fileSources = OptionsParser.getSourcePathList(); + std::string dst = OutputFilename; + if (!dst.empty() && fileSources.size() > 1) { + llvm::errs() << "[HIPIFY] conflict: -o and multiple source files are specified.\n"; + return 1; + } + + if (NoOutput) { + if (Inplace) { + llvm::errs() << "[HIPIFY] conflict: both -no-output and -inplace options are specified.\n"; + return 1; + } + if (!dst.empty()) { + llvm::errs() << "[HIPIFY] conflict: both -no-output and -o options are specified.\n"; + return 1; + } + } + + if (Examine) { + NoOutput = PrintStats = true; + } + + int Result = 0; + + // Arguments for the Statistics print routines. + std::unique_ptr csv = nullptr; + llvm::raw_ostream* statPrint = nullptr; + if (!OutputStatsFilename.empty()) { + csv = std::unique_ptr(new std::ofstream(OutputStatsFilename, std::ios_base::trunc)); + } + if (PrintStats) { + statPrint = &llvm::errs(); + } + + for (const auto & src : fileSources) { + if (dst.empty()) { + if (Inplace) { + dst = src; + } else { + dst = src + ".hip"; + } + } else if (Inplace) { + llvm::errs() << "[HIPIFY] conflict: both -o and -inplace options are specified.\n"; + return 1; + } + + std::string tmpFile = src + ".hipify-tmp"; + + // Create a copy of the file to work on. When we're done, we'll move this onto the + // output (which may mean overwriting the input, if we're in-place). + // Should we fail for some reason, we'll just leak this file and not corrupt the input. + copyFile(src, tmpFile); + + // Initialise the statistics counters for this file. + Statistics::setActive(src); + + // RefactoringTool operates on the file in-place. Giving it the output path is no good, + // because that'll break relative includes, and we don't want to overwrite the input file. + // So what we do is operate on a copy, which we then move to the output. + ct::RefactoringTool Tool(OptionsParser.getCompilations(), tmpFile); + ct::Replacements& replacementsToUse = llcompat::getReplacements(Tool, tmpFile); + + ReplacementsFrontendActionFactory actionFactory(&replacementsToUse); + + Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("--cuda-host-only", ct::ArgumentInsertPosition::BEGIN)); + + // Ensure at least c++11 is used. + Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("-std=c++11", ct::ArgumentInsertPosition::BEGIN)); +#if defined(HIPIFY_CLANG_RES) + Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("-resource-dir=" HIPIFY_CLANG_RES)); +#endif + Tool.appendArgumentsAdjuster(ct::getClangSyntaxOnlyAdjuster()); + + // Hipify _all_ the things! + if (Tool.runAndSave(&actionFactory)) { + DEBUG(llvm::dbgs() << "Skipped some replacements.\n"); + } + + // Either move the tmpfile to the output, or remove it. + if (!NoOutput) { + rename(tmpFile.c_str(), dst.c_str()); + } else { + remove(tmpFile.c_str()); + } + + Statistics::current().markCompletion(); + Statistics::current().print(csv.get(), statPrint); + + dst.clear(); + } + + if (fileSources.size() > 1) { + Statistics::printAggregate(csv.get(), statPrint); + } + + return Result; +}