Merge pull request #1642 from emankov/hipify

[HIPIFY] Clang style formatting

[ROCm/clr commit: 806ae1f43e]
This commit is contained in:
Evgeny Mankov
2019-11-07 12:49:52 +03:00
committed by GitHub
melakukan cdff675646
8 mengubah file dengan 51 tambahan dan 51 penghapusan
@@ -117,7 +117,7 @@ namespace perl {
{sCudaInGauge}, {sCudaColorSpinorField}, {sCudaSiteLink}, {sCudaFatLink}, {sCudaStaple}, {sCudaCloverField}, {sCudaParam} {sCudaInGauge}, {sCudaColorSpinorField}, {sCudaSiteLink}, {sCudaFatLink}, {sCudaStaple}, {sCudaCloverField}, {sCudaParam}
}; };
void generateHeader(unique_ptr<ostream>& streamPtr) { void generateHeader(unique_ptr<ostream> &streamPtr) {
*streamPtr.get() << "#!/usr/bin/perl -w" << endl_2; *streamPtr.get() << "#!/usr/bin/perl -w" << endl_2;
*streamPtr.get() << sCopyright << endl; *streamPtr.get() << sCopyright << endl;
*streamPtr.get() << sImportant << endl_2; *streamPtr.get() << sImportant << endl_2;
@@ -148,7 +148,7 @@ namespace perl {
*streamPtr.get() << "push(@whitelist, split(',', $whitelist));" << endl_2; *streamPtr.get() << "push(@whitelist, split(',', $whitelist));" << endl_2;
} }
void generateStatFunctions(unique_ptr<ostream>& streamPtr) { void generateStatFunctions(unique_ptr<ostream> &streamPtr) {
*streamPtr.get() << endl << sub << "totalStats" << " {" << endl; *streamPtr.get() << endl << sub << "totalStats" << " {" << endl;
*streamPtr.get() << tab << my << "%count = %{ shift() };" << endl; *streamPtr.get() << tab << my << "%count = %{ shift() };" << endl;
*streamPtr.get() << tab << my << "$total = 0;" << endl; *streamPtr.get() << tab << my << "$total = 0;" << endl;
@@ -175,11 +175,11 @@ namespace perl {
} }
} }
void generateSimpleSubstitutions(unique_ptr<ostream>& streamPtr) { void generateSimpleSubstitutions(unique_ptr<ostream> &streamPtr) {
*streamPtr.get() << endl << sub << "simpleSubstitutions" << " {" << endl; *streamPtr.get() << endl << sub << "simpleSubstitutions" << " {" << endl;
for (int i = 0; i < NUM_CONV_TYPES; ++i) { for (int i = 0; i < NUM_CONV_TYPES; ++i) {
if (i == CONV_INCLUDE_CUDA_MAIN_H || i == CONV_INCLUDE) { if (i == CONV_INCLUDE_CUDA_MAIN_H || i == CONV_INCLUDE) {
for (auto& ma : CUDA_INCLUDE_MAP) { for (auto &ma : CUDA_INCLUDE_MAP) {
if (Statistics::isUnsupported(ma.second)) continue; if (Statistics::isUnsupported(ma.second)) continue;
if (i == ma.second.type) { if (i == ma.second.type) {
string sCUDA = ma.first.str(); string sCUDA = ma.first.str();
@@ -190,7 +190,7 @@ namespace perl {
} }
} }
} else { } else {
for (auto& ma : CUDA_RENAMES_MAP()) { for (auto &ma : CUDA_RENAMES_MAP()) {
if (Statistics::isUnsupported(ma.second)) continue; if (Statistics::isUnsupported(ma.second)) continue;
if (i == ma.second.type) { if (i == ma.second.type) {
*streamPtr.get() << tab << "$ft{'" << counterNames[ma.second.type] << "'} += s/\\b" << ma.first.str() << "\\b/" << ma.second.hipName.str() << "/g;" << endl; *streamPtr.get() << tab << "$ft{'" << counterNames[ma.second.type] << "'} += s/\\b" << ma.first.str() << "\\b/" << ma.second.hipName.str() << "/g;" << endl;
@@ -201,7 +201,7 @@ namespace perl {
*streamPtr.get() << "}" << endl; *streamPtr.get() << "}" << endl;
} }
void generateExternShared(unique_ptr<ostream>& streamPtr) { void generateExternShared(unique_ptr<ostream> &streamPtr) {
*streamPtr.get() << endl << "# CUDA extern __shared__ syntax replace with HIP_DYNAMIC_SHARED() macro" << endl; *streamPtr.get() << endl << "# CUDA extern __shared__ syntax replace with HIP_DYNAMIC_SHARED() macro" << endl;
*streamPtr.get() << sub << "transformExternShared" << " {" << endl; *streamPtr.get() << sub << "transformExternShared" << " {" << endl;
*streamPtr.get() << tab << no_warns << endl; *streamPtr.get() << tab << no_warns << endl;
@@ -210,7 +210,7 @@ namespace perl {
*streamPtr.get() << tab << "$ft{'extern_shared'} += $k;" << endl << "}" << endl; *streamPtr.get() << tab << "$ft{'extern_shared'} += $k;" << endl << "}" << endl;
} }
void generateKernelLaunch(unique_ptr<ostream>& streamPtr) { void generateKernelLaunch(unique_ptr<ostream> &streamPtr) {
*streamPtr.get() << endl << "# CUDA Kernel Launch Syntax" << endl << sub << "transformKernelLaunch" << " {" << endl; *streamPtr.get() << endl << "# CUDA Kernel Launch Syntax" << endl << sub << "transformKernelLaunch" << " {" << endl;
*streamPtr.get() << tab << no_warns << endl; *streamPtr.get() << tab << no_warns << endl;
*streamPtr.get() << tab << my_k << endl_2; *streamPtr.get() << tab << my_k << endl_2;
@@ -251,13 +251,13 @@ namespace perl {
*streamPtr.get() << tab_2 << "$Tkernels{$1}++;" << endl_tab << "}" << endl << "}" << endl; *streamPtr.get() << tab_2 << "$Tkernels{$1}++;" << endl_tab << "}" << endl << "}" << endl;
} }
void generateCubNamespace(unique_ptr<ostream>& streamPtr) { void generateCubNamespace(unique_ptr<ostream> &streamPtr) {
*streamPtr.get() << endl << sub << "transformCubNamespace" << " {" << endl_tab << my_k << endl; *streamPtr.get() << endl << sub << "transformCubNamespace" << " {" << endl_tab << my_k << endl;
*streamPtr.get() << tab << "$k += s/using\\s*namespace\\s*cub/using namespace hipcub/g;" << endl; *streamPtr.get() << tab << "$k += s/using\\s*namespace\\s*cub/using namespace hipcub/g;" << endl;
*streamPtr.get() << tab << "$k += s/\\bcub::\\b/hipcub::/g;" << endl << tab << return_k << "}" << endl; *streamPtr.get() << tab << "$k += s/\\bcub::\\b/hipcub::/g;" << endl << tab << return_k << "}" << endl;
} }
void generateHostFunctions(unique_ptr<ostream>& streamPtr) { void generateHostFunctions(unique_ptr<ostream> &streamPtr) {
*streamPtr.get() << endl << sub << "transformHostFunctions" << " {" << endl_tab << my_k << endl; *streamPtr.get() << endl << sub << "transformHostFunctions" << " {" << endl_tab << my_k << endl;
set<string> &funcSet = DeviceSymbolFunctions0; set<string> &funcSet = DeviceSymbolFunctions0;
const string s0 = "$k += s/(?<!\\/\\/ CHECK: )($func)\\s*\\(\\s*([^,]+)\\s*,/$func\\("; const string s0 = "$k += s/(?<!\\/\\/ CHECK: )($func)\\s*\\(\\s*([^,]+)\\s*,/$func\\(";
@@ -271,7 +271,7 @@ namespace perl {
default: funcSet = DeviceSymbolFunctions0; default: funcSet = DeviceSymbolFunctions0;
} }
unsigned int count = 0; unsigned int count = 0;
for (auto& f : funcSet) { for (auto &f : funcSet) {
const auto found = CUDA_RUNTIME_FUNCTION_MAP.find(f); const auto found = CUDA_RUNTIME_FUNCTION_MAP.find(f);
if (found != CUDA_RUNTIME_FUNCTION_MAP.end()) { if (found != CUDA_RUNTIME_FUNCTION_MAP.end()) {
*streamPtr.get() << (count ? ",\n" : "") << tab_2 << "\"" << found->second.hipName.str() << "\""; *streamPtr.get() << (count ? ",\n" : "") << tab_2 << "\"" << found->second.hipName.str() << "\"";
@@ -291,12 +291,12 @@ namespace perl {
*streamPtr.get() << tab << return_k << "}" << endl; *streamPtr.get() << tab << return_k << "}" << endl;
} }
void generateDeviceFunctions(unique_ptr<ostream>& streamPtr) { void generateDeviceFunctions(unique_ptr<ostream> &streamPtr) {
unsigned int countUnsupported = 0; unsigned int countUnsupported = 0;
unsigned int countSupported = 0; unsigned int countSupported = 0;
stringstream sSupported; stringstream sSupported;
stringstream sUnsupported; stringstream sUnsupported;
for (auto& ma : CUDA_DEVICE_FUNC_MAP) { for (auto &ma : CUDA_DEVICE_FUNC_MAP) {
bool isUnsupported = Statistics::isUnsupported(ma.second); bool isUnsupported = Statistics::isUnsupported(ma.second);
(isUnsupported ? sUnsupported : sSupported) << ((isUnsupported && countUnsupported) || (!isUnsupported && countSupported) ? ",\n" : "") << tab_2 << "\"" << ma.first.str() << "\""; (isUnsupported ? sUnsupported : sSupported) << ((isUnsupported && countUnsupported) || (!isUnsupported && countSupported) ? ",\n" : "") << tab_2 << "\"" << ma.first.str() << "\"";
if (isUnsupported) countUnsupported++; if (isUnsupported) countUnsupported++;
@@ -60,9 +60,9 @@ namespace python {
*pythonStreamPtr.get() << "from pyHIPIFY.constants import *\n\n"; *pythonStreamPtr.get() << "from pyHIPIFY.constants import *\n\n";
*pythonStreamPtr.get() << "CUDA_RENAMES_MAP = collections.OrderedDict([\n"; *pythonStreamPtr.get() << "CUDA_RENAMES_MAP = collections.OrderedDict([\n";
const std::string sHIP_UNS = ", HIP_UNSUPPORTED"; const std::string sHIP_UNS = ", HIP_UNSUPPORTED";
for (int i = 0; i < NUM_CONV_TYPES; i++) { for (int i = 0; i < NUM_CONV_TYPES; ++i) {
if (i == CONV_INCLUDE_CUDA_MAIN_H || i == CONV_INCLUDE) { if (i == CONV_INCLUDE_CUDA_MAIN_H || i == CONV_INCLUDE) {
for (auto& ma : CUDA_INCLUDE_MAP) { for (auto &ma : CUDA_INCLUDE_MAP) {
if (i == ma.second.type) { if (i == ma.second.type) {
std::string sUnsupported; std::string sUnsupported;
if (Statistics::isUnsupported(ma.second)) { if (Statistics::isUnsupported(ma.second)) {
@@ -74,7 +74,7 @@ namespace python {
} }
} }
else { else {
for (auto& ma : CUDA_RENAMES_MAP()) { for (auto &ma : CUDA_RENAMES_MAP()) {
if (i == ma.second.type) { if (i == ma.second.type) {
std::string sUnsupported; std::string sUnsupported;
if (Statistics::isUnsupported(ma.second)) { if (Statistics::isUnsupported(ma.second)) {
@@ -150,7 +150,7 @@ void HipifyAction::RewriteToken(const clang::Token &t) {
void HipifyAction::FindAndReplace(StringRef name, void HipifyAction::FindAndReplace(StringRef name,
clang::SourceLocation sl, clang::SourceLocation sl,
const std::map<StringRef, hipCounter> &repMap, const std::map<StringRef, hipCounter> &repMap,
bool bReplace) { bool bReplace) {
const auto found = repMap.find(name); const auto found = repMap.find(name);
if (found == repMap.end()) { if (found == repMap.end()) {
@@ -36,15 +36,15 @@ namespace ct = clang::tooling;
*/ */
template <typename T> template <typename T>
class ReplacementsFrontendActionFactory : public ct::FrontendActionFactory { class ReplacementsFrontendActionFactory : public ct::FrontendActionFactory {
ct::Replacements* replacements; ct::Replacements *replacements;
public: public:
explicit ReplacementsFrontendActionFactory(ct::Replacements* r): explicit ReplacementsFrontendActionFactory(ct::Replacements *r):
ct::FrontendActionFactory(), ct::FrontendActionFactory(),
replacements(r) {} replacements(r) {}
#if LLVM_VERSION_MAJOR < 10 #if LLVM_VERSION_MAJOR < 10
clang::FrontendAction* create() override { clang::FrontendAction *create() override {
return new T(replacements); return new T(replacements);
} }
#else #else
@@ -163,51 +163,51 @@ void printStat(std::ostream *csv, llvm::raw_ostream* printOut, const std::string
} // Anonymous namespace } // Anonymous namespace
void StatCounter::incrementCounter(const hipCounter& counter, const std::string& name) { void StatCounter::incrementCounter(const hipCounter &counter, const std::string &name) {
counters[name]++; counters[name]++;
apiCounters[(int) counter.apiType]++; apiCounters[(int) counter.apiType]++;
convTypeCounters[(int) counter.type]++; convTypeCounters[(int) counter.type]++;
} }
void StatCounter::add(const StatCounter& other) { void StatCounter::add(const StatCounter &other) {
for (const auto& p : other.counters) { for (const auto &p : other.counters) {
counters[p.first] += p.second; counters[p.first] += p.second;
} }
for (int i = 0; i < NUM_API_TYPES; i++) { for (int i = 0; i < NUM_API_TYPES; ++i) {
apiCounters[i] += other.apiCounters[i]; apiCounters[i] += other.apiCounters[i];
} }
for (int i = 0; i < NUM_CONV_TYPES; i++) { for (int i = 0; i < NUM_CONV_TYPES; ++i) {
convTypeCounters[i] += other.convTypeCounters[i]; convTypeCounters[i] += other.convTypeCounters[i];
} }
} }
int StatCounter::getConvSum() { int StatCounter::getConvSum() {
int acc = 0; int acc = 0;
for (const int& i : convTypeCounters) { for (const int &i : convTypeCounters) {
acc += i; acc += i;
} }
return acc; return acc;
} }
void StatCounter::print(std::ostream* csv, llvm::raw_ostream* printOut, const std::string& prefix) { void StatCounter::print(std::ostream* csv, llvm::raw_ostream* printOut, const std::string &prefix) {
for (int i = 0; i < NUM_CONV_TYPES; i++) { for (int i = 0; i < NUM_CONV_TYPES; ++i) {
if (convTypeCounters[i] > 0) { if (convTypeCounters[i] > 0) {
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");
break; break;
} }
} }
for (int i = 0; i < NUM_CONV_TYPES; i++) { for (int i = 0; i < NUM_CONV_TYPES; ++i) {
if (convTypeCounters[i] > 0) { if (convTypeCounters[i] > 0) {
printStat(csv, printOut, counterNames[i], convTypeCounters[i]); printStat(csv, printOut, counterNames[i], convTypeCounters[i]);
} }
} }
for (int i = 0; i < NUM_API_TYPES; i++) { for (int i = 0; i < NUM_API_TYPES; ++i) {
if (apiCounters[i] > 0) { if (apiCounters[i] > 0) {
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");
break; break;
} }
} }
for (int i = 0; i < NUM_API_TYPES; i++) { for (int i = 0; i < NUM_API_TYPES; ++i) {
if (apiCounters[i] > 0) { if (apiCounters[i] > 0) {
printStat(csv, printOut, apiNames[i], apiCounters[i]); printStat(csv, printOut, apiNames[i], apiCounters[i]);
} }
@@ -220,7 +220,7 @@ void StatCounter::print(std::ostream* csv, llvm::raw_ostream* printOut, const st
} }
} }
Statistics::Statistics(const std::string& name): fileName(name) { Statistics::Statistics(const std::string &name): fileName(name) {
// Compute the total bytes/lines in the input file. // Compute the total bytes/lines in the input file.
std::ifstream src_file(name, std::ios::binary | std::ios::ate); std::ifstream src_file(name, std::ios::binary | std::ios::ate);
src_file.clear(); src_file.clear();
@@ -235,7 +235,7 @@ Statistics::Statistics(const std::string& name): fileName(name) {
///////// Counter update routines ////////// ///////// Counter update routines //////////
void Statistics::incrementCounter(const hipCounter &counter, const std::string& name) { void Statistics::incrementCounter(const hipCounter &counter, const std::string &name) {
if (Statistics::isUnsupported(counter)) { if (Statistics::isUnsupported(counter)) {
unsupported.incrementCounter(counter, name); unsupported.incrementCounter(counter, name);
} else { } else {
@@ -308,7 +308,7 @@ void Statistics::printAggregate(std::ostream *csv, llvm::raw_ostream* printOut)
Statistics globalStats = getAggregate(); Statistics globalStats = getAggregate();
// A file is considered "converted" if we made any changes to it. // A file is considered "converted" if we made any changes to it.
int convertedFiles = 0; int convertedFiles = 0;
for (const auto& p : stats) { for (const auto &p : stats) {
if (p.second.touchedLines && p.second.totalBytes && if (p.second.touchedLines && p.second.totalBytes &&
p.second.totalLines && !p.second.hasErrors) { p.second.totalLines && !p.second.hasErrors) {
convertedFiles++; convertedFiles++;
@@ -326,18 +326,18 @@ void Statistics::printAggregate(std::ostream *csv, llvm::raw_ostream* printOut)
Statistics Statistics::getAggregate() { Statistics Statistics::getAggregate() {
Statistics globalStats("GLOBAL"); Statistics globalStats("GLOBAL");
for (const auto& p : stats) { for (const auto &p : stats) {
globalStats.add(p.second); globalStats.add(p.second);
} }
return globalStats; return globalStats;
} }
Statistics& Statistics::current() { Statistics &Statistics::current() {
assert(Statistics::currentStatistics); assert(Statistics::currentStatistics);
return *Statistics::currentStatistics; return *Statistics::currentStatistics;
} }
void Statistics::setActive(const std::string& name) { void Statistics::setActive(const std::string &name) {
stats.emplace(std::make_pair(name, Statistics{name})); stats.emplace(std::make_pair(name, Statistics{name}));
Statistics::currentStatistics = &stats.at(name); Statistics::currentStatistics = &stats.at(name);
} }
@@ -172,11 +172,11 @@ private:
int convTypeCounters[NUM_CONV_TYPES] = {}; int convTypeCounters[NUM_CONV_TYPES] = {};
public: public:
void incrementCounter(const hipCounter& counter, const std::string& name); void incrementCounter(const hipCounter &counter, const std::string &name);
// Add the counters from `other` onto the counters of this object. // Add the counters from `other` onto the counters of this object.
void add(const StatCounter& other); void add(const StatCounter &other);
int getConvSum(); int getConvSum();
void print(std::ostream* csv, llvm::raw_ostream* printOut, const std::string& prefix); void print(std::ostream* csv, llvm::raw_ostream* printOut, const std::string &prefix);
}; };
/** /**
@@ -195,8 +195,8 @@ class Statistics {
chr::steady_clock::time_point completionTime; chr::steady_clock::time_point completionTime;
public: public:
Statistics(const std::string& name); Statistics(const std::string &name);
void incrementCounter(const hipCounter &counter, const std::string& name); void incrementCounter(const hipCounter &counter, const std::string &name);
// Add the counters from `other` onto the counters of this object. // 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 lineTouched(int lineNumber);
@@ -226,12 +226,12 @@ public:
* processing one file at a time, this allows us to simply expose the stats for the current file globally, * processing one file at a time, this allows us to simply expose the stats for the current file globally,
* simplifying things. * simplifying things.
*/ */
static Statistics& current(); static Statistics &current();
/** /**
* Set the active Statistics object to the named one, creating it if necessary, and write the completion * Set the active Statistics object to the named one, creating it if necessary, and write the completion
* timestamp into the currently active one. * timestamp into the currently active one.
*/ */
static void setActive(const std::string& name); static void setActive(const std::string &name);
// Check the counter and option TranslateToRoc whether it should be translated to Roc or not. // Check the counter and option TranslateToRoc whether it should be translated to Roc or not.
static bool isToRoc(const hipCounter &counter); static bool isToRoc(const hipCounter &counter);
// Check whether the counter is HIP_UNSUPPORTED or not. // Check whether the counter is HIP_UNSUPPORTED or not.
@@ -33,14 +33,14 @@ llvm::StringRef unquoteStr(llvm::StringRef s) {
return s; return s;
} }
void removePrefixIfPresent(std::string &s, const std::string& prefix) { void removePrefixIfPresent(std::string &s, const std::string &prefix) {
if (s.find(prefix) != 0) { if (s.find(prefix) != 0) {
return; return;
} }
s.erase(0, prefix.size()); s.erase(0, prefix.size());
} }
std::string getAbsoluteFilePath(const std::string& sFile, std::error_code& EC) { std::string getAbsoluteFilePath(const std::string &sFile, std::error_code &EC) {
if (sFile.empty()) { if (sFile.empty()) {
return sFile; return sFile;
} }
@@ -59,8 +59,8 @@ std::string getAbsoluteFilePath(const std::string& sFile, std::error_code& EC) {
return fileAbsPath.c_str(); return fileAbsPath.c_str();
} }
std::string getAbsoluteDirectoryPath(const std::string& sDir, std::error_code& EC, std::string getAbsoluteDirectoryPath(const std::string &sDir, std::error_code &EC,
const std::string& sDirType, bool bCreateDir) { const std::string &sDirType, bool bCreateDir) {
if (sDir.empty()) { if (sDir.empty()) {
return sDir; return sDir;
} }
@@ -33,16 +33,16 @@ llvm::StringRef unquoteStr(llvm::StringRef s);
/** /**
* If `s` starts with `prefix`, remove it. Otherwise, does nothing. * If `s` starts with `prefix`, remove it. Otherwise, does nothing.
*/ */
void removePrefixIfPresent(std::string &s, const std::string& prefix); void removePrefixIfPresent(std::string &s, const std::string &prefix);
/** /**
* Returns Absolute File Path based on filename, otherwise - error. * Returns Absolute File Path based on filename, otherwise - error.
*/ */
std::string getAbsoluteFilePath(const std::string& sFile, std::error_code& EC); std::string getAbsoluteFilePath(const std::string &sFile, std::error_code &EC);
/** /**
* Returns Absolute Directory Path based on directory name, otherwise - error; * Returns Absolute Directory Path based on directory name, otherwise - error;
* by default the directory is temporary and created. * by default the directory is temporary and created.
*/ */
std::string getAbsoluteDirectoryPath(const std::string& sDir, std::error_code& EC, std::string getAbsoluteDirectoryPath(const std::string &sDir, std::error_code &EC,
const std::string& sDirType = "temporary", bool bCreateDir = true); const std::string &sDirType = "temporary", bool bCreateDir = true);