AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

[HIPIFY] undo apply .clangformat to HIPIFY source files

Browse Source 
[ROCm/hip commit: 8499e79e2b]
This commit is contained in:
Evgeny Mankov
2018-04-23 20:22:15 +03:00
parent f117777f77
commit 540fc92b1c
12 changed files with 2302 additions and 3995 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/hip/hipify-clang/src/HipifyAction.cpp
+55 -79
View File
@@ -25,8 +25,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 = found->second.hipName;
hipCounter counter = {"[string literal]", ConvTypes::CONV_LITERAL,
ApiTypes::API_RUNTIME, found->second.unsupported};
hipCounter counter = {"[string literal]", ConvTypes::CONV_LITERAL, ApiTypes::API_RUNTIME, found->second.unsupported};
Statistics::current().incrementCounter(counter, name.str());
if (!counter.unsupported) {
@@ -78,8 +77,7 @@ void HipifyAction::RewriteToken(const clang::Token& t) {
if (found->second.unsupported) {
// An unsupported identifier? Curses! Warn the user.
clang::DiagnosticsEngine& DE = getCompilerInstance().getDiagnostics();
const auto ID = DE.getCustomDiagID(clang::DiagnosticsEngine::Warning,
"CUDA identifier unsupported in hip");
const auto ID = DE.getCustomDiagID(clang::DiagnosticsEngine::Warning, "CUDA identifier unsupported in hip");
DE.Report(sl, ID);
return;
}
@@ -96,10 +94,8 @@ clang::SourceRange getReadRange(clang::SourceManager& SM, const clang::SourceRan
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{});
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)};
@@ -124,9 +120,7 @@ clang::SourceRange getWriteRange(clang::SourceManager& SM, const clang::SourceRa
StringRef readSourceText(clang::SourceManager& SM, const clang::SourceRange& exprRange) {
return clang::Lexer::getSourceText(
clang::CharSourceRange::getTokenRange(getReadRange(SM, exprRange)), SM,
clang::LangOptions(), nullptr);
return clang::Lexer::getSourceText(clang::CharSourceRange::getTokenRange(getReadRange(SM, exprRange)), SM, clang::LangOptions(), nullptr);
}
/**
@@ -141,36 +135,28 @@ std::string stringifyZeroDefaultedArg(clang::SourceManager& SM, const clang::Exp
}
}
} // anonymous namespace
} // 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 (insertedRAND_kernelHeader) { return true; }
insertedRAND_kernelHeader = true;
return false;
} else if (hipToken.hipName == "hiprand.h") {
if (insertedRANDHeader) {
return true;
}
if (insertedRANDHeader) { return true; }
insertedRANDHeader = true;
return false;
}
@@ -181,9 +167,7 @@ bool HipifyAction::Exclude(const hipCounter& hipToken) {
case CONV_INCLUDE:
switch (hipToken.apiType) {
case API_RAND:
if (insertedRAND_kernelHeader) {
return true;
}
if (insertedRAND_kernelHeader) { return true; }
insertedRAND_kernelHeader = true;
return false;
default:
@@ -196,11 +180,13 @@ bool HipifyAction::Exclude(const hipCounter& hipToken) {
return false;
}
void HipifyAction::InclusionDirective(clang::SourceLocation hash_loc, const clang::Token&,
StringRef file_name, bool is_angled,
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*) {
const clang::FileEntry*, StringRef,
StringRef, const clang::Module*) {
clang::SourceManager& SM = getCompilerInstance().getSourceManager();
if (!SM.isWrittenInMainFile(hash_loc)) {
return;
@@ -222,8 +208,7 @@ void HipifyAction::InclusionDirective(clang::SourceLocation hash_loc, const clan
clang::SourceLocation sl = filename_range.getBegin();
if (found->second.unsupported) {
clang::DiagnosticsEngine& DE = getCompilerInstance().getDiagnostics();
DE.Report(sl,
DE.getCustomDiagID(clang::DiagnosticsEngine::Warning, "Unsupported CUDA header"));
DE.Report(sl, DE.getCustomDiagID(clang::DiagnosticsEngine::Warning, "Unsupported CUDA header"));
return;
}
@@ -235,22 +220,19 @@ void HipifyAction::InclusionDirective(clang::SourceLocation hash_loc, const clan
if (is_angled) {
newInclude = llvm::Twine("<" + found->second.hipName + ">").toStringRef(includeBuffer);
} else {
newInclude =
llvm::Twine("\"" + found->second.hipName + "\"").toStringRef(includeBuffer);
newInclude = llvm::Twine("\"" + found->second.hipName + "\"").toStringRef(includeBuffer);
}
} else {
// hashLoc is location of the '#', thus replacing the whole include directive by empty
// newInclude starting with '#'.
// hashLoc is location of the '#', thus replacing the whole include directive by empty newInclude starting with '#'.
sl = hash_loc;
}
const char* B = SM.getCharacterData(sl);
const char* E = SM.getCharacterData(filename_range.getEnd());
const char *B = SM.getCharacterData(sl);
const char *E = SM.getCharacterData(filename_range.getEnd());
ct::Replacement Rep(SM, sl, E - B, newInclude);
insertReplacement(Rep, clang::FullSourceLoc{sl, SM});
}
void HipifyAction::PragmaDirective(clang::SourceLocation Loc,
clang::PragmaIntroducerKind Introducer) {
void HipifyAction::PragmaDirective(clang::SourceLocation Loc, clang::PragmaIntroducerKind Introducer) {
if (pragmaOnce) {
return;
}
@@ -284,8 +266,7 @@ bool HipifyAction::cudaLaunchKernel(const clang::ast_matchers::MatchFinder::Matc
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.
// 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.
@@ -312,14 +293,11 @@ bool HipifyAction::cudaLaunchKernel(const clang::ast_matchers::MatchFinder::Matc
OS << ")";
clang::SourceRange replacementRange =
getWriteRange(*SM, {launchKernel->getLocStart(), launchKernel->getLocEnd()});
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);
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);
@@ -330,8 +308,7 @@ bool HipifyAction::cudaLaunchKernel(const clang::ast_matchers::MatchFinder::Matc
return true;
}
bool HipifyAction::cudaSharedIncompleteArrayVar(
const clang::ast_matchers::MatchFinder::MatchResult& Result) {
bool HipifyAction::cudaSharedIncompleteArrayVar(const clang::ast_matchers::MatchFinder::MatchResult& Result) {
StringRef refName = "cudaSharedIncompleteArrayVar";
auto* sharedVar = Result.Nodes.getNodeAs<clang::VarDecl>(refName);
if (!sharedVar) {
@@ -379,8 +356,7 @@ bool HipifyAction::cudaSharedIncompleteArrayVar(
return true;
}
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();
@@ -389,19 +365,22 @@ void HipifyAction::insertReplacement(const ct::Replacement& rep,
}
}
std::unique_ptr<clang::ASTConsumer> HipifyAction::CreateASTConsumer(clang::CompilerInstance& CI,
llvm::StringRef) {
std::unique_ptr<clang::ASTConsumer> 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);
Finder->addMatcher(mat::cudaKernelCallExpr(mat::isExpansionInMainFile()).bind("cudaLaunchKernel"), this);
Finder->addMatcher(mat::varDecl(mat::isExpansionInMainFile(),
mat::allOf(mat::hasAttr(clang::attr::CUDAShared),
mat::hasType(mat::incompleteArrayType())))
.bind("cudaSharedIncompleteArrayVar"),
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();
@@ -410,9 +389,9 @@ std::unique_ptr<clang::ASTConsumer> HipifyAction::CreateASTConsumer(clang::Compi
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.
// 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;
if (pragmaOnce) {
@@ -439,25 +418,22 @@ namespace {
class PPCallbackProxy : public clang::PPCallbacks {
HipifyAction& hipifyAction;
public:
explicit PPCallbackProxy(HipifyAction& action) : hipifyAction(action) {}
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,
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);
hipifyAction.InclusionDirective(hash_loc, include_token, file_name, is_angled, filename_range, file, search_path, relative_path, imported);
}
void PragmaDirective(clang::SourceLocation Loc,
clang::PragmaIntroducerKind Introducer) override {
void PragmaDirective(clang::SourceLocation Loc, clang::PragmaIntroducerKind Introducer) override {
hipifyAction.PragmaDirective(Loc, Introducer);
}
};
} // namespace
}
void HipifyAction::ExecuteAction() {
clang::Preprocessor& PP = getCompilerInstance().getPreprocessor();
@@ -468,10 +444,10 @@ void HipifyAction::ExecuteAction() {
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.
// 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)) {