From af81909cabb7ee9d058036f7327d5c7b37cef566 Mon Sep 17 00:00:00 2001 From: Chris Kitching Date: Wed, 18 Oct 2017 10:32:06 +0100 Subject: [PATCH 1/6] Prefer early-return to deep nesting A chain of 7 closing braces is never a great sign :D In the process it became apparant that the unsupported flag was being silently ignored, causing users to be left with cuda API calls in their programs with no warning given. This has been rectified for consistency. [ROCm/clr commit: c1f4612176968dae85ef06b26f89e5ec827e99dd] --- .../clr/hipamd/hipify-clang/src/Cuda2Hip.cpp | 268 ++++++++++-------- 1 file changed, 148 insertions(+), 120 deletions(-) diff --git a/projects/clr/hipamd/hipify-clang/src/Cuda2Hip.cpp b/projects/clr/hipamd/hipify-clang/src/Cuda2Hip.cpp index 4bffcdaed9..046b3823ec 100644 --- a/projects/clr/hipamd/hipify-clang/src/Cuda2Hip.cpp +++ b/projects/clr/hipamd/hipify-clang/src/Cuda2Hip.cpp @@ -291,147 +291,175 @@ public: const FileEntry *file, StringRef search_path, StringRef relative_path, const clang::Module *imported) override { - if (_sm->isWrittenInMainFile(hash_loc)) { - if (is_angled) { - const auto found = CUDA_INCLUDE_MAP.find(file_name); - if (found != CUDA_INCLUDE_MAP.end()) { - updateCounters(found->second, file_name.str()); - if (!found->second.unsupported) { - 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); - } - } else { -// llvm::outs() << "[HIPIFY] warning: the following reference is not handled: '" << file_name << "' [inclusion directive].\n"; - } - } + 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; + } + + updateCounters(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); } virtual void MacroDefined(const Token &MacroNameTok, const MacroDirective *MD) override { - if (_sm->isWrittenInMainFile(MD->getLocation()) && - MD->getKind() == MacroDirective::MD_Define) { - for (auto T : MD->getMacroInfo()->tokens()) { - if (T.isAnyIdentifier()) { - StringRef name = T.getIdentifierInfo()->getName(); - const auto found = CUDA_RENAMES_MAP().find(name); - if (found != CUDA_RENAMES_MAP().end()) { - updateCounters(found->second, name.str()); - if (!found->second.unsupported) { - StringRef repName = found->second.hipName; - SourceLocation sl = T.getLocation(); - DEBUG(dbgs() << "Identifier " << name << " found in definition of macro " - << MacroNameTok.getIdentifierInfo()->getName() << "\n" - << "will be replaced with: " << repName << "\n" - << "SourceLocation: " << sl.printToString(*_sm) << "\n"); - Replacement Rep(*_sm, sl, name.size(), repName); - FullSourceLoc fullSL(sl, *_sm); - insertReplacement(Rep, fullSL); - } - } else { - // llvm::outs() << "[HIPIFY] warning: the following reference is not handled: '" << name << "' [macro].\n"; - } - } + if (!_sm->isWrittenInMainFile(MD->getLocation()) || + MD->getKind() != MacroDirective::MD_Define) { + return; + } + + for (auto T : MD->getMacroInfo()->tokens()) { + // We're looking for CUDA identifiers in the macro definition to rewrite... + if (!T.isAnyIdentifier()) { + continue; } + + StringRef name = T.getIdentifierInfo()->getName(); + const auto found = CUDA_RENAMES_MAP().find(name); + if (found == CUDA_RENAMES_MAP().end()) { + // So it's an identifier that isn't CUDA? Boring. + continue; + } + + updateCounters(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()); + continue; + } + + StringRef repName = found->second.hipName; + DEBUG(dbgs() << "Identifier " << name << " found in definition of macro " + << MacroNameTok.getIdentifierInfo()->getName() << "\n" + << "will be replaced with: " << repName << "\n" + << "SourceLocation: " << sl.printToString(*_sm) << "\n"); + Replacement Rep(*_sm, sl, name.size(), repName); + FullSourceLoc fullSL(sl, *_sm); + insertReplacement(Rep, fullSL); } } virtual void MacroExpands(const Token &MacroNameTok, const MacroDefinition &MD, SourceRange Range, const MacroArgs *Args) override { - if (_sm->isWrittenInMainFile(MacroNameTok.getLocation())) { - // The getNumArgs function was rather unhelpfully renamed in clang 4.0. Its semantics - // remain unchanged. + + if (!_sm->isWrittenInMainFile(MacroNameTok.getLocation())) { + return; // Macros in headers are not our concern. + } + + // The getNumArgs function was rather unhelpfully renamed in clang 4.0. Its semantics + // remain unchanged. #if LLVM_VERSION_MAJOR > 4 - #define GET_NUM_ARGS() getNumParams() + #define GET_NUM_ARGS() getNumParams() #else - #define GET_NUM_ARGS() getNumArgs() + #define GET_NUM_ARGS() getNumArgs() #endif - 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; + + 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; #if (LLVM_VERSION_MAJOR == 3) && (LLVM_VERSION_MINOR == 8) - _pp->EnterTokenStream(start, len, false, false); + _pp->EnterTokenStream(start, len, false, false); #else - _pp->EnterTokenStream(ArrayRef(start, len), false); + _pp->EnterTokenStream(ArrayRef(start, len), false); #endif - 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) { - if (tok.isAnyIdentifier()) { - StringRef name = tok.getIdentifierInfo()->getName(); + 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) { + if (tok.isAnyIdentifier()) { + StringRef name = tok.getIdentifierInfo()->getName(); + SourceLocation sl = tok.getLocation(); + + const auto found = CUDA_RENAMES_MAP().find(name); + if (found == CUDA_RENAMES_MAP().end()) { + // It's not a CUDA identifier. We have nothing to do. + continue; + } + + updateCounters(found->second, name.str()); + if (found->second.unsupported) { + // We know about it, but it isn't supported. Warn the user. + printHipifyMessage(*_sm, sl, "Unsupported CUDA identifier: " + name.str()); + continue; + } + + StringRef repName = found->second.hipName; + DEBUG(dbgs() << "Identifier " << name + << " found as an actual argument in expansion of macro " + << MacroNameTok.getIdentifierInfo()->getName() << "\n" + << "will be replaced with: " << repName << "\n"); + size_t length = name.size(); + if (_sm->isMacroBodyExpansion(sl)) { + LangOptions DefaultLangOptions; + SourceLocation sl_macro = _sm->getExpansionLoc(sl); + SourceLocation sl_end = Lexer::getLocForEndOfToken(sl_macro, 0, *_sm, DefaultLangOptions); + length = _sm->getCharacterData(sl_end) - _sm->getCharacterData(sl_macro); + sl = sl_macro; + } + Replacement Rep(*_sm, sl, length, repName); + FullSourceLoc fullSL(sl, *_sm); + insertReplacement(Rep, fullSL); + } else if (tok.isLiteral()) { + SourceLocation sl = tok.getLocation(); + if (_sm->isMacroBodyExpansion(sl)) { + LangOptions DefaultLangOptions; + SourceLocation sl_macro = _sm->getExpansionLoc(sl); + SourceLocation sl_end = Lexer::getLocForEndOfToken(sl_macro, 0, *_sm, DefaultLangOptions); + size_t length = _sm->getCharacterData(sl_end) - _sm->getCharacterData(sl_macro); + StringRef name = StringRef(_sm->getCharacterData(sl_macro), length); + const auto found = CUDA_RENAMES_MAP().find(name); - if (found != CUDA_RENAMES_MAP().end()) { - updateCounters(found->second, name.str()); - if (!found->second.unsupported) { - StringRef repName = found->second.hipName; - DEBUG(dbgs() << "Identifier " << name - << " found as an actual argument in expansion of macro " - << MacroNameTok.getIdentifierInfo()->getName() << "\n" - << "will be replaced with: " << repName << "\n"); - size_t length = name.size(); - SourceLocation sl = tok.getLocation(); - if (_sm->isMacroBodyExpansion(sl)) { - LangOptions DefaultLangOptions; - SourceLocation sl_macro = _sm->getExpansionLoc(sl); - SourceLocation sl_end = Lexer::getLocForEndOfToken(sl_macro, 0, *_sm, DefaultLangOptions); - length = _sm->getCharacterData(sl_end) - _sm->getCharacterData(sl_macro); - name = StringRef(_sm->getCharacterData(sl_macro), length); - sl = sl_macro; - } - Replacement Rep(*_sm, sl, length, repName); - FullSourceLoc fullSL(sl, *_sm); - insertReplacement(Rep, fullSL); - } - } else { - // llvm::outs() << "[HIPIFY] warning: the following reference is not handled: '" << name << "' [macro expansion].\n"; + if (found == CUDA_RENAMES_MAP().end()) { + continue; // Not CUDA, we don't care. } - } else if (tok.isLiteral()) { - SourceLocation sl = tok.getLocation(); - if (_sm->isMacroBodyExpansion(sl)) { - LangOptions DefaultLangOptions; - SourceLocation sl_macro = _sm->getExpansionLoc(sl); - SourceLocation sl_end = Lexer::getLocForEndOfToken(sl_macro, 0, *_sm, DefaultLangOptions); - size_t length = _sm->getCharacterData(sl_end) - _sm->getCharacterData(sl_macro); - StringRef name = StringRef(_sm->getCharacterData(sl_macro), length); - const auto found = CUDA_RENAMES_MAP().find(name); - if (found != CUDA_RENAMES_MAP().end()) { - updateCounters(found->second, name.str()); - if (!found->second.unsupported) { - StringRef repName = found->second.hipName; - sl = sl_macro; - Replacement Rep(*_sm, sl, length, repName); - FullSourceLoc fullSL(sl, *_sm); - insertReplacement(Rep, fullSL); - } - } else { - // llvm::outs() << "[HIPIFY] warning: the following reference is not handled: '" << name << "' [literal macro expansion].\n"; - } - } else { - if (tok.is(tok::string_literal)) { - StringRef s(tok.getLiteralData(), tok.getLength()); - processString(unquoteStr(s), *_sm, tok.getLocation()); - } + + updateCounters(found->second, name.str()); + if (found->second.unsupported) { + printHipifyMessage(*_sm, sl, "Unsupported CUDA identifier: " + name.str()); + continue; } + + StringRef repName = found->second.hipName; + sl = sl_macro; + Replacement Rep(*_sm, sl, length, repName); + FullSourceLoc fullSL(sl, *_sm); + insertReplacement(Rep, fullSL); + } else if (tok.is(tok::string_literal)) { + StringRef s(tok.getLiteralData(), tok.getLength()); + processString(unquoteStr(s), *_sm, tok.getLocation()); } } } From cf50b4f97ae35a20fd84f7f4660229ff21e276e1 Mon Sep 17 00:00:00 2001 From: Chris Kitching Date: Wed, 18 Oct 2017 11:25:17 +0100 Subject: [PATCH 2/6] Don't special-case source locations for calls in macros The source location for a call that's inside a macro body will, by default, point into the macro definition itself. The original logic was causing macro invocations to be overwritten, as I explain here: https://github.com/ROCm-Developer-Tools/HIP/issues/207#issuecomment-337521851 The existing PPCallbacks code is correctly rewriting macro definitions, so the practical effect of this change is that AST rewrites on code that's expanded from macros are no-ops. It might be a performance optimisation to put a short-circiut at the top of the AST callbacks to abort when faced with code that was expanded from macros. It might yet prove wise to do absolutely everything at lex-time... [ROCm/clr commit: f7e65c5334c41d8206e81a600bfb66967ea35750] --- .../clr/hipamd/hipify-clang/src/Cuda2Hip.cpp | 26 +++---------------- 1 file changed, 4 insertions(+), 22 deletions(-) diff --git a/projects/clr/hipamd/hipify-clang/src/Cuda2Hip.cpp b/projects/clr/hipamd/hipify-clang/src/Cuda2Hip.cpp index 046b3823ec..a3ec11ff1a 100644 --- a/projects/clr/hipamd/hipify-clang/src/Cuda2Hip.cpp +++ b/projects/clr/hipamd/hipify-clang/src/Cuda2Hip.cpp @@ -531,28 +531,10 @@ private: } size_t length = name.size(); - bool bReplace = true; - if (SM->isMacroArgExpansion(sl)) { - sl = SM->getImmediateSpellingLoc(sl); - } else if (SM->isMacroBodyExpansion(sl)) { - LangOptions DefaultLangOptions; - SourceLocation sl_macro = SM->getExpansionLoc(sl); - SourceLocation sl_end = Lexer::getLocForEndOfToken(sl_macro, 0, *SM, DefaultLangOptions); - length = SM->getCharacterData(sl_end) - SM->getCharacterData(sl_macro); - StringRef macroName = StringRef(SM->getCharacterData(sl_macro), length); - if (CUDA_EXCLUDES.end() != CUDA_EXCLUDES.find(macroName)) { - bReplace = false; - } else { - sl = sl_macro; - } - } - - if (bReplace) { - updateCounters(found->second, name); - Replacement Rep(*SM, sl, length, hipCtr.hipName); - FullSourceLoc fullSL(sl, *SM); - insertReplacement(Rep, fullSL); - } + updateCounters(found->second, name); + Replacement Rep(*SM, sl, length, hipCtr.hipName); + FullSourceLoc fullSL(sl, *SM); + insertReplacement(Rep, fullSL); return true; } From b25b17b6b30c960a378acce09ccb1fae42fb6e36 Mon Sep 17 00:00:00 2001 From: Chris Kitching Date: Wed, 18 Oct 2017 14:25:14 +0100 Subject: [PATCH 3/6] Rewrite _all_ CUDA macro identifiers in the preprocessor Calls to macros that were themselves CUDA API calls were often being missed - this applies the identifier transform to macro names at the callsites, too. [ROCm/clr commit: 30e7e7d9190b4ea68aea0b1c1d5afb59d5860162] --- .../clr/hipamd/hipify-clang/src/Cuda2Hip.cpp | 27 +++++++++++++++++++ 1 file changed, 27 insertions(+) diff --git a/projects/clr/hipamd/hipify-clang/src/Cuda2Hip.cpp b/projects/clr/hipamd/hipify-clang/src/Cuda2Hip.cpp index a3ec11ff1a..8edc87bc2f 100644 --- a/projects/clr/hipamd/hipify-clang/src/Cuda2Hip.cpp +++ b/projects/clr/hipamd/hipify-clang/src/Cuda2Hip.cpp @@ -362,6 +362,29 @@ public: } } + void RewriteMacroIdentifier(const Token &MacroNameTok) { + std::string macroName = MacroNameTok.getIdentifierInfo()->getName(); + + // TODO: LUT just for macro names, to improve performance. + const auto found = CUDA_RENAMES_MAP().find(macroName); + if (found == CUDA_RENAMES_MAP().end()) { + // Not a CUDA macro. Moving on... + return; + } + + SourceLocation sl = MacroNameTok.getLocation(); + if (found->second.unsupported) { + // We know about it, but it isn't supported. Warn the user. + printHipifyMessage(*_sm, sl, "Unsupported CUDA macro: " + macroName); + return; + } + + StringRef repName = found->second.hipName; + Replacement Rep(*_sm, sl, macroName.size(), repName); + FullSourceLoc fullSL(sl, *_sm); + insertReplacement(Rep, fullSL); + } + virtual void MacroExpands(const Token &MacroNameTok, const MacroDefinition &MD, SourceRange Range, const MacroArgs *Args) override { @@ -370,6 +393,9 @@ public: return; // Macros in headers are not our concern. } + // Is the macro itself a CUDA identifier? If so, rewrite it + RewriteMacroIdentifier(MacroNameTok); + // The getNumArgs function was rather unhelpfully renamed in clang 4.0. Its semantics // remain unchanged. #if LLVM_VERSION_MAJOR > 4 @@ -378,6 +404,7 @@ public: #define GET_NUM_ARGS() getNumArgs() #endif + // 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' From 74af29d66abac888fc30d9345e462053f4598b8e Mon Sep 17 00:00:00 2001 From: Chris Kitching Date: Wed, 18 Oct 2017 15:16:02 +0100 Subject: [PATCH 4/6] Deduplicate preprocessor code There's three functions here that all do the same thing... There was also logic that looks for numeric literals and works backwards to find the macro name from which they are expanded. I previously introduced code that rewrites macro references at expand-time in the `MacroExpands` callback, so that code is no longer doing anything useful. [ROCm/clr commit: 1ef68090ae3e69eaa3bba491d080d109cb163b55] --- .../clr/hipamd/hipify-clang/src/Cuda2Hip.cpp | 157 +++++------------- 1 file changed, 42 insertions(+), 115 deletions(-) diff --git a/projects/clr/hipamd/hipify-clang/src/Cuda2Hip.cpp b/projects/clr/hipamd/hipify-clang/src/Cuda2Hip.cpp index 8edc87bc2f..882b90f4d1 100644 --- a/projects/clr/hipamd/hipify-clang/src/Cuda2Hip.cpp +++ b/projects/clr/hipamd/hipify-clang/src/Cuda2Hip.cpp @@ -323,6 +323,45 @@ public: 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; + } + updateCounters(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()) || @@ -331,60 +370,10 @@ public: } for (auto T : MD->getMacroInfo()->tokens()) { - // We're looking for CUDA identifiers in the macro definition to rewrite... - if (!T.isAnyIdentifier()) { - continue; - } - - StringRef name = T.getIdentifierInfo()->getName(); - const auto found = CUDA_RENAMES_MAP().find(name); - if (found == CUDA_RENAMES_MAP().end()) { - // So it's an identifier that isn't CUDA? Boring. - continue; - } - - updateCounters(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()); - continue; - } - - StringRef repName = found->second.hipName; - DEBUG(dbgs() << "Identifier " << name << " found in definition of macro " - << MacroNameTok.getIdentifierInfo()->getName() << "\n" - << "will be replaced with: " << repName << "\n" - << "SourceLocation: " << sl.printToString(*_sm) << "\n"); - Replacement Rep(*_sm, sl, name.size(), repName); - FullSourceLoc fullSL(sl, *_sm); - insertReplacement(Rep, fullSL); + RewriteToken(T); } } - void RewriteMacroIdentifier(const Token &MacroNameTok) { - std::string macroName = MacroNameTok.getIdentifierInfo()->getName(); - - // TODO: LUT just for macro names, to improve performance. - const auto found = CUDA_RENAMES_MAP().find(macroName); - if (found == CUDA_RENAMES_MAP().end()) { - // Not a CUDA macro. Moving on... - return; - } - - SourceLocation sl = MacroNameTok.getLocation(); - if (found->second.unsupported) { - // We know about it, but it isn't supported. Warn the user. - printHipifyMessage(*_sm, sl, "Unsupported CUDA macro: " + macroName); - return; - } - - StringRef repName = found->second.hipName; - Replacement Rep(*_sm, sl, macroName.size(), repName); - FullSourceLoc fullSL(sl, *_sm); - insertReplacement(Rep, fullSL); - } - virtual void MacroExpands(const Token &MacroNameTok, const MacroDefinition &MD, SourceRange Range, const MacroArgs *Args) override { @@ -394,7 +383,7 @@ public: } // Is the macro itself a CUDA identifier? If so, rewrite it - RewriteMacroIdentifier(MacroNameTok); + RewriteToken(MacroNameTok); // The getNumArgs function was rather unhelpfully renamed in clang 4.0. Its semantics // remain unchanged. @@ -426,69 +415,7 @@ public: // end of stolen code for (auto tok : toks) { - if (tok.isAnyIdentifier()) { - StringRef name = tok.getIdentifierInfo()->getName(); - SourceLocation sl = tok.getLocation(); - - const auto found = CUDA_RENAMES_MAP().find(name); - if (found == CUDA_RENAMES_MAP().end()) { - // It's not a CUDA identifier. We have nothing to do. - continue; - } - - updateCounters(found->second, name.str()); - if (found->second.unsupported) { - // We know about it, but it isn't supported. Warn the user. - printHipifyMessage(*_sm, sl, "Unsupported CUDA identifier: " + name.str()); - continue; - } - - StringRef repName = found->second.hipName; - DEBUG(dbgs() << "Identifier " << name - << " found as an actual argument in expansion of macro " - << MacroNameTok.getIdentifierInfo()->getName() << "\n" - << "will be replaced with: " << repName << "\n"); - size_t length = name.size(); - if (_sm->isMacroBodyExpansion(sl)) { - LangOptions DefaultLangOptions; - SourceLocation sl_macro = _sm->getExpansionLoc(sl); - SourceLocation sl_end = Lexer::getLocForEndOfToken(sl_macro, 0, *_sm, DefaultLangOptions); - length = _sm->getCharacterData(sl_end) - _sm->getCharacterData(sl_macro); - sl = sl_macro; - } - Replacement Rep(*_sm, sl, length, repName); - FullSourceLoc fullSL(sl, *_sm); - insertReplacement(Rep, fullSL); - } else if (tok.isLiteral()) { - SourceLocation sl = tok.getLocation(); - if (_sm->isMacroBodyExpansion(sl)) { - LangOptions DefaultLangOptions; - SourceLocation sl_macro = _sm->getExpansionLoc(sl); - SourceLocation sl_end = Lexer::getLocForEndOfToken(sl_macro, 0, *_sm, DefaultLangOptions); - size_t length = _sm->getCharacterData(sl_end) - _sm->getCharacterData(sl_macro); - StringRef name = StringRef(_sm->getCharacterData(sl_macro), length); - - const auto found = CUDA_RENAMES_MAP().find(name); - if (found == CUDA_RENAMES_MAP().end()) { - continue; // Not CUDA, we don't care. - } - - updateCounters(found->second, name.str()); - if (found->second.unsupported) { - printHipifyMessage(*_sm, sl, "Unsupported CUDA identifier: " + name.str()); - continue; - } - - StringRef repName = found->second.hipName; - sl = sl_macro; - Replacement Rep(*_sm, sl, length, repName); - FullSourceLoc fullSL(sl, *_sm); - insertReplacement(Rep, fullSL); - } else if (tok.is(tok::string_literal)) { - StringRef s(tok.getLiteralData(), tok.getLength()); - processString(unquoteStr(s), *_sm, tok.getLocation()); - } - } + RewriteToken(tok); } } } From 58428d739c800110a9474f4e3cd2ddf03eabad0d Mon Sep 17 00:00:00 2001 From: Chris Kitching Date: Fri, 20 Oct 2017 12:46:39 +0100 Subject: [PATCH 5/6] Simplify how kernel launch expressions get translated It seems like there was a lot of machinery here that is no longer needed now we have hipLaunchKernelGGL (which doesn't require us to insert an extra argument into kernel functions). We no longer need to waste cycles scanning the AST for callees. We can literally just do "Take the callee expression, and dump it into the first argument of hipLaunchKernelGGL()". [ROCm/clr commit: a35d30e0b76fac4f3f35c91f893949baa68578a4] --- .../clr/hipamd/hipify-clang/src/Cuda2Hip.cpp | 65 ++++--------------- 1 file changed, 11 insertions(+), 54 deletions(-) diff --git a/projects/clr/hipamd/hipify-clang/src/Cuda2Hip.cpp b/projects/clr/hipamd/hipify-clang/src/Cuda2Hip.cpp index 882b90f4d1..d5fd29688f 100644 --- a/projects/clr/hipamd/hipify-clang/src/Cuda2Hip.cpp +++ b/projects/clr/hipamd/hipify-clang/src/Cuda2Hip.cpp @@ -435,28 +435,6 @@ private: class Cuda2HipCallback : public MatchFinder::MatchCallback, public Cuda2Hip { private: - void convertKernelDecl(const FunctionDecl *kernelDecl, const MatchFinder::MatchResult &Result) { - SourceManager *SM = Result.SourceManager; - LangOptions DefaultLangOptions; - SmallString<40> XStr; - raw_svector_ostream OS(XStr); - SourceLocation sl = kernelDecl->getNameInfo().getEndLoc(); - SourceLocation kernelArgListStart = Lexer::findLocationAfterToken(sl, tok::l_paren, *SM, DefaultLangOptions, true); - DEBUG(dbgs() << kernelArgListStart.printToString(*SM)); - if (kernelDecl->getNumParams() > 0) { - const ParmVarDecl *pvdFirst = kernelDecl->getParamDecl(0); - const ParmVarDecl *pvdLast = kernelDecl->getParamDecl(kernelDecl->getNumParams() - 1); - SourceLocation kernelArgListStart(pvdFirst->getLocStart()); - SourceLocation kernelArgListEnd(pvdLast->getLocEnd()); - SourceLocation stop = Lexer::getLocForEndOfToken(kernelArgListEnd, 0, *SM, DefaultLangOptions); - size_t repLength = SM->getCharacterData(stop) - SM->getCharacterData(kernelArgListStart); - OS << StringRef(SM->getCharacterData(kernelArgListStart), repLength); - Replacement Rep0(*(Result.SourceManager), kernelArgListStart, repLength, OS.str()); - FullSourceLoc fullSL(sl, *(Result.SourceManager)); - insertReplacement(Rep0, fullSL); - } - } - bool cudaCall(const MatchFinder::MatchResult &Result) { const CallExpr *call = Result.Nodes.getNodeAs("cudaCall"); if (!call) { @@ -498,30 +476,19 @@ private: if (const CUDAKernelCallExpr *launchKernel = Result.Nodes.getNodeAs(refName)) { SmallString<40> XStr; raw_svector_ostream OS(XStr); - StringRef calleeName; - const FunctionDecl *kernelDecl = launchKernel->getDirectCallee(); - if (kernelDecl) { - calleeName = kernelDecl->getName(); - convertKernelDecl(kernelDecl, Result); - } else { - const Expr *e = launchKernel->getCallee(); - if (const UnresolvedLookupExpr *ule = - dyn_cast(e)) { - calleeName = ule->getName().getAsIdentifierInfo()->getName(); - owner->addMatcher(functionTemplateDecl(hasName(calleeName)) - .bind("unresolvedTemplateName"), this); - } - } - XStr.clear(); - if (calleeName.find(',') != StringRef::npos) { - SmallString<128> tmpData; - calleeName = Twine("(" + calleeName + ")").toStringRef(tmpData); - } - OS << "hipLaunchKernelGGL(" << calleeName << ","; + + LangOptions DefaultLangOptions; + SourceManager *SM = Result.SourceManager; + + const Expr *e = launchKernel->getCallee(); + + // Grab the characters for the callee expression and dump them into hipLaunchKernelGGL's + // first argument. + StringRef exprSource = Lexer::getSourceText(CharSourceRange::getTokenRange(e->getSourceRange()), *SM, LangOptions(), 0); + + OS << "hipLaunchKernelGGL(" << exprSource << ","; const CallExpr *config = launchKernel->getConfig(); DEBUG(dbgs() << "Kernel config arguments:" << "\n"); - SourceManager *SM = Result.SourceManager; - LangOptions DefaultLangOptions; for (unsigned argno = 0; argno < config->getNumArgs(); argno++) { const Expr *arg = config->getArg(argno); if (!isa(arg)) { @@ -724,15 +691,6 @@ private: return false; } - bool unresolvedTemplateName(const MatchFinder::MatchResult &Result) { - if (const FunctionTemplateDecl *templateDecl = Result.Nodes.getNodeAs("unresolvedTemplateName")) { - FunctionDecl *kernelDecl = templateDecl->getTemplatedDecl(); - convertKernelDecl(kernelDecl, Result); - return true; - } - return false; - } - bool stringLiteral(const MatchFinder::MatchResult &Result) { if (const clang::StringLiteral *sLiteral = Result.Nodes.getNodeAs("stringLiteral")) { if (sLiteral->getCharByteWidth() == 1) { @@ -759,7 +717,6 @@ public: if (cudaLaunchKernel(Result)) return; if (cudaSharedIncompleteArrayVar(Result)) return; if (stringLiteral(Result)) return; - if (unresolvedTemplateName(Result)) return; } private: From 7360326705fa0a0713fec452b1caba2081421dfd Mon Sep 17 00:00:00 2001 From: Chris Kitching Date: Fri, 20 Oct 2017 13:10:31 +0100 Subject: [PATCH 6/6] Greatly enhance handling of macros in kernel launches All but the most contrived use of macros is now properly handled - have a look at the new testcases this commit adds. You can have macros in kernel calls, macros spanning chunks of your arguments, the call, call parameters, or callee can all be macros or partially macros. [ROCm/clr commit: 6491c2c3eb78f29876a5415fadc168fa5aefef7c] --- .../clr/hipamd/hipify-clang/src/Cuda2Hip.cpp | 126 +++++++++++------- .../clr/hipamd/tests/hipify-clang/axpy.cu | 35 ++++- 2 files changed, 111 insertions(+), 50 deletions(-) diff --git a/projects/clr/hipamd/hipify-clang/src/Cuda2Hip.cpp b/projects/clr/hipamd/hipify-clang/src/Cuda2Hip.cpp index d5fd29688f..fd23344f9f 100644 --- a/projects/clr/hipamd/hipify-clang/src/Cuda2Hip.cpp +++ b/projects/clr/hipamd/hipify-clang/src/Cuda2Hip.cpp @@ -471,6 +471,49 @@ private: 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)) { @@ -480,59 +523,46 @@ private: LangOptions DefaultLangOptions; SourceManager *SM = Result.SourceManager; - const Expr *e = launchKernel->getCallee(); + const Expr& calleeExpr = *(launchKernel->getCallee()); + OS << "hipLaunchKernelGGL(" << readSourceText(*SM, calleeExpr.getSourceRange()) << ", "; - // Grab the characters for the callee expression and dump them into hipLaunchKernelGGL's - // first argument. - StringRef exprSource = Lexer::getSourceText(CharSourceRange::getTokenRange(e->getSourceRange()), *SM, LangOptions(), 0); + // Next up are the four kernel configuration parameters, the last two of which are optional and default to zero. + const CallExpr& config = *(launchKernel->getConfig()); - OS << "hipLaunchKernelGGL(" << exprSource << ","; - const CallExpr *config = launchKernel->getConfig(); - DEBUG(dbgs() << "Kernel config arguments:" << "\n"); - for (unsigned argno = 0; argno < config->getNumArgs(); argno++) { - const Expr *arg = config->getArg(argno); - if (!isa(arg)) { - const ParmVarDecl *pvd = config->getDirectCallee()->getParamDecl(argno); - SourceLocation sl(arg->getLocStart()); - SourceLocation el(arg->getLocEnd()); - SourceLocation stop = Lexer::getLocForEndOfToken(el, 0, *SM, DefaultLangOptions); - StringRef outs(SM->getCharacterData(sl), SM->getCharacterData(stop) - SM->getCharacterData(sl)); - DEBUG(dbgs() << "args[ " << argno << "]" << outs << " <" << pvd->getType().getAsString() << ">\n"); - if (pvd->getType().getAsString().compare("dim3") == 0) { - OS << " dim3(" << outs << "),"; - } else { - OS << " " << outs << ","; - } - } else { - OS << " 0,"; - } + // 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}); } - for (unsigned argno = 0; argno < launchKernel->getNumArgs(); argno++) { - const Expr *arg = launchKernel->getArg(argno); - SourceLocation sl(arg->getLocStart()); - if (SM->isMacroBodyExpansion(sl)) { - sl = SM->getExpansionLoc(sl); - } else if (SM->isMacroArgExpansion(sl)) { - sl = SM->getImmediateSpellingLoc(sl); - } - SourceLocation el(arg->getLocEnd()); - if (SM->isMacroBodyExpansion(el)) { - el = SM->getExpansionLoc(el); - } else if (SM->isMacroArgExpansion(el)) { - el = SM->getImmediateSpellingLoc(el); - } - SourceLocation stop = Lexer::getLocForEndOfToken(el, 0, *SM, DefaultLangOptions); - std::string outs(SM->getCharacterData(sl), SM->getCharacterData(stop) - SM->getCharacterData(sl)); - DEBUG(dbgs() << outs << "\n"); - OS << " " << outs << ","; - } - XStr.pop_back(); + OS << ")"; + + SourceRange replacementRange = getWriteRange(*SM, {launchKernel->getLocStart(), launchKernel->getLocEnd()}); + SourceLocation launchStart = replacementRange.getBegin(); + SourceLocation launchEnd = replacementRange.getEnd(); + size_t length = SM->getCharacterData(Lexer::getLocForEndOfToken( - launchKernel->getLocEnd(), 0, *SM, DefaultLangOptions)) - - SM->getCharacterData(launchKernel->getLocStart()); - Replacement Rep(*SM, launchKernel->getLocStart(), length, OS.str()); - FullSourceLoc fullSL(launchKernel->getLocStart(), *SM); + launchEnd, 0, *SM, DefaultLangOptions)) - + SM->getCharacterData(launchStart); + + Replacement Rep(*SM, launchStart, length, OS.str()); + FullSourceLoc fullSL(launchStart, *SM); insertReplacement(Rep, fullSL); hipCounter counter = {"hipLaunchKernelGGL", CONV_KERN, API_RUNTIME}; updateCounters(counter, refName.str()); diff --git a/projects/clr/hipamd/tests/hipify-clang/axpy.cu b/projects/clr/hipamd/tests/hipify-clang/axpy.cu index 8c6b0e0d8d..2fd62ac344 100644 --- a/projects/clr/hipamd/tests/hipify-clang/axpy.cu +++ b/projects/clr/hipamd/tests/hipify-clang/axpy.cu @@ -2,11 +2,23 @@ #include -__global__ void axpy(float a, float* x, float* y) { + +#define TOKEN_PASTE(X, Y) X ## Y +#define ARG_LIST_AS_MACRO a, device_x, device_y +#define KERNEL_CALL_AS_MACRO axpy<<<1, kDataLen>>> +#define KERNEL_NAME_MACRO axpy + +// CHECK: #define COMPLETE_LAUNCH hipLaunchKernelGGL(axpy, dim3(1), dim3(kDataLen), 0, 0, a, device_x, device_y) +#define COMPLETE_LAUNCH axpy<<<1, kDataLen>>>(a, device_x, device_y) + + +template +__global__ void axpy(T a, T *x, T *y) { // CHECK: y[hipThreadIdx_x] = a * x[hipThreadIdx_x]; y[threadIdx.x] = a * x[threadIdx.x]; } + int main(int argc, char* argv[]) { const int kDataLen = 4; @@ -27,10 +39,29 @@ int main(int argc, char* argv[]) { // CHECK: hipMemcpy(device_x, host_x, kDataLen * sizeof(float), hipMemcpyHostToDevice); cudaMemcpy(device_x, host_x, kDataLen * sizeof(float), cudaMemcpyHostToDevice); - // Launch the kernel. + // Launch the kernel in numerous different strange ways to exercise the prerocessor. // CHECK: hipLaunchKernelGGL(axpy, dim3(1), dim3(kDataLen), 0, 0, a, device_x, device_y); axpy<<<1, kDataLen>>>(a, device_x, device_y); + // CHECK: hipLaunchKernelGGL(axpy, dim3(1), dim3(kDataLen), 0, 0, a, device_x, device_y); + axpy<<<1, kDataLen>>>(a, device_x, device_y); + + // CHECK: hipLaunchKernelGGL(axpy, dim3(1), dim3(kDataLen), 0, 0, a, TOKEN_PASTE(device, _x), device_y); + axpy<<<1, kDataLen>>>(a, TOKEN_PASTE(device, _x), device_y); + + // CHECK: hipLaunchKernelGGL(axpy, dim3(1), dim3(kDataLen), 0, 0, ARG_LIST_AS_MACRO); + axpy<<<1, kDataLen>>>(ARG_LIST_AS_MACRO); + + // CHECK: hipLaunchKernelGGL(KERNEL_NAME_MACRO, dim3(1), dim3(kDataLen), 0, 0, ARG_LIST_AS_MACRO); + KERNEL_NAME_MACRO<<<1, kDataLen>>>(ARG_LIST_AS_MACRO); + + // CHECK: hipLaunchKernelGGL(axpy, dim3(1), dim3(kDataLen), 0, 0, ARG_LIST_AS_MACRO); + KERNEL_CALL_AS_MACRO(ARG_LIST_AS_MACRO); + + // CHECK: COMPLETE_LAUNCH; + COMPLETE_LAUNCH; + + // Copy output data to host. // CHECK: hipDeviceSynchronize(); cudaDeviceSynchronize();