[HIPIFY] undo apply .clangformat to HIPIFY source files

Этот коммит содержится в:
Evgeny Mankov
2018-04-23 20:22:15 +03:00
родитель 35a8deec00
Коммит 8499e79e2b
12 изменённых файлов: 2302 добавлений и 3995 удалений
+27 -15
Просмотреть файл
@@ -2,27 +2,39 @@
cl::OptionCategory ToolTemplateCategory("CUDA to HIP source translator options");
cl::opt<std::string> OutputFilename("o", cl::desc("Output filename"), cl::value_desc("filename"),
cl::cat(ToolTemplateCategory));
cl::opt<std::string> OutputFilename("o",
cl::desc("Output filename"),
cl::value_desc("filename"),
cl::cat(ToolTemplateCategory));
cl::opt<bool> 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::desc("Modify input file inplace, replacing input with hipified output, save backup in .prehip file"),
cl::value_desc("inplace"),
cl::cat(ToolTemplateCategory));
cl::opt<bool> 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<bool> 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<bool> NoOutput("no-output", cl::desc("Don't write any translated output to stdout"),
cl::value_desc("no-output"), cl::cat(ToolTemplateCategory));
cl::opt<bool> NoOutput("no-output",
cl::desc("Don't write any translated output to stdout"),
cl::value_desc("no-output"),
cl::cat(ToolTemplateCategory));
cl::opt<bool> PrintStats("print-stats", cl::desc("Print translation statistics"),
cl::value_desc("print-stats"), cl::cat(ToolTemplateCategory));
cl::opt<bool> PrintStats("print-stats",
cl::desc("Print translation statistics"),
cl::value_desc("print-stats"),
cl::cat(ToolTemplateCategory));
cl::opt<std::string> OutputStatsFilename("o-stats", cl::desc("Output filename for statistics"),
cl::value_desc("filename"), cl::cat(ToolTemplateCategory));
cl::opt<std::string> OutputStatsFilename("o-stats",
cl::desc("Output filename for statistics"),
cl::value_desc("filename"),
cl::cat(ToolTemplateCategory));
cl::opt<bool> Examine("examine", cl::desc("Combines -no-output and -print-stats options"),
cl::value_desc("examine"), cl::cat(ToolTemplateCategory));
cl::opt<bool> Examine("examine",
cl::desc("Combines -no-output and -print-stats options"),
cl::value_desc("examine"),
cl::cat(ToolTemplateCategory));
cl::extrahelp CommonHelp(ct::CommonOptionsParser::HelpMessage);
Разница между файлами не показана из-за своего большого размера Загрузить разницу
+55 -79
Просмотреть файл
@@ -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)) {
+27 -25
Просмотреть файл
@@ -15,15 +15,14 @@ namespace ct = clang::tooling;
*/
class HipifyAction : public clang::ASTFrontendAction,
public clang::ast_matchers::MatchFinder::MatchCallback {
private:
private:
ct::Replacements* replacements;
std::unique_ptr<clang::ast_matchers::MatchFinder> 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.
/// 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;
bool insertedBLASHeader = false;
bool insertedRANDHeader = false;
@@ -41,11 +40,12 @@ class HipifyAction : public clang::ASTFrontendAction,
/**
* Replace a CUDA identifier with the corresponding hip identifier, if applicable.
*/
void RewriteToken(const clang::Token& t);
void RewriteToken(const clang::Token &t);
public:
explicit HipifyAction(ct::Replacements* replacements)
: clang::ASTFrontendAction(), replacements(replacements) {}
public:
explicit HipifyAction(ct::Replacements *replacements):
clang::ASTFrontendAction(),
replacements(replacements) {}
// MatchCallback listeners
bool cudaBuiltin(const clang::ast_matchers::MatchFinder::MatchResult& Result);
@@ -55,21 +55,24 @@ class HipifyAction : public clang::ASTFrontendAction,
/**
* 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);
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);
/**
* Called by the preprocessor for each pragma directive during the non-raw lexing pass.
*/
* Called by the preprocessor for each pragma directive during the non-raw lexing pass.
*/
void PragmaDirective(clang::SourceLocation Loc, clang::PragmaIntroducerKind Introducer);
protected:
protected:
/**
* Add a Replacement for the current file. These will all be applied after executing the
* FrontendAction.
* 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);
@@ -84,13 +87,12 @@ class HipifyAction : public clang::ASTFrontendAction,
void EndSourceFileAction() override;
/**
* MatchCallback API entry point. Called by the AST visitor while searching the AST for things
* we registered an interest for.
* 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<clang::ASTConsumer> CreateASTConsumer(clang::CompilerInstance& CI,
llvm::StringRef InFile) override;
std::unique_ptr<clang::ASTConsumer> CreateASTConsumer(clang::CompilerInstance &CI, llvm::StringRef InFile) override;
bool Exclude(const hipCounter& hipToken);
bool Exclude(const hipCounter & hipToken);
};
+2 -3
Просмотреть файл
@@ -32,8 +32,7 @@ void insertReplacement(ct::Replacements& replacements, const ct::Replacement& re
#endif
}
void EnterPreprocessorTokenStream(clang::Preprocessor& _pp, const clang::Token* start, size_t len,
bool DisableMacroExpansion) {
void EnterPreprocessorTokenStream(clang::Preprocessor& _pp, const clang::Token *start, size_t len, bool DisableMacroExpansion) {
#if (LLVM_VERSION_MAJOR == 3) && (LLVM_VERSION_MINOR == 8)
_pp.EnterTokenStream(start, len, false, DisableMacroExpansion);
#else
@@ -41,4 +40,4 @@ void EnterPreprocessorTokenStream(clang::Preprocessor& _pp, const clang::Token*
#endif
}
} // namespace llcompat
} // namespace llcompat
+6 -4
Просмотреть файл
@@ -18,9 +18,9 @@ namespace llcompat {
* remain unchanged, so let's be slightly ugly about it here. :D
*/
#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
void PrintStackTraceOnErrorSignal();
@@ -41,7 +41,9 @@ void insertReplacement(ct::Replacements& replacements, const ct::Replacement& re
/**
* Version-agnostic version of Preprocessor::EnterTokenStream().
*/
void EnterPreprocessorTokenStream(clang::Preprocessor& _pp, const clang::Token* start, size_t len,
void EnterPreprocessorTokenStream(clang::Preprocessor& _pp,
const clang::Token *start,
size_t len,
bool DisableMacroExpansion);
} // namespace llcompat
} // namespace llcompat
+7 -4
Просмотреть файл
@@ -17,9 +17,12 @@ template <typename T>
class ReplacementsFrontendActionFactory : public ct::FrontendActionFactory {
ct::Replacements* replacements;
public:
explicit ReplacementsFrontendActionFactory(ct::Replacements* r)
: ct::FrontendActionFactory(), replacements(r) {}
public:
explicit ReplacementsFrontendActionFactory(ct::Replacements* r):
ct::FrontendActionFactory(),
replacements(r) {}
clang::FrontendAction* create() override { return new T(replacements); }
clang::FrontendAction* create() override {
return new T(replacements);
}
};
+69 -72
Просмотреть файл
@@ -4,26 +4,26 @@
#include <iomanip>
const char* counterNames[NUM_CONV_TYPES] = {"version", "init", "device",
"mem", "kern", "coord_func",
"math_func", "device_func", "special_func",
"stream", "event", "occupancy",
"ctx", "peer", "module",
"cache", "exec", "err",
"def", "tex", "gl",
"graphics", "surface", "jit",
"d3d9", "d3d10", "d3d11",
"vdpau", "egl", "thread",
"other", "include", "include_cuda_main_header",
"type", "literal", "numeric_literal"};
const char *counterNames[NUM_CONV_TYPES] = {
"version", "init", "device", "mem", "kern", "coord_func", "math_func", "device_func",
"special_func", "stream", "event", "occupancy", "ctx", "peer", "module",
"cache", "exec", "err", "def", "tex", "gl", "graphics",
"surface", "jit", "d3d9", "d3d10", "d3d11", "vdpau", "egl",
"thread", "other", "include", "include_cuda_main_header", "type", "literal",
"numeric_literal"
};
const char* apiNames[NUM_API_TYPES] = {"CUDA Driver API", "CUDA RT API", "CUBLAS API",
"CURAND API"};
const char *apiNames[NUM_API_TYPES] = {
"CUDA Driver API", "CUDA RT API", "CUBLAS API", "CURAND API"
};
namespace {
template <typename ST, typename ST2>
void conditionalPrint(ST* stream1, ST2* stream2, const std::string& s1, const std::string& s2) {
template<typename ST, typename ST2>
void conditionalPrint(ST *stream1,
ST2* stream2,
const std::string& s1,
const std::string& s2) {
if (stream1) {
*stream1 << s1;
}
@@ -37,8 +37,8 @@ void conditionalPrint(ST* stream1, ST2* stream2, const std::string& s1, const st
/**
* Print a named stat value to both the terminal and the CSV file.
*/
template <typename T>
void printStat(std::ostream* csv, llvm::raw_ostream* printOut, const std::string& name, T value) {
template<typename T>
void printStat(std::ostream *csv, llvm::raw_ostream* printOut, const std::string &name, T value) {
if (printOut) {
*printOut << " " << name << ": " << value << "\n";
}
@@ -49,12 +49,12 @@ void printStat(std::ostream* csv, llvm::raw_ostream* printOut, const std::string
}
} // Anonymous namespace
} // Anonymous namespace
void StatCounter::incrementCounter(const hipCounter& counter, std::string name) {
counters[name]++;
apiCounters[(int)counter.apiType]++;
convTypeCounters[(int)counter.type]++;
apiCounters[(int) counter.apiType]++;
convTypeCounters[(int) counter.type]++;
}
void StatCounter::add(const StatCounter& other) {
@@ -81,36 +81,32 @@ int StatCounter::getConvSum() {
}
void StatCounter::print(std::ostream* csv, llvm::raw_ostream* printOut, std::string prefix) {
conditionalPrint(csv, printOut, "\nCUDA ref type;Count\n",
"[HIPIFY] info: " + prefix + " refs by type:\n");
conditionalPrint(csv, printOut, "\nCUDA ref type;Count\n", "[HIPIFY] info: " + prefix + " refs by type:\n");
for (int i = 0; i < NUM_CONV_TYPES; i++) {
if (convTypeCounters[i] > 0) {
printStat(csv, printOut, counterNames[i], convTypeCounters[i]);
}
}
conditionalPrint(csv, printOut, "\nCUDA API;Count\n",
"[HIPIFY] info: " + prefix + " refs by API:\n");
conditionalPrint(csv, printOut, "\nCUDA API;Count\n", "[HIPIFY] info: " + prefix + " refs by API:\n");
for (int i = 0; i < NUM_API_TYPES; i++) {
printStat(csv, printOut, apiNames[i], apiCounters[i]);
}
conditionalPrint(csv, printOut, "\nCUDA ref name;Count\n",
"[HIPIFY] info: " + prefix + " refs by names:\n");
for (const auto& it : counters) {
conditionalPrint(csv, printOut, "\nCUDA ref name;Count\n", "[HIPIFY] info: " + prefix + " refs by names:\n");
for (const auto &it : counters) {
printStat(csv, printOut, it.first, it.second);
}
}
Statistics::Statistics(std::string name) : fileName(name) {
Statistics::Statistics(std::string name): fileName(name) {
// Compute the total bytes/lines in the input file.
std::ifstream src_file(name, std::ios::binary | std::ios::ate);
src_file.clear();
src_file.seekg(0);
totalLines = (int)std::count(std::istreambuf_iterator<char>(src_file),
std::istreambuf_iterator<char>(), '\n');
totalBytes = (int)src_file.tellg();
totalLines = (int) std::count(std::istreambuf_iterator<char>(src_file), std::istreambuf_iterator<char>(), '\n');
totalBytes = (int) src_file.tellg();
// Mark the start time...
startTime = chr::steady_clock::now();
@@ -119,7 +115,7 @@ Statistics::Statistics(std::string name) : fileName(name) {
///////// Counter update routines //////////
void Statistics::incrementCounter(const hipCounter& counter, std::string name) {
void Statistics::incrementCounter(const hipCounter &counter, std::string name) {
if (counter.unsupported) {
unsupported.incrementCounter(counter, name);
} else {
@@ -127,7 +123,7 @@ void Statistics::incrementCounter(const hipCounter& counter, std::string name) {
}
}
void Statistics::add(const Statistics& other) {
void Statistics::add(const Statistics &other) {
supported.add(other.supported);
unsupported.add(other.unsupported);
totalBytes += other.totalBytes;
@@ -135,60 +131,61 @@ void Statistics::add(const Statistics& other) {
touchedBytes += other.touchedBytes;
}
void Statistics::lineTouched(int lineNumber) { touchedLines.insert(lineNumber); }
void Statistics::bytesChanged(int bytes) { touchedBytes += bytes; }
void Statistics::markCompletion() { completionTime = chr::steady_clock::now(); }
void Statistics::lineTouched(int lineNumber) {
touchedLines.insert(lineNumber);
}
void Statistics::bytesChanged(int bytes) {
touchedBytes += bytes;
}
void Statistics::markCompletion() {
completionTime = chr::steady_clock::now();
}
///////// Output functions //////////
void Statistics::print(std::ostream* csv, llvm::raw_ostream* printOut, bool skipHeader) {
if (!skipHeader) {
std::string str = "file \'" + fileName + "\' statistics:\n";
conditionalPrint(csv, printOut, "\n" + str, "\n[HIPIFY] info: " + str);
}
if (!skipHeader) {
std::string str = "file \'" + fileName + "\' statistics:\n";
conditionalPrint(csv, printOut, "\n" + str, "\n[HIPIFY] info: " + str);
}
size_t changedLines = touchedLines.size();
size_t changedLines = touchedLines.size();
// Total number of (un)supported refs that were converted.
int supportedSum = supported.getConvSum();
int unsupportedSum = unsupported.getConvSum();
// Total number of (un)supported refs that were converted.
int supportedSum = supported.getConvSum();
int unsupportedSum = unsupported.getConvSum();
printStat(csv, printOut, "CONVERTED refs count", supportedSum);
printStat(csv, printOut, "UNCONVERTED refs count", unsupportedSum);
printStat(
csv, printOut, "CONVERSION %",
100 - std::lround(double(unsupportedSum * 100) / double(supportedSum + unsupportedSum)));
printStat(csv, printOut, "REPLACED bytes", touchedBytes);
printStat(csv, printOut, "TOTAL bytes", totalBytes);
printStat(csv, printOut, "CHANGED lines of code", changedLines);
printStat(csv, printOut, "TOTAL lines of code", totalLines);
printStat(csv, printOut, "CONVERTED refs count", supportedSum);
printStat(csv, printOut, "UNCONVERTED refs count", unsupportedSum);
printStat(csv, printOut, "CONVERSION %", 100 - std::lround(double(unsupportedSum * 100) / double(supportedSum + unsupportedSum)));
printStat(csv, printOut, "REPLACED bytes", touchedBytes);
printStat(csv, printOut, "TOTAL bytes", totalBytes);
printStat(csv, printOut, "CHANGED lines of code", changedLines);
printStat(csv, printOut, "TOTAL lines of code", totalLines);
if (totalBytes > 0) {
printStat(csv, printOut, "CODE CHANGED (in bytes) %",
std::lround(double(touchedBytes * 100) / double(totalBytes)));
}
if (totalBytes > 0) {
printStat(csv, printOut, "CODE CHANGED (in bytes) %", std::lround(double(touchedBytes * 100) / double(totalBytes)));
}
if (totalLines > 0) {
printStat(csv, printOut, "CODE CHANGED (in lines) %",
std::lround(double(changedLines * 100) / double(totalLines)));
}
if (totalLines > 0) {
printStat(csv, printOut, "CODE CHANGED (in lines) %", std::lround(double(changedLines * 100) / double(totalLines)));
}
typedef std::chrono::duration<double, std::milli> duration;
duration elapsed = completionTime - startTime;
std::stringstream stream;
stream << std::fixed << std::setprecision(2) << elapsed.count() / 1000;
printStat(csv, printOut, "TIME ELAPSED s", stream.str());
typedef std::chrono::duration<double, std::milli> duration;
duration elapsed = completionTime - startTime;
std::stringstream stream;
stream << std::fixed << std::setprecision(2) << elapsed.count() / 1000;
printStat(csv, printOut, "TIME ELAPSED s", stream.str());
supported.print(csv, printOut, "CONVERTED");
unsupported.print(csv, printOut, "UNCONVERTED");
supported.print(csv, printOut, "CONVERTED");
unsupported.print(csv, printOut, "UNCONVERTED");
}
void Statistics::printAggregate(std::ostream* csv, llvm::raw_ostream* printOut) {
void Statistics::printAggregate(std::ostream *csv, llvm::raw_ostream* printOut) {
Statistics globalStats = getAggregate();
conditionalPrint(csv, printOut, "\nTOTAL statistics:\n",
"\n[HIPIFY] info: TOTAL statistics:\n");
conditionalPrint(csv, printOut, "\nTOTAL statistics:\n", "\n[HIPIFY] info: TOTAL statistics:\n");
// A file is considered "converted" if we made any changes to it.
int convertedFiles = 0;
+26 -20
Просмотреть файл
@@ -49,14 +49,20 @@ enum ConvTypes {
CONV_NUMERIC_LITERAL,
CONV_LAST
};
constexpr int NUM_CONV_TYPES = (int)ConvTypes::CONV_LAST;
constexpr int NUM_CONV_TYPES = (int) ConvTypes::CONV_LAST;
enum ApiTypes { API_DRIVER = 0, API_RUNTIME, API_BLAS, API_RAND, API_LAST };
constexpr int NUM_API_TYPES = (int)ApiTypes::API_LAST;
enum ApiTypes {
API_DRIVER = 0,
API_RUNTIME,
API_BLAS,
API_RAND,
API_LAST
};
constexpr int NUM_API_TYPES = (int) ApiTypes::API_LAST;
// The names of various fields in in the statistics reports.
extern const char* counterNames[NUM_CONV_TYPES];
extern const char* apiNames[NUM_API_TYPES];
extern const char *counterNames[NUM_CONV_TYPES];
extern const char *apiNames[NUM_API_TYPES];
struct hipCounter {
@@ -71,14 +77,14 @@ struct hipCounter {
* Tracks a set of named counters, as well as counters for each of the type enums defined above.
*/
class StatCounter {
private:
private:
// Each thing we track is either "supported" or "unsupported"...
std::map<std::string, int> counters;
int apiCounters[NUM_API_TYPES] = {};
int convTypeCounters[NUM_CONV_TYPES] = {};
public:
public:
void incrementCounter(const hipCounter& counter, std::string name);
/**
@@ -109,15 +115,15 @@ class Statistics {
chr::steady_clock::time_point startTime;
chr::steady_clock::time_point completionTime;
public:
public:
Statistics(std::string name);
void incrementCounter(const hipCounter& counter, std::string name);
void incrementCounter(const hipCounter &counter, std::string name);
/**
* Add the counters from `other` onto the counters of this object.
*/
void add(const Statistics& other);
void add(const Statistics &other);
void lineTouched(int lineNumber);
void bytesChanged(int bytes);
@@ -129,18 +135,18 @@ class Statistics {
/////// Output functions ///////
public:
/**
public:
/**
* Pretty-print the statistics stored in this object.
*
* @param csv Pointer to an output stream for the CSV to write. If null, no CSV is written
* @param printOut Pointer to an output stream to print human-readable textual stats to. If
* null, no such stats are produced.
* @param printOut Pointer to an output stream to print human-readable textual stats to. If null, no
* such stats are produced.
*/
void print(std::ostream* csv, llvm::raw_ostream* printOut, bool skipHeader = false);
/// Print aggregated statistics for all registered counters.
static void printAggregate(std::ostream* csv, llvm::raw_ostream* printOut);
static void printAggregate(std::ostream *csv, llvm::raw_ostream* printOut);
/////// Static nonsense ///////
@@ -156,15 +162,15 @@ class Statistics {
static Statistics getAggregate();
/**
* Convenient global entry point for updating the "active" Statistics. Since we operate
* single-threadedly processing one file at a time, this allows us to simply expose the stats
* for the current file globally, simplifying things.
* Convenient global entry point for updating the "active" Statistics. Since we operate single-threadedly
* processing one file at a time, this allows us to simply expose the stats for the current file globally,
* simplifying things.
*/
static Statistics& current();
/**
* Set the active Statistics object to the named one, creating it if necessary, and write the
* completion timestamp into the currently active one.
* Set the active Statistics object to the named one, creating it if necessary, and write the completion
* timestamp into the currently active one.
*/
static void setActive(std::string name);
};
+1 -1
Просмотреть файл
@@ -8,7 +8,7 @@ llvm::StringRef unquoteStr(llvm::StringRef s) {
return s;
}
void removePrefixIfPresent(std::string& s, std::string prefix) {
void removePrefixIfPresent(std::string &s, std::string prefix) {
if (s.find(prefix) != 0) {
return;
}
+1 -1
Просмотреть файл
@@ -11,4 +11,4 @@ llvm::StringRef unquoteStr(llvm::StringRef s);
/**
* If `s` starts with `prefix`, remove it. Otherwise, does nothing.
*/
void removePrefixIfPresent(std::string& s, std::string prefix);
void removePrefixIfPresent(std::string &s, std::string prefix);
+88 -93
Просмотреть файл
@@ -46,116 +46,111 @@ 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();
std::ifstream source(src, std::ios::binary);
std::ofstream dest(dst, std::ios::binary);
dest << source.rdbuf();
}
} // anonymous namespace
} // anonymous namespace
int main(int argc, const char** argv) {
llcompat::PrintStackTraceOnErrorSignal();
int main(int argc, const char **argv) {
llcompat::PrintStackTraceOnErrorSignal();
ct::CommonOptionsParser OptionsParser(argc, argv, ToolTemplateCategory, llvm::cl::OneOrMore);
std::vector<std::string> 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;
ct::CommonOptionsParser OptionsParser(argc, argv, ToolTemplateCategory, llvm::cl::OneOrMore);
std::vector<std::string> 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<std::ostream> csv = nullptr;
llvm::raw_ostream* statPrint = nullptr;
if (!OutputStatsFilename.empty()) {
csv = std::unique_ptr<std::ostream>(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;
}
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;
}
}
std::string tmpFile = src + ".hipify-tmp";
if (Examine) {
NoOutput = PrintStats = true;
}
// 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);
int Result = 0;
// Initialise the statistics counters for this file.
Statistics::setActive(src);
// Arguments for the Statistics print routines.
std::unique_ptr<std::ostream> csv = nullptr;
llvm::raw_ostream* statPrint = nullptr;
if (!OutputStatsFilename.empty()) {
csv = std::unique_ptr<std::ostream>(
new std::ofstream(OutputStatsFilename, std::ios_base::trunc));
}
if (PrintStats) {
statPrint = &llvm::errs();
}
// 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);
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;
}
ReplacementsFrontendActionFactory<HipifyAction> actionFactory(&replacementsToUse);
std::string tmpFile = src + ".hipify-tmp";
Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("--cuda-host-only", ct::ArgumentInsertPosition::BEGIN));
// 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<HipifyAction> 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));
// 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));
Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("-resource-dir=" HIPIFY_CLANG_RES));
#endif
Tool.appendArgumentsAdjuster(ct::getClangSyntaxOnlyAdjuster());
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();
// Hipify _all_ the things!
if (Tool.runAndSave(&actionFactory)) {
DEBUG(llvm::dbgs() << "Skipped some replacements.\n");
}
if (fileSources.size() > 1) {
Statistics::printAggregate(csv.get(), statPrint);
// Either move the tmpfile to the output, or remove it.
if (!NoOutput) {
rename(tmpFile.c_str(), dst.c_str());
} else {
remove(tmpFile.c_str());
}
return Result;
Statistics::current().markCompletion();
Statistics::current().print(csv.get(), statPrint);
dst.clear();
}
if (fileSources.size() > 1) {
Statistics::printAggregate(csv.get(), statPrint);
}
return Result;
}