From c0f62157f917ef9aab4bd23798fe4135cf8a2929 Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Mon, 14 Oct 2019 18:08:29 +0300 Subject: [PATCH] [HIPIFY] HipifyAction code clean-up --- hipify-clang/src/HipifyAction.cpp | 301 +++++++++++++----------------- hipify-clang/src/HipifyAction.h | 28 +-- 2 files changed, 141 insertions(+), 188 deletions(-) diff --git a/hipify-clang/src/HipifyAction.cpp b/hipify-clang/src/HipifyAction.cpp index c6d10f0cdc..930f3ec8c5 100644 --- a/hipify-clang/src/HipifyAction.cpp +++ b/hipify-clang/src/HipifyAction.cpp @@ -33,16 +33,20 @@ THE SOFTWARE. #include "StringUtils.h" #include "ArgParse.h" -namespace ct = clang::tooling; -namespace mat = clang::ast_matchers; - +const std::string sHIP = "HIP"; +const std::string sROC = "ROC"; +const std::string sCub = "cub"; const std::string sHIP_DYNAMIC_SHARED = "HIP_DYNAMIC_SHARED"; +const std::string sHIP_KERNEL_NAME = "HIP_KERNEL_NAME"; std::string sHIP_SYMBOL = "HIP_SYMBOL"; -std::string sHIP_KERNEL_NAME = "HIP_KERNEL_NAME"; std::string s_reinterpret_cast = "reinterpret_cast"; -const std::string sHipLaunchKernelGGL = "hipLaunchKernelGGL("; +const std::string sHipLaunchKernelGGL = "hipLaunchKernelGGL"; const std::string sDim3 = "dim3("; - +const std::string s_hiprand_kernel_h = "hiprand_kernel.h"; +const std::string s_hiprand_h = "hiprand.h"; +const std::string sOnce = "once"; +const std::string s_string_literal = "[string literal]"; +// CUDA identifiers, used in matchers const std::string sCudaMemcpyToSymbol = "cudaMemcpyToSymbol"; const std::string sCudaMemcpyToSymbolAsync = "cudaMemcpyToSymbolAsync"; const std::string sCudaGetSymbolSize = "cudaGetSymbolSize"; @@ -51,6 +55,12 @@ const std::string sCudaMemcpyFromSymbol = "cudaMemcpyFromSymbol"; const std::string sCudaMemcpyFromSymbolAsync = "cudaMemcpyFromSymbolAsync"; const std::string sCudaFuncSetCacheConfig = "cudaFuncSetCacheConfig"; const std::string sCudaFuncGetAttributes = "cudaFuncGetAttributes"; +// Matchers' names +const StringRef sCudaSharedIncompleteArrayVar = "cudaSharedIncompleteArrayVar"; +const StringRef sCudaLaunchKernel = "cudaLaunchKernel"; +const StringRef sCudaHostFuncCall = "cudaHostFuncCall"; +const StringRef sCudaDeviceFuncCall = "cudaDeviceFuncCall"; +const StringRef sCubNamespacePrefix = "cubNamespacePrefix"; std::set DeviceSymbolFunctions0 { {sCudaMemcpyToSymbol}, @@ -78,7 +88,7 @@ std::set ReinterpretFunctions1{ }; void HipifyAction::RewriteString(StringRef s, clang::SourceLocation start) { - clang::SourceManager& SM = getCompilerInstance().getSourceManager(); + auto &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); @@ -86,7 +96,7 @@ void HipifyAction::RewriteString(StringRef s, clang::SourceLocation start) { const auto found = CUDA_RENAMES_MAP().find(name); if (found != CUDA_RENAMES_MAP().end()) { StringRef repName = Statistics::isToRoc(found->second) ? found->second.rocName : found->second.hipName; - hipCounter counter = {"[string literal]", "", ConvTypes::CONV_LITERAL, ApiTypes::API_RUNTIME, found->second.supportDegree}; + hipCounter counter = {s_string_literal, "", ConvTypes::CONV_LITERAL, ApiTypes::API_RUNTIME, found->second.supportDegree}; Statistics::current().incrementCounter(counter, name.str()); if (!Statistics::isUnsupported(counter)) { clang::SourceLocation sl = start.getLocWithOffset(begin + 1); @@ -95,9 +105,7 @@ void HipifyAction::RewriteString(StringRef s, clang::SourceLocation start) { insertReplacement(Rep, fullSL); } } - if (end == StringRef::npos) { - break; - } + if (end == StringRef::npos) break; begin = end + 1; } } @@ -109,7 +117,7 @@ void HipifyAction::RewriteString(StringRef s, clang::SourceLocation start) { * 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) { +void HipifyAction::RewriteToken(const clang::Token &t) { // String literals containing CUDA references need fixing. if (t.is(clang::tok::string_literal)) { StringRef s(t.getLiteralData(), t.getLength()); @@ -124,20 +132,21 @@ void HipifyAction::RewriteToken(const clang::Token& t) { FindAndReplace(name, sl, CUDA_RENAMES_MAP()); } -void HipifyAction::FindAndReplace(llvm::StringRef name, +void HipifyAction::FindAndReplace(StringRef name, clang::SourceLocation sl, - const std::map& repMap, bool bReplace) { + const std::map &repMap, + bool bReplace) { const auto found = repMap.find(name); if (found == repMap.end()) { // So it's an identifier, but not CUDA? Boring. return; } Statistics::current().incrementCounter(found->second, name.str()); - clang::DiagnosticsEngine& DE = getCompilerInstance().getDiagnostics(); + clang::DiagnosticsEngine &DE = getCompilerInstance().getDiagnostics(); // Warn the user about unsupported identifier. if (Statistics::isUnsupported(found->second)) { std::string sWarn; - Statistics::isToRoc(found->second) ? sWarn = "ROC" : sWarn = "HIP"; + Statistics::isToRoc(found->second) ? sWarn = sROC : sWarn = sHIP; sWarn = "" + sWarn; const auto ID = DE.getCustomDiagID(clang::DiagnosticsEngine::Warning, "CUDA identifier is unsupported in %0."); DE.Report(sl, ID) << sWarn; @@ -147,7 +156,7 @@ void HipifyAction::FindAndReplace(llvm::StringRef name, return; } StringRef repName = Statistics::isToRoc(found->second) ? found->second.rocName : found->second.hipName; - clang::SourceManager& SM = getCompilerInstance().getSourceManager(); + auto &SM = getCompilerInstance().getSourceManager(); ct::Replacement Rep(SM, sl, name.size(), repName.str()); clang::FullSourceLoc fullSL(sl, SM); insertReplacement(Rep, fullSL); @@ -155,7 +164,7 @@ void HipifyAction::FindAndReplace(llvm::StringRef name, namespace { -clang::SourceRange getReadRange(clang::SourceManager& SM, const clang::SourceRange& exprRange) { +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{}); @@ -167,7 +176,7 @@ clang::SourceRange getReadRange(clang::SourceManager& SM, const clang::SourceRan } } -clang::SourceRange getWriteRange(clang::SourceManager& SM, const clang::SourceRange& exprRange) { +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. @@ -178,7 +187,7 @@ clang::SourceRange getWriteRange(clang::SourceManager& SM, const clang::SourceRa return {SM.getSpellingLoc(begin), SM.getSpellingLoc(end)}; } -StringRef readSourceText(clang::SourceManager& SM, const clang::SourceRange& exprRange) { +StringRef readSourceText(clang::SourceManager &SM, const clang::SourceRange &exprRange) { return clang::Lexer::getSourceText(clang::CharSourceRange::getTokenRange(getReadRange(SM, exprRange)), SM, clang::LangOptions(), nullptr); } @@ -186,53 +195,50 @@ StringRef readSourceText(clang::SourceManager& SM, const clang::SourceRange& exp * 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()); - } +std::string stringifyZeroDefaultedArg(clang::SourceManager &SM, const clang::Expr *arg) { + if (clang::isa(arg)) return "0"; + else return readSourceText(SM, arg->getSourceRange()); } } // anonymous namespace -bool HipifyAction::Exclude(const hipCounter & hipToken) { +bool HipifyAction::Exclude(const hipCounter &hipToken) { switch (hipToken.type) { case CONV_INCLUDE_CUDA_MAIN_H: switch (hipToken.apiType) { case API_DRIVER: case API_RUNTIME: - if (insertedRuntimeHeader) { return true; } + if (insertedRuntimeHeader) return true; insertedRuntimeHeader = true; return false; case API_BLAS: - if (insertedBLASHeader) { return true; } + if (insertedBLASHeader) return true; insertedBLASHeader = true; return false; case API_RAND: - if (hipToken.hipName == "hiprand_kernel.h") { - if (insertedRAND_kernelHeader) { return true; } + if (hipToken.hipName == s_hiprand_kernel_h) { + if (insertedRAND_kernelHeader) return true; insertedRAND_kernelHeader = true; return false; - } else if (hipToken.hipName == "hiprand.h") { - if (insertedRANDHeader) { return true; } + } else if (hipToken.hipName == s_hiprand_h) { + if (insertedRANDHeader) return true; insertedRANDHeader = true; return false; } case API_DNN: - if (insertedDNNHeader) { return true; } + if (insertedDNNHeader) return true; insertedDNNHeader = true; return false; case API_FFT: - if (insertedFFTHeader) { return true; } + if (insertedFFTHeader) return true; insertedFFTHeader = true; return false; case API_COMPLEX: - if (insertedComplexHeader) { return true; } + if (insertedComplexHeader) return true; insertedComplexHeader = true; return false; case API_SPARSE: - if (insertedSPARSEHeader) { return true; } + if (insertedSPARSEHeader) return true; insertedSPARSEHeader = true; return false; default: @@ -240,13 +246,11 @@ bool HipifyAction::Exclude(const hipCounter & hipToken) { } return false; case CONV_INCLUDE: - if (hipToken.hipName.empty()) { - return true; - } + if (hipToken.hipName.empty()) return true; switch (hipToken.apiType) { case API_RAND: - if (hipToken.hipName == "hiprand_kernel.h") { - if (insertedRAND_kernelHeader) { return true; } + if (hipToken.hipName == s_hiprand_kernel_h) { + if (insertedRAND_kernelHeader) return true; insertedRAND_kernelHeader = true; } return false; @@ -267,24 +271,19 @@ void HipifyAction::InclusionDirective(clang::SourceLocation hash_loc, clang::CharSourceRange filename_range, const clang::FileEntry*, StringRef, StringRef, const clang::Module*) { - clang::SourceManager& SM = getCompilerInstance().getSourceManager(); - if (!SM.isWrittenInMainFile(hash_loc)) { - return; - } + auto &SM = getCompilerInstance().getSourceManager(); + if (!SM.isWrittenInMainFile(hash_loc)) return; if (!firstHeader) { firstHeader = true; firstHeaderLoc = hash_loc; } const auto found = CUDA_INCLUDE_MAP.find(file_name); - if (found == CUDA_INCLUDE_MAP.end()) { - return; - } + if (found == CUDA_INCLUDE_MAP.end()) return; bool exclude = Exclude(found->second); Statistics::current().incrementCounter(found->second, file_name.str()); - clang::SourceLocation sl = filename_range.getBegin(); if (Statistics::isUnsupported(found->second)) { - clang::DiagnosticsEngine& DE = getCompilerInstance().getDiagnostics(); + clang::DiagnosticsEngine &DE = getCompilerInstance().getDiagnostics(); DE.Report(sl, DE.getCustomDiagID(clang::DiagnosticsEngine::Warning, "Unsupported CUDA header.")); return; } @@ -293,11 +292,8 @@ void HipifyAction::InclusionDirective(clang::SourceLocation hash_loc, if (!exclude) { clang::SmallString<128> includeBuffer; llvm::StringRef name = Statistics::isToRoc(found->second) ? found->second.rocName : found->second.hipName; - if (is_angled) { - newInclude = llvm::Twine("<" + name+ ">").toStringRef(includeBuffer); - } else { - newInclude = llvm::Twine("\"" + name + "\"").toStringRef(includeBuffer); - } + if (is_angled) newInclude = llvm::Twine("<" + name+ ">").toStringRef(includeBuffer); + else newInclude = llvm::Twine("\"" + name + "\"").toStringRef(includeBuffer); } else { // hashLoc is location of the '#', thus replacing the whole include directive by empty newInclude starting with '#'. sl = hash_loc; @@ -309,46 +305,33 @@ void HipifyAction::InclusionDirective(clang::SourceLocation hash_loc, } void HipifyAction::PragmaDirective(clang::SourceLocation Loc, clang::PragmaIntroducerKind Introducer) { - if (pragmaOnce) { - return; - } - clang::SourceManager& SM = getCompilerInstance().getSourceManager(); - if (!SM.isWrittenInMainFile(Loc)) { - return; - } - clang::Preprocessor& PP = getCompilerInstance().getPreprocessor(); + if (pragmaOnce) return; + auto &SM = getCompilerInstance().getSourceManager(); + if (!SM.isWrittenInMainFile(Loc)) return; + clang::Preprocessor &PP = getCompilerInstance().getPreprocessor(); clang::Token tok; PP.Lex(tok); StringRef Text(SM.getCharacterData(tok.getLocation()), tok.getLength()); - if (Text == "once") { + if (Text == sOnce) { pragmaOnce = true; pragmaOnceLoc = tok.getEndLoc(); } } -bool HipifyAction::cudaLaunchKernel(const clang::ast_matchers::MatchFinder::MatchResult& Result) { - StringRef refName = "cudaLaunchKernel"; - const auto* launchKernel = Result.Nodes.getNodeAs(refName); - if (!launchKernel) { - return false; - } - const clang::Expr* calleeExpr = launchKernel->getCallee(); - if (!calleeExpr) { - return false; - } - const clang::FunctionDecl *caleeDecl = launchKernel->getDirectCallee(); - if (!caleeDecl) { - return false; - } - const clang::CallExpr* config = launchKernel->getConfig(); - if (!config) { - return false; - } +bool HipifyAction::cudaLaunchKernel(const mat::MatchFinder::MatchResult &Result) { + auto *launchKernel = Result.Nodes.getNodeAs(sCudaLaunchKernel); + if (!launchKernel) return false; + auto *calleeExpr = launchKernel->getCallee(); + if (!calleeExpr) return false; + auto *caleeDecl = launchKernel->getDirectCallee(); + if (!caleeDecl) return false; + auto *config = launchKernel->getConfig(); + if (!config) return false; clang::SmallString<40> XStr; llvm::raw_svector_ostream OS(XStr); clang::LangOptions DefaultLangOptions; - clang::SourceManager* SM = Result.SourceManager; - OS << sHipLaunchKernelGGL; + auto *SM = Result.SourceManager; + OS << sHipLaunchKernelGGL << "("; if (caleeDecl->isTemplateInstantiation()) OS << sHIP_KERNEL_NAME << "("; OS << readSourceText(*SM, calleeExpr->getSourceRange()); if (caleeDecl->isTemplateInstantiation()) OS << ")"; @@ -381,29 +364,24 @@ bool HipifyAction::cudaLaunchKernel(const clang::ast_matchers::MatchFinder::Matc ct::Replacement Rep(*SM, launchStart, length, OS.str()); clang::FullSourceLoc fullSL(launchStart, *SM); insertReplacement(Rep, fullSL); - hipCounter counter = {"hipLaunchKernelGGL", "", ConvTypes::CONV_KERNEL_LAUNCH, ApiTypes::API_RUNTIME}; - Statistics::current().incrementCounter(counter, refName.str()); + hipCounter counter = {sHipLaunchKernelGGL, "", ConvTypes::CONV_KERNEL_LAUNCH, ApiTypes::API_RUNTIME}; + Statistics::current().incrementCounter(counter, sCudaLaunchKernel.str()); 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; - } +bool HipifyAction::cudaSharedIncompleteArrayVar(const mat::MatchFinder::MatchResult &Result) { + auto *sharedVar = Result.Nodes.getNodeAs(sCudaSharedIncompleteArrayVar); + if (!sharedVar) return false; // Example: extern __shared__ uint sRadix1[]; - if (!sharedVar->hasExternalFormalLinkage()) { - return false; - } + if (!sharedVar->hasExternalFormalLinkage()) return false; clang::QualType QT = sharedVar->getType(); std::string typeName; if (QT->isIncompleteArrayType()) { - const clang::ArrayType* AT = QT.getTypePtr()->getAsArrayTypeUnsafe(); + const clang::ArrayType *AT = QT.getTypePtr()->getAsArrayTypeUnsafe(); QT = AT->getElementType(); if (QT.getTypePtr()->isBuiltinType()) { QT = QT.getCanonicalType(); - const auto* BT = clang::dyn_cast(QT); + auto *BT = clang::dyn_cast(QT); if (BT) { clang::LangOptions LO; LO.CUDA = true; @@ -417,7 +395,7 @@ bool HipifyAction::cudaSharedIncompleteArrayVar(const clang::ast_matchers::Match if (!typeName.empty()) { clang::SourceLocation slStart = sharedVar->getOuterLocStart(); clang::SourceLocation slEnd = llcompat::getEndLoc(sharedVar->getTypeSourceInfo()->getTypeLoc()); - clang::SourceManager* SM = Result.SourceManager; + auto *SM = Result.SourceManager; size_t repLength = SM->getCharacterData(slEnd) - SM->getCharacterData(slStart) + 1; std::string varName = sharedVar->getNameAsString(); std::string repName = sHIP_DYNAMIC_SHARED + "(" + typeName + ", " + varName + ")"; @@ -425,52 +403,39 @@ bool HipifyAction::cudaSharedIncompleteArrayVar(const clang::ast_matchers::Match clang::FullSourceLoc fullSL(slStart, *SM); insertReplacement(Rep, fullSL); hipCounter counter = {sHIP_DYNAMIC_SHARED, "", ConvTypes::CONV_EXTERN_SHARED, ApiTypes::API_RUNTIME}; - Statistics::current().incrementCounter(counter, refName.str()); + Statistics::current().incrementCounter(counter, sCudaSharedIncompleteArrayVar.str()); return true; } return false; } -bool HipifyAction::cudaDeviceFuncCall(const clang::ast_matchers::MatchFinder::MatchResult& Result) { - if (const clang::CallExpr *call = Result.Nodes.getNodeAs("cudaDeviceFuncCall")) { - const clang::FunctionDecl *funcDcl = call->getDirectCallee(); - if (!funcDcl) { - return false; - } +bool HipifyAction::cudaDeviceFuncCall(const mat::MatchFinder::MatchResult &Result) { + if (const clang::CallExpr *call = Result.Nodes.getNodeAs(sCudaDeviceFuncCall)) { + auto *funcDcl = call->getDirectCallee(); + if (!funcDcl) return false; FindAndReplace(funcDcl->getDeclName().getAsString(), llcompat::getBeginLoc(call), CUDA_DEVICE_FUNC_MAP, false); return true; } return false; } -bool HipifyAction::cubNamespacePrefix(const clang::ast_matchers::MatchFinder::MatchResult& Result) { - if (const clang::TypedefNameDecl *decl = Result.Nodes.getNodeAs("cubNamespacePrefix")) { - if (!decl) { - return false; - } +bool HipifyAction::cubNamespacePrefix(const mat::MatchFinder::MatchResult &Result) { + if (auto *decl = Result.Nodes.getNodeAs(sCubNamespacePrefix)) { clang::QualType QT = decl->getUnderlyingType(); - const clang::Type* t = QT.getTypePtr(); - if (!t) { - return false; - } - const clang::ElaboratedType* et = t->getAs(); - if (!et) { - return false; - } + auto *t = QT.getTypePtr(); + if (!t) return false; + const clang::ElaboratedType *et = t->getAs(); + if (!et) return false; const clang::NestedNameSpecifier *nns = et->getQualifier(); - if (!nns) { - return false; - } + if (!nns) return false; const clang::NamespaceDecl *nsd = nns->getAsNamespace(); - if (!nsd) { - return false; - } + if (!nsd) return false; const clang::TypeSourceInfo *si = decl->getTypeSourceInfo(); const clang::TypeLoc tloc = si->getTypeLoc(); const clang::SourceRange sr = tloc.getSourceRange(); clang::SourceLocation sl(sr.getBegin()); clang::SourceLocation end(sr.getEnd()); - clang::SourceManager& SM = getCompilerInstance().getSourceManager(); + auto &SM = getCompilerInstance().getSourceManager(); size_t length = SM.getCharacterData(end) - SM.getCharacterData(sl); StringRef sfull = StringRef(SM.getCharacterData(sl), length); std::string name = nsd->getDeclName().getAsString(); @@ -484,15 +449,11 @@ bool HipifyAction::cubNamespacePrefix(const clang::ast_matchers::MatchFinder::Ma return false; } -bool HipifyAction::cudaHostFuncCall(const clang::ast_matchers::MatchFinder::MatchResult& Result) { - if (const clang::CallExpr * call = Result.Nodes.getNodeAs("cudaHostFuncCall")) { - if (!call->getNumArgs()) { - return false; - } - const clang::FunctionDecl* funcDcl = call->getDirectCallee(); - if (!funcDcl) { - return false; - } +bool HipifyAction::cudaHostFuncCall(const mat::MatchFinder::MatchResult &Result) { + if (auto *call = Result.Nodes.getNodeAs(sCudaHostFuncCall)) { + if (!call->getNumArgs()) return false; + auto *funcDcl = call->getDirectCallee(); + if (!funcDcl) return false; std::string sName = funcDcl->getDeclName().getAsString(); unsigned int argNum = 0; bool b_reinterpret = (ReinterpretFunctions.find(sName) != ReinterpretFunctions.end()) ? true : false; @@ -506,7 +467,7 @@ bool HipifyAction::cudaHostFuncCall(const clang::ast_matchers::MatchFinder::Matc clang::SmallString<40> XStr; llvm::raw_svector_ostream OS(XStr); clang::SourceRange sr = call->getArg(argNum)->getSourceRange(); - clang::SourceManager* SM = Result.SourceManager; + auto *SM = Result.SourceManager; OS << (b_reinterpret ? s_reinterpret_cast : sHIP_SYMBOL) << "(" << readSourceText(*SM, sr) << ")"; clang::SourceRange replacementRange = getWriteRange(*SM, { sr.getBegin(), sr.getEnd() }); clang::SourceLocation s = replacementRange.getBegin(); @@ -521,7 +482,7 @@ bool HipifyAction::cudaHostFuncCall(const clang::ast_matchers::MatchFinder::Matc return false; } -void HipifyAction::insertReplacement(const ct::Replacement& rep, const clang::FullSourceLoc& fullSL) { +void HipifyAction::insertReplacement(const ct::Replacement &rep, const clang::FullSourceLoc &fullSL) { llcompat::insertReplacement(*replacements, rep); if (PrintStats) { rep.getLength(); @@ -530,10 +491,10 @@ void HipifyAction::insertReplacement(const ct::Replacement& rep, const clang::Fu } } -std::unique_ptr HipifyAction::CreateASTConsumer(clang::CompilerInstance& CI, llvm::StringRef) { - Finder.reset(new clang::ast_matchers::MatchFinder); +std::unique_ptr HipifyAction::CreateASTConsumer(clang::CompilerInstance &CI, StringRef) { + Finder.reset(new mat::MatchFinder); // Replace the <<<...>>> language extension with a hip kernel launch - Finder->addMatcher(mat::cudaKernelCallExpr(mat::isExpansionInMainFile()).bind("cudaLaunchKernel"), this); + Finder->addMatcher(mat::cudaKernelCallExpr(mat::isExpansionInMainFile()).bind(sCudaLaunchKernel), this); Finder->addMatcher( mat::varDecl( mat::isExpansionInMainFile(), @@ -541,7 +502,7 @@ std::unique_ptr HipifyAction::CreateASTConsumer(clang::Compi mat::hasAttr(clang::attr::CUDAShared), mat::hasType(mat::incompleteArrayType()) ) - ).bind("cudaSharedIncompleteArrayVar"), + ).bind(sCudaSharedIncompleteArrayVar), this ); Finder->addMatcher( @@ -561,7 +522,7 @@ std::unique_ptr HipifyAction::CreateASTConsumer(clang::Compi ) ) ) - ).bind("cudaHostFuncCall"), + ).bind(sCudaHostFuncCall), this ); Finder->addMatcher( @@ -576,7 +537,7 @@ std::unique_ptr HipifyAction::CreateASTConsumer(clang::Compi mat::unless(mat::hasAttr(clang::attr::CUDAHost)) ) ) - ).bind("cudaDeviceFuncCall"), + ).bind(sCudaDeviceFuncCall), this ); Finder->addMatcher( @@ -586,12 +547,12 @@ std::unique_ptr HipifyAction::CreateASTConsumer(clang::Compi mat::elaboratedType( mat::hasQualifier( mat::specifiesNamespace( - mat::hasName("cub") + mat::hasName(sCub) ) ) ) ) - ).bind("cubNamespacePrefix"), + ).bind(sCubNamespacePrefix), this ); // Ownership is transferred to the caller. @@ -599,10 +560,8 @@ std::unique_ptr HipifyAction::CreateASTConsumer(clang::Compi } void HipifyAction::Ifndef(clang::SourceLocation Loc, const clang::Token &MacroNameTok, const clang::MacroDefinition &MD) { - clang::SourceManager& SM = getCompilerInstance().getSourceManager(); - if (!SM.isWrittenInMainFile(Loc)) { - return; - } + auto &SM = getCompilerInstance().getSourceManager(); + if (!SM.isWrittenInMainFile(Loc)) return; StringRef Text(SM.getCharacterData(MacroNameTok.getLocation()), MacroNameTok.getLength()); Ifndefs.insert(std::make_pair(Text.str(), MacroNameTok.getEndLoc())); } @@ -615,12 +574,12 @@ void HipifyAction::EndSourceFileAction() { // one copy of the hip include into every file. bool placeForIncludeCalculated = false; clang::SourceLocation sl, controllingMacroLoc; - clang::SourceManager& SM = getCompilerInstance().getSourceManager(); - clang::Preprocessor& PP = getCompilerInstance().getPreprocessor(); - clang::HeaderSearch& HS = PP.getHeaderSearchInfo(); - clang::ExternalPreprocessorSource* EPL = HS.getExternalLookup(); - const clang::FileEntry* FE = SM.getFileEntryForID(SM.getMainFileID()); - const clang::IdentifierInfo* controllingMacro = HS.getFileInfo(FE).getControllingMacro(EPL); + auto &SM = getCompilerInstance().getSourceManager(); + clang::Preprocessor &PP = getCompilerInstance().getPreprocessor(); + clang::HeaderSearch &HS = PP.getHeaderSearchInfo(); + clang::ExternalPreprocessorSource *EPL = HS.getExternalLookup(); + const clang::FileEntry *FE = SM.getFileEntryForID(SM.getMainFileID()); + const clang::IdentifierInfo *controllingMacro = HS.getFileInfo(FE).getControllingMacro(EPL); if (controllingMacro) { auto found = Ifndefs.find(controllingMacro->getName().str()); if (found != Ifndefs.end()) { @@ -629,19 +588,13 @@ void HipifyAction::EndSourceFileAction() { } } if (pragmaOnce) { - if (placeForIncludeCalculated) { - sl = pragmaOnceLoc < controllingMacroLoc ? pragmaOnceLoc : controllingMacroLoc; - } else { - sl = pragmaOnceLoc; - } + if (placeForIncludeCalculated) sl = pragmaOnceLoc < controllingMacroLoc ? pragmaOnceLoc : controllingMacroLoc; + else sl = pragmaOnceLoc; placeForIncludeCalculated = true; } if (!placeForIncludeCalculated) { - if (firstHeader) { - sl = firstHeaderLoc; - } else { - sl = SM.getLocForStartOfFile(SM.getMainFileID()); - } + if (firstHeader) sl = firstHeaderLoc; + else sl = SM.getLocForStartOfFile(SM.getMainFileID()); } clang::FullSourceLoc fullSL(sl, SM); ct::Replacement Rep(SM, sl, 0, "\n#include \n"); @@ -656,15 +609,15 @@ namespace { * A silly little class to proxy PPCallbacks back to the HipifyAction class. */ class PPCallbackProxy : public clang::PPCallbacks { - HipifyAction& hipifyAction; + HipifyAction &hipifyAction; public: - explicit PPCallbackProxy(HipifyAction& action): hipifyAction(action) {} + explicit PPCallbackProxy(HipifyAction &action): hipifyAction(action) {} - void InclusionDirective(clang::SourceLocation hash_loc, const clang::Token& include_token, + 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 + const clang::FileEntry *file, StringRef search_path, StringRef relative_path, + const clang::Module *imported #if LLVM_VERSION_MAJOR > 6 , clang::SrcMgr::CharacteristicKind FileType #endif @@ -688,10 +641,10 @@ bool HipifyAction::BeginInvocation(clang::CompilerInstance &CI) { } void HipifyAction::ExecuteAction() { - clang::Preprocessor& PP = getCompilerInstance().getPreprocessor(); - clang::SourceManager& SM = getCompilerInstance().getSourceManager(); + clang::Preprocessor &PP = getCompilerInstance().getPreprocessor(); + auto &SM = getCompilerInstance().getSourceManager(); // Start lexing the specified input file. - const llvm::MemoryBuffer* FromFile = SM.getBuffer(SM.getMainFileID()); + 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 @@ -709,7 +662,7 @@ void HipifyAction::ExecuteAction() { clang::ASTFrontendAction::ExecuteAction(); } -void HipifyAction::run(const clang::ast_matchers::MatchFinder::MatchResult& Result) { +void HipifyAction::run(const mat::MatchFinder::MatchResult &Result) { if (cudaLaunchKernel(Result)) return; if (cudaSharedIncompleteArrayVar(Result)) return; if (cudaHostFuncCall(Result)) return; diff --git a/hipify-clang/src/HipifyAction.h b/hipify-clang/src/HipifyAction.h index a24404deee..78532449ee 100644 --- a/hipify-clang/src/HipifyAction.h +++ b/hipify-clang/src/HipifyAction.h @@ -31,17 +31,18 @@ THE SOFTWARE. #include "Statistics.h" namespace ct = clang::tooling; +namespace mat = clang::ast_matchers; using namespace llvm; /** * A FrontendAction that hipifies CUDA programs. */ class HipifyAction : public clang::ASTFrontendAction, - public clang::ast_matchers::MatchFinder::MatchCallback { + public mat::MatchFinder::MatchCallback { private: - ct::Replacements* replacements; + ct::Replacements *replacements; std::map Ifndefs; - std::unique_ptr Finder; + 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. @@ -67,12 +68,11 @@ 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); - bool cudaDeviceFuncCall(const clang::ast_matchers::MatchFinder::MatchResult& Result); - bool cudaHostFuncCall(const clang::ast_matchers::MatchFinder::MatchResult& Result); - bool cubNamespacePrefix(const clang::ast_matchers::MatchFinder::MatchResult& Result); + bool cudaLaunchKernel(const mat::MatchFinder::MatchResult &Result); + bool cudaSharedIncompleteArrayVar(const mat::MatchFinder::MatchResult &Result); + bool cudaDeviceFuncCall(const mat::MatchFinder::MatchResult &Result); + bool cudaHostFuncCall(const mat::MatchFinder::MatchResult &Result); + bool cubNamespacePrefix(const mat::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, @@ -91,7 +91,7 @@ public: 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); + void insertReplacement(const ct::Replacement &rep, const clang::FullSourceLoc &fullSL); // FrontendAction entry point. void ExecuteAction() override; // Callback before starting processing a single input; used by hipify-clang for setting Preprocessor options. @@ -99,8 +99,8 @@ protected: // 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; - bool Exclude(const hipCounter & hipToken); - void FindAndReplace(llvm::StringRef name, clang::SourceLocation sl, const std::map& repMap, bool bReplace = true); + void run(const mat::MatchFinder::MatchResult &Result) override; + std::unique_ptr CreateASTConsumer(clang::CompilerInstance &CI, StringRef InFile) override; + bool Exclude(const hipCounter &hipToken); + void FindAndReplace(StringRef name, clang::SourceLocation sl, const std::map &repMap, bool bReplace = true); };