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

[HIPIFY] Initial support of translation to ROCm

Browse Source 
+ Option '-roc' is added to translate to ROCm instead of HIP
+ Update hipification mechanism accordingly
+ Populate all translation maps with additional field for ROCm identifiers
+ There are 4 degrees of CUDA support now: full support, only by HIP, only by ROCm, unsupported


[ROCm/hip commit: c5ed25da22]
This commit is contained in:
Evgeny Mankov
2018-12-06 21:52:54 +03:00
parent 42f97d2e46
commit 3621afd4e6
23 changed files with 3735 additions and 3705 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/hip/hipify-clang/src/HipifyAction.cpp
+31 -15
View File
@@ -19,10 +19,11 @@ void HipifyAction::RewriteString(StringRef s, clang::SourceLocation start) {
StringRef name = s.slice(begin, end);
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};
StringRef repName = TranslateToRoc ? found->second.rocName : found->second.hipName;
hipCounter counter = {"[string literal]", "", ConvTypes::CONV_LITERAL, ApiTypes::API_RUNTIME, found->second.supportDegree};
Statistics::current().incrementCounter(counter, name.str());
if (!counter.unsupported) {
if ((!TranslateToRoc && (HIP_UNSUPPORTED != (counter.supportDegree & HIP_UNSUPPORTED))) ||
(TranslateToRoc && (ROC_UNSUPPORTED != (counter.supportDegree & ROC_UNSUPPORTED)))) {
clang::SourceLocation sl = start.getLocWithOffset(begin + 1);
ct::Replacement Rep(SM, sl, name.size(), repName.str());
clang::FullSourceLoc fullSL(sl, SM);
@@ -62,14 +63,27 @@ void HipifyAction::RewriteToken(const clang::Token& t) {
}
Statistics::current().incrementCounter(found->second, name.str());
clang::SourceLocation sl = t.getLocation();
if (found->second.unsupported) {
// Warn the user about unsupported identifier.
clang::DiagnosticsEngine& DE = getCompilerInstance().getDiagnostics();
const auto ID = DE.getCustomDiagID(clang::DiagnosticsEngine::Warning, "CUDA identifier unsupported in hip");
DE.Report(sl, ID);
clang::DiagnosticsEngine& DE = getCompilerInstance().getDiagnostics();
bool bWarn = false;
std::string sWarn = "HIP";
if (TranslateToRoc) {
if ((found->second.supportDegree & ROC_UNSUPPORTED) == ROC_UNSUPPORTED) {
sWarn = "ROCm";
bWarn = true;
}
} else {
if ((found->second.supportDegree & HIP_UNSUPPORTED) == HIP_UNSUPPORTED) {
bWarn = true;
}
}
// Warn the user about unsupported identifier.
if (bWarn) {
sWarn = "" + sWarn;
const auto ID = DE.getCustomDiagID(clang::DiagnosticsEngine::Warning, "CUDA identifier is unsupported in %0.");
DE.Report(sl, ID) << sWarn;
return;
}
StringRef repName = found->second.hipName;
StringRef repName = TranslateToRoc ? found->second.rocName : found->second.hipName;
ct::Replacement Rep(SM, sl, name.size(), repName.str());
clang::FullSourceLoc fullSL(sl, SM);
insertReplacement(Rep, fullSL);
@@ -202,9 +216,10 @@ void HipifyAction::InclusionDirective(clang::SourceLocation hash_loc,
Statistics::current().incrementCounter(found->second, file_name.str());
clang::SourceLocation sl = filename_range.getBegin();
if (found->second.unsupported) {
if ((!TranslateToRoc && (HIP_UNSUPPORTED == (found->second.supportDegree & HIP_UNSUPPORTED))) ||
(TranslateToRoc && (ROC_UNSUPPORTED == (found->second.supportDegree & ROC_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;
}
@@ -212,10 +227,11 @@ void HipifyAction::InclusionDirective(clang::SourceLocation hash_loc,
// Keep the same include type that the user gave.
if (!exclude) {
clang::SmallString<128> includeBuffer;
llvm::StringRef name = TranslateToRoc ? found->second.rocName : found->second.hipName;
if (is_angled) {
newInclude = llvm::Twine("<" + found->second.hipName + ">").toStringRef(includeBuffer);
newInclude = llvm::Twine("<" + name+ ">").toStringRef(includeBuffer);
} else {
newInclude = llvm::Twine("\"" + found->second.hipName + "\"").toStringRef(includeBuffer);
newInclude = llvm::Twine("\"" + name + "\"").toStringRef(includeBuffer);
}
} else {
// hashLoc is location of the '#', thus replacing the whole include directive by empty newInclude starting with '#'.
@@ -288,7 +304,7 @@ bool HipifyAction::cudaLaunchKernel(const clang::ast_matchers::MatchFinder::Matc
ct::Replacement Rep(*SM, launchStart, length, OS.str());
clang::FullSourceLoc fullSL(launchStart, *SM);
insertReplacement(Rep, fullSL);
hipCounter counter = {"hipLaunchKernelGGL", ConvTypes::CONV_EXECUTION, ApiTypes::API_RUNTIME};
hipCounter counter = {"hipLaunchKernelGGL", "", ConvTypes::CONV_EXECUTION, ApiTypes::API_RUNTIME};
Statistics::current().incrementCounter(counter, refName.str());
return true;
}
@@ -333,7 +349,7 @@ bool HipifyAction::cudaSharedIncompleteArrayVar(const clang::ast_matchers::Match
ct::Replacement Rep(*SM, slStart, repLength, repName);
clang::FullSourceLoc fullSL(slStart, *SM);
insertReplacement(Rep, fullSL);
hipCounter counter = {"HIP_DYNAMIC_SHARED", ConvTypes::CONV_MEMORY, ApiTypes::API_RUNTIME};
hipCounter counter = {"HIP_DYNAMIC_SHARED", "", ConvTypes::CONV_MEMORY, ApiTypes::API_RUNTIME};
Statistics::current().incrementCounter(counter, refName.str());
}
return true;