From 6160cff6ff81ed57e85a75ae16bed4957bc2cd15 Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Thu, 7 Feb 2019 13:17:05 +0300 Subject: [PATCH] [HIPIFY][BLAS] Restrict hipification to 'ROC' to BLAS library only + Add utility functions isToRoc, isHipUnsupported, isRocUnsupported, isUnsupported --- hipamd/hipify-clang/src/CUDA2HIP.cpp | 4 +-- hipamd/hipify-clang/src/HipifyAction.cpp | 32 ++++++++---------------- hipamd/hipify-clang/src/Statistics.cpp | 25 +++++++++++++++--- hipamd/hipify-clang/src/Statistics.h | 11 ++++++++ 4 files changed, 45 insertions(+), 27 deletions(-) diff --git a/hipamd/hipify-clang/src/CUDA2HIP.cpp b/hipamd/hipify-clang/src/CUDA2HIP.cpp index 0dcef4da35..920b310da8 100644 --- a/hipamd/hipify-clang/src/CUDA2HIP.cpp +++ b/hipamd/hipify-clang/src/CUDA2HIP.cpp @@ -37,8 +37,8 @@ const std::map CUDA_INCLUDE_MAP{ // cuComplex includes {"cuComplex.h", {"hip/hip_complex.h", "", CONV_INCLUDE_CUDA_MAIN_H, API_COMPLEX}}, // cuBLAS includes - {"cublas.h", {"hipblas.h", "", CONV_INCLUDE_CUDA_MAIN_H, API_BLAS}}, - {"cublas_v2.h", {"hipblas.h", "", CONV_INCLUDE_CUDA_MAIN_H, API_BLAS}}, + {"cublas.h", {"hipblas.h", "rocblas.h", CONV_INCLUDE_CUDA_MAIN_H, API_BLAS}}, + {"cublas_v2.h", {"hipblas.h", "rocblas.h", CONV_INCLUDE_CUDA_MAIN_H, API_BLAS}}, // cuRAND includes {"curand.h", {"hiprand.h", "", CONV_INCLUDE_CUDA_MAIN_H, API_RAND}}, {"curand_kernel.h", {"hiprand_kernel.h", "", CONV_INCLUDE, API_RAND}}, diff --git a/hipamd/hipify-clang/src/HipifyAction.cpp b/hipamd/hipify-clang/src/HipifyAction.cpp index a021fb3f74..f4d68cd820 100644 --- a/hipamd/hipify-clang/src/HipifyAction.cpp +++ b/hipamd/hipify-clang/src/HipifyAction.cpp @@ -41,11 +41,10 @@ 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 = TranslateToRoc ? found->second.rocName : found->second.hipName; + StringRef repName = Statistics::isToRoc(found->second) ? 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 ((!TranslateToRoc && (HIP_UNSUPPORTED != (counter.supportDegree & HIP_UNSUPPORTED))) || - (TranslateToRoc && (ROC_UNSUPPORTED != (counter.supportDegree & ROC_UNSUPPORTED)))) { + if (!Statistics::isUnsupported(counter)) { clang::SourceLocation sl = start.getLocWithOffset(begin + 1); ct::Replacement Rep(SM, sl, name.size(), repName.str()); clang::FullSourceLoc fullSL(sl, SM); @@ -60,7 +59,7 @@ void HipifyAction::RewriteString(StringRef s, clang::SourceLocation start) { } /** - * Look at, and consider altering, a given token. + * 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. @@ -68,7 +67,7 @@ void HipifyAction::RewriteString(StringRef s, clang::SourceLocation start) { */ void HipifyAction::RewriteToken(const clang::Token& t) { clang::SourceManager& SM = getCompilerInstance().getSourceManager(); - // String literals containing CUDA references need fixing... + // String literals containing CUDA references need fixing. if (t.is(clang::tok::string_literal)) { StringRef s(t.getLiteralData(), t.getLength()); RewriteString(unquoteStr(s), t.getLocation()); @@ -86,26 +85,16 @@ void HipifyAction::RewriteToken(const clang::Token& t) { Statistics::current().incrementCounter(found->second, name.str()); clang::SourceLocation sl = t.getLocation(); 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) { + if (Statistics::isUnsupported(found->second)) { + std::string sWarn; + Statistics::isToRoc(found->second) ? sWarn = "ROC" : sWarn = "HIP"; sWarn = "" + sWarn; const auto ID = DE.getCustomDiagID(clang::DiagnosticsEngine::Warning, "CUDA identifier is unsupported in %0."); DE.Report(sl, ID) << sWarn; return; } - StringRef repName = TranslateToRoc ? found->second.rocName : found->second.hipName; + StringRef repName = Statistics::isToRoc(found->second) ? found->second.rocName : found->second.hipName; ct::Replacement Rep(SM, sl, name.size(), repName.str()); clang::FullSourceLoc fullSL(sl, SM); insertReplacement(Rep, fullSL); @@ -238,8 +227,7 @@ void HipifyAction::InclusionDirective(clang::SourceLocation hash_loc, Statistics::current().incrementCounter(found->second, file_name.str()); clang::SourceLocation sl = filename_range.getBegin(); - if ((!TranslateToRoc && (HIP_UNSUPPORTED == (found->second.supportDegree & HIP_UNSUPPORTED))) || - (TranslateToRoc && (ROC_UNSUPPORTED == (found->second.supportDegree & ROC_UNSUPPORTED)))) { + if (Statistics::isUnsupported(found->second)) { clang::DiagnosticsEngine& DE = getCompilerInstance().getDiagnostics(); DE.Report(sl, DE.getCustomDiagID(clang::DiagnosticsEngine::Warning, "Unsupported CUDA header.")); return; @@ -249,7 +237,7 @@ 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; + llvm::StringRef name = Statistics::isToRoc(found->second) ? found->second.rocName : found->second.hipName; if (is_angled) { newInclude = llvm::Twine("<" + name+ ">").toStringRef(includeBuffer); } else { diff --git a/hipamd/hipify-clang/src/Statistics.cpp b/hipamd/hipify-clang/src/Statistics.cpp index 1943d82f98..39f70e9d8a 100644 --- a/hipamd/hipify-clang/src/Statistics.cpp +++ b/hipamd/hipify-clang/src/Statistics.cpp @@ -177,8 +177,7 @@ Statistics::Statistics(const std::string& name): fileName(name) { ///////// Counter update routines ////////// void Statistics::incrementCounter(const hipCounter &counter, const std::string& name) { - if ((!TranslateToRoc && (HIP_UNSUPPORTED == (counter.supportDegree & HIP_UNSUPPORTED))) || - (TranslateToRoc && (ROC_UNSUPPORTED == (counter.supportDegree & ROC_UNSUPPORTED)))) { + if (Statistics::isUnsupported(counter)) { unsupported.incrementCounter(counter, name); } else { supported.incrementCounter(counter, name); @@ -202,7 +201,7 @@ void Statistics::add(const Statistics &other) { void Statistics::lineTouched(int lineNumber) { touchedLinesSet.insert(lineNumber); - touchedLines = touchedLinesSet.size(); + touchedLines = unsigned(touchedLinesSet.size()); } void Statistics::bytesChanged(int bytes) { @@ -284,5 +283,25 @@ void Statistics::setActive(const std::string& name) { Statistics::currentStatistics = &stats.at(name); } +bool Statistics::isToRoc(const hipCounter &counter) { + return TranslateToRoc && counter.apiType == API_BLAS; +} + +bool Statistics::isHipUnsupported(const hipCounter &counter) { + return HIP_UNSUPPORTED == (counter.supportDegree & HIP_UNSUPPORTED); +} + +bool Statistics::isRocUnsupported(const hipCounter &counter) { + return ROC_UNSUPPORTED == (counter.supportDegree & ROC_UNSUPPORTED); +} + +bool Statistics::isUnsupported(const hipCounter &counter) { + if (Statistics::isToRoc(counter)) { + return Statistics::isRocUnsupported(counter); + } else { + return Statistics::isHipUnsupported(counter); + } +} + std::map Statistics::stats = {}; Statistics* Statistics::currentStatistics = nullptr; diff --git a/hipamd/hipify-clang/src/Statistics.h b/hipamd/hipify-clang/src/Statistics.h index 4d51b71697..6d8986bc52 100644 --- a/hipamd/hipify-clang/src/Statistics.h +++ b/hipamd/hipify-clang/src/Statistics.h @@ -225,6 +225,17 @@ public: * timestamp into the currently active one. */ static void setActive(const std::string& name); + // Check the counter and option TranslateToRoc whether it should be translated to Roc or not. + static bool isToRoc(const hipCounter &counter); + // Check whether the counter is HIP_UNSUPPORTED or not. + static bool isHipUnsupported(const hipCounter &counter); + // Check whether the counter is ROC_UNSUPPORTED or not. + static bool isRocUnsupported(const hipCounter &counter); + /** + * Check whether the counter is ROC_UNSUPPORTED/HIP_UNSUPPORTED/UNSUPPORTED or not + * based on counter's API_TYPE and option TranslateToRoc. + */ + static bool isUnsupported(const hipCounter &counter); // Set this flag in case of hipification errors bool hasErrors = false; };