Merge pull request #1536 from emankov/doc
[HIPIFY] HipifyAction code clean-up
This commit is contained in:
Evgeny Mankov
GitHub
@@ -33,16 +33,20 @@ THE SOFTWARE.
|
||||
#include "StringUtils.h"
|
||||
#include "ArgParse.h"
|
||||
|
||||
namespace ct = clang::tooling;
|
||||
namespace mat = clang::ast_matchers;
|
||||
|
||||
const std::string sHIP = "HIP";
|
||||
const std::string sROC = "ROC";
|
||||
const std::string sCub = "cub";
|
||||
const std::string sHIP_DYNAMIC_SHARED = "HIP_DYNAMIC_SHARED";
|
||||
const std::string sHIP_KERNEL_NAME = "HIP_KERNEL_NAME";
|
||||
std::string sHIP_SYMBOL = "HIP_SYMBOL";
|
||||
std::string sHIP_KERNEL_NAME = "HIP_KERNEL_NAME";
|
||||
std::string s_reinterpret_cast = "reinterpret_cast<const void*>";
|
||||
const std::string sHipLaunchKernelGGL = "hipLaunchKernelGGL(";
|
||||
const std::string sHipLaunchKernelGGL = "hipLaunchKernelGGL";
|
||||
const std::string sDim3 = "dim3(";
|
||||
|
||||
const std::string s_hiprand_kernel_h = "hiprand_kernel.h";
|
||||
const std::string s_hiprand_h = "hiprand.h";
|
||||
const std::string sOnce = "once";
|
||||
const std::string s_string_literal = "[string literal]";
|
||||
// CUDA identifiers, used in matchers
|
||||
const std::string sCudaMemcpyToSymbol = "cudaMemcpyToSymbol";
|
||||
const std::string sCudaMemcpyToSymbolAsync = "cudaMemcpyToSymbolAsync";
|
||||
const std::string sCudaGetSymbolSize = "cudaGetSymbolSize";
|
||||
@@ -51,6 +55,12 @@ const std::string sCudaMemcpyFromSymbol = "cudaMemcpyFromSymbol";
|
||||
const std::string sCudaMemcpyFromSymbolAsync = "cudaMemcpyFromSymbolAsync";
|
||||
const std::string sCudaFuncSetCacheConfig = "cudaFuncSetCacheConfig";
|
||||
const std::string sCudaFuncGetAttributes = "cudaFuncGetAttributes";
|
||||
// Matchers' names
|
||||
const StringRef sCudaSharedIncompleteArrayVar = "cudaSharedIncompleteArrayVar";
|
||||
const StringRef sCudaLaunchKernel = "cudaLaunchKernel";
|
||||
const StringRef sCudaHostFuncCall = "cudaHostFuncCall";
|
||||
const StringRef sCudaDeviceFuncCall = "cudaDeviceFuncCall";
|
||||
const StringRef sCubNamespacePrefix = "cubNamespacePrefix";
|
||||
|
||||
std::set<std::string> DeviceSymbolFunctions0 {
|
||||
{sCudaMemcpyToSymbol},
|
||||
@@ -78,7 +88,7 @@ std::set<std::string> ReinterpretFunctions1{
|
||||
};
|
||||
|
||||
void HipifyAction::RewriteString(StringRef s, clang::SourceLocation start) {
|
||||
clang::SourceManager& SM = getCompilerInstance().getSourceManager();
|
||||
auto &SM = getCompilerInstance().getSourceManager();
|
||||
size_t begin = 0;
|
||||
while ((begin = s.find("cu", begin)) != StringRef::npos) {
|
||||
const size_t end = s.find_first_of(" ", begin + 4);
|
||||
@@ -86,7 +96,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 = Statistics::isToRoc(found->second) ? found->second.rocName : found->second.hipName;
|
||||
hipCounter counter = {"[string literal]", "", ConvTypes::CONV_LITERAL, ApiTypes::API_RUNTIME, found->second.supportDegree};
|
||||
hipCounter counter = {s_string_literal, "", ConvTypes::CONV_LITERAL, ApiTypes::API_RUNTIME, found->second.supportDegree};
|
||||
Statistics::current().incrementCounter(counter, name.str());
|
||||
if (!Statistics::isUnsupported(counter)) {
|
||||
clang::SourceLocation sl = start.getLocWithOffset(begin + 1);
|
||||
@@ -95,9 +105,7 @@ void HipifyAction::RewriteString(StringRef s, clang::SourceLocation start) {
|
||||
insertReplacement(Rep, fullSL);
|
||||
}
|
||||
}
|
||||
if (end == StringRef::npos) {
|
||||
break;
|
||||
}
|
||||
if (end == StringRef::npos) break;
|
||||
begin = end + 1;
|
||||
}
|
||||
}
|
||||
@@ -109,7 +117,7 @@ void HipifyAction::RewriteString(StringRef s, clang::SourceLocation start) {
|
||||
* If it's an unsupported CUDA identifier, a warning is emitted.
|
||||
* Otherwise, the source file is updated with the corresponding hipification.
|
||||
*/
|
||||
void HipifyAction::RewriteToken(const clang::Token& t) {
|
||||
void HipifyAction::RewriteToken(const clang::Token &t) {
|
||||
// String literals containing CUDA references need fixing.
|
||||
if (t.is(clang::tok::string_literal)) {
|
||||
StringRef s(t.getLiteralData(), t.getLength());
|
||||
@@ -124,20 +132,21 @@ void HipifyAction::RewriteToken(const clang::Token& t) {
|
||||
FindAndReplace(name, sl, CUDA_RENAMES_MAP());
|
||||
}
|
||||
|
||||
void HipifyAction::FindAndReplace(llvm::StringRef name,
|
||||
void HipifyAction::FindAndReplace(StringRef name,
|
||||
clang::SourceLocation sl,
|
||||
const std::map<llvm::StringRef, hipCounter>& repMap, bool bReplace) {
|
||||
const std::map<StringRef, hipCounter> &repMap,
|
||||
bool bReplace) {
|
||||
const auto found = repMap.find(name);
|
||||
if (found == repMap.end()) {
|
||||
// So it's an identifier, but not CUDA? Boring.
|
||||
return;
|
||||
}
|
||||
Statistics::current().incrementCounter(found->second, name.str());
|
||||
clang::DiagnosticsEngine& DE = getCompilerInstance().getDiagnostics();
|
||||
clang::DiagnosticsEngine &DE = getCompilerInstance().getDiagnostics();
|
||||
// Warn the user about unsupported identifier.
|
||||
if (Statistics::isUnsupported(found->second)) {
|
||||
std::string sWarn;
|
||||
Statistics::isToRoc(found->second) ? sWarn = "ROC" : sWarn = "HIP";
|
||||
Statistics::isToRoc(found->second) ? sWarn = sROC : sWarn = sHIP;
|
||||
sWarn = "" + sWarn;
|
||||
const auto ID = DE.getCustomDiagID(clang::DiagnosticsEngine::Warning, "CUDA identifier is unsupported in %0.");
|
||||
DE.Report(sl, ID) << sWarn;
|
||||
@@ -147,7 +156,7 @@ void HipifyAction::FindAndReplace(llvm::StringRef name,
|
||||
return;
|
||||
}
|
||||
StringRef repName = Statistics::isToRoc(found->second) ? found->second.rocName : found->second.hipName;
|
||||
clang::SourceManager& SM = getCompilerInstance().getSourceManager();
|
||||
auto &SM = getCompilerInstance().getSourceManager();
|
||||
ct::Replacement Rep(SM, sl, name.size(), repName.str());
|
||||
clang::FullSourceLoc fullSL(sl, SM);
|
||||
insertReplacement(Rep, fullSL);
|
||||
@@ -155,7 +164,7 @@ void HipifyAction::FindAndReplace(llvm::StringRef name,
|
||||
|
||||
namespace {
|
||||
|
||||
clang::SourceRange getReadRange(clang::SourceManager& SM, const clang::SourceRange& exprRange) {
|
||||
clang::SourceRange getReadRange(clang::SourceManager &SM, const clang::SourceRange &exprRange) {
|
||||
clang::SourceLocation begin = exprRange.getBegin();
|
||||
clang::SourceLocation end = exprRange.getEnd();
|
||||
bool beginSafe = !SM.isMacroBodyExpansion(begin) || clang::Lexer::isAtStartOfMacroExpansion(begin, SM, clang::LangOptions{});
|
||||
@@ -167,7 +176,7 @@ clang::SourceRange getReadRange(clang::SourceManager& SM, const clang::SourceRan
|
||||
}
|
||||
}
|
||||
|
||||
clang::SourceRange getWriteRange(clang::SourceManager& SM, const clang::SourceRange& exprRange) {
|
||||
clang::SourceRange getWriteRange(clang::SourceManager &SM, const clang::SourceRange &exprRange) {
|
||||
clang::SourceLocation begin = exprRange.getBegin();
|
||||
clang::SourceLocation end = exprRange.getEnd();
|
||||
// If the range is contained within a macro, update the macro definition.
|
||||
@@ -178,7 +187,7 @@ clang::SourceRange getWriteRange(clang::SourceManager& SM, const clang::SourceRa
|
||||
return {SM.getSpellingLoc(begin), SM.getSpellingLoc(end)};
|
||||
}
|
||||
|
||||
StringRef readSourceText(clang::SourceManager& SM, const clang::SourceRange& exprRange) {
|
||||
StringRef readSourceText(clang::SourceManager &SM, const clang::SourceRange &exprRange) {
|
||||
return clang::Lexer::getSourceText(clang::CharSourceRange::getTokenRange(getReadRange(SM, exprRange)), SM, clang::LangOptions(), nullptr);
|
||||
}
|
||||
|
||||
@@ -186,53 +195,50 @@ StringRef readSourceText(clang::SourceManager& SM, const clang::SourceRange& exp
|
||||
* Get a string representation of the expression `arg`, unless it's a defaulting function
|
||||
* call argument, in which case get a 0. Used for building argument lists to kernel calls.
|
||||
*/
|
||||
std::string stringifyZeroDefaultedArg(clang::SourceManager& SM, const clang::Expr* arg) {
|
||||
if (clang::isa<clang::CXXDefaultArgExpr>(arg)) {
|
||||
return "0";
|
||||
} else {
|
||||
return readSourceText(SM, arg->getSourceRange());
|
||||
}
|
||||
std::string stringifyZeroDefaultedArg(clang::SourceManager &SM, const clang::Expr *arg) {
|
||||
if (clang::isa<clang::CXXDefaultArgExpr>(arg)) return "0";
|
||||
else return readSourceText(SM, arg->getSourceRange());
|
||||
}
|
||||
|
||||
} // 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 (hipToken.hipName == s_hiprand_kernel_h) {
|
||||
if (insertedRAND_kernelHeader) return true;
|
||||
insertedRAND_kernelHeader = true;
|
||||
return false;
|
||||
} else if (hipToken.hipName == "hiprand.h") {
|
||||
if (insertedRANDHeader) { return true; }
|
||||
} else if (hipToken.hipName == s_hiprand_h) {
|
||||
if (insertedRANDHeader) return true;
|
||||
insertedRANDHeader = true;
|
||||
return false;
|
||||
}
|
||||
case API_DNN:
|
||||
if (insertedDNNHeader) { return true; }
|
||||
if (insertedDNNHeader) return true;
|
||||
insertedDNNHeader = true;
|
||||
return false;
|
||||
case API_FFT:
|
||||
if (insertedFFTHeader) { return true; }
|
||||
if (insertedFFTHeader) return true;
|
||||
insertedFFTHeader = true;
|
||||
return false;
|
||||
case API_COMPLEX:
|
||||
if (insertedComplexHeader) { return true; }
|
||||
if (insertedComplexHeader) return true;
|
||||
insertedComplexHeader = true;
|
||||
return false;
|
||||
case API_SPARSE:
|
||||
if (insertedSPARSEHeader) { return true; }
|
||||
if (insertedSPARSEHeader) return true;
|
||||
insertedSPARSEHeader = true;
|
||||
return false;
|
||||
default:
|
||||
@@ -240,13 +246,11 @@ bool HipifyAction::Exclude(const hipCounter & hipToken) {
|
||||
}
|
||||
return false;
|
||||
case CONV_INCLUDE:
|
||||
if (hipToken.hipName.empty()) {
|
||||
return true;
|
||||
}
|
||||
if (hipToken.hipName.empty()) return true;
|
||||
switch (hipToken.apiType) {
|
||||
case API_RAND:
|
||||
if (hipToken.hipName == "hiprand_kernel.h") {
|
||||
if (insertedRAND_kernelHeader) { return true; }
|
||||
if (hipToken.hipName == s_hiprand_kernel_h) {
|
||||
if (insertedRAND_kernelHeader) return true;
|
||||
insertedRAND_kernelHeader = true;
|
||||
}
|
||||
return false;
|
||||
@@ -267,24 +271,19 @@ void HipifyAction::InclusionDirective(clang::SourceLocation hash_loc,
|
||||
clang::CharSourceRange filename_range,
|
||||
const clang::FileEntry*, StringRef,
|
||||
StringRef, const clang::Module*) {
|
||||
clang::SourceManager& SM = getCompilerInstance().getSourceManager();
|
||||
if (!SM.isWrittenInMainFile(hash_loc)) {
|
||||
return;
|
||||
}
|
||||
auto &SM = getCompilerInstance().getSourceManager();
|
||||
if (!SM.isWrittenInMainFile(hash_loc)) return;
|
||||
if (!firstHeader) {
|
||||
firstHeader = true;
|
||||
firstHeaderLoc = hash_loc;
|
||||
}
|
||||
const auto found = CUDA_INCLUDE_MAP.find(file_name);
|
||||
if (found == CUDA_INCLUDE_MAP.end()) {
|
||||
return;
|
||||
}
|
||||
if (found == CUDA_INCLUDE_MAP.end()) return;
|
||||
bool exclude = Exclude(found->second);
|
||||
Statistics::current().incrementCounter(found->second, file_name.str());
|
||||
|
||||
clang::SourceLocation sl = filename_range.getBegin();
|
||||
if (Statistics::isUnsupported(found->second)) {
|
||||
clang::DiagnosticsEngine& DE = getCompilerInstance().getDiagnostics();
|
||||
clang::DiagnosticsEngine &DE = getCompilerInstance().getDiagnostics();
|
||||
DE.Report(sl, DE.getCustomDiagID(clang::DiagnosticsEngine::Warning, "Unsupported CUDA header."));
|
||||
return;
|
||||
}
|
||||
@@ -293,11 +292,8 @@ void HipifyAction::InclusionDirective(clang::SourceLocation hash_loc,
|
||||
if (!exclude) {
|
||||
clang::SmallString<128> includeBuffer;
|
||||
llvm::StringRef name = Statistics::isToRoc(found->second) ? found->second.rocName : found->second.hipName;
|
||||
if (is_angled) {
|
||||
newInclude = llvm::Twine("<" + name+ ">").toStringRef(includeBuffer);
|
||||
} else {
|
||||
newInclude = llvm::Twine("\"" + name + "\"").toStringRef(includeBuffer);
|
||||
}
|
||||
if (is_angled) newInclude = llvm::Twine("<" + name+ ">").toStringRef(includeBuffer);
|
||||
else newInclude = llvm::Twine("\"" + name + "\"").toStringRef(includeBuffer);
|
||||
} else {
|
||||
// hashLoc is location of the '#', thus replacing the whole include directive by empty newInclude starting with '#'.
|
||||
sl = hash_loc;
|
||||
@@ -309,46 +305,33 @@ void HipifyAction::InclusionDirective(clang::SourceLocation hash_loc,
|
||||
}
|
||||
|
||||
void HipifyAction::PragmaDirective(clang::SourceLocation Loc, clang::PragmaIntroducerKind Introducer) {
|
||||
if (pragmaOnce) {
|
||||
return;
|
||||
}
|
||||
clang::SourceManager& SM = getCompilerInstance().getSourceManager();
|
||||
if (!SM.isWrittenInMainFile(Loc)) {
|
||||
return;
|
||||
}
|
||||
clang::Preprocessor& PP = getCompilerInstance().getPreprocessor();
|
||||
if (pragmaOnce) return;
|
||||
auto &SM = getCompilerInstance().getSourceManager();
|
||||
if (!SM.isWrittenInMainFile(Loc)) return;
|
||||
clang::Preprocessor &PP = getCompilerInstance().getPreprocessor();
|
||||
clang::Token tok;
|
||||
PP.Lex(tok);
|
||||
StringRef Text(SM.getCharacterData(tok.getLocation()), tok.getLength());
|
||||
if (Text == "once") {
|
||||
if (Text == sOnce) {
|
||||
pragmaOnce = true;
|
||||
pragmaOnceLoc = tok.getEndLoc();
|
||||
}
|
||||
}
|
||||
|
||||
bool HipifyAction::cudaLaunchKernel(const clang::ast_matchers::MatchFinder::MatchResult& Result) {
|
||||
StringRef refName = "cudaLaunchKernel";
|
||||
const auto* launchKernel = Result.Nodes.getNodeAs<clang::CUDAKernelCallExpr>(refName);
|
||||
if (!launchKernel) {
|
||||
return false;
|
||||
}
|
||||
const clang::Expr* calleeExpr = launchKernel->getCallee();
|
||||
if (!calleeExpr) {
|
||||
return false;
|
||||
}
|
||||
const clang::FunctionDecl *caleeDecl = launchKernel->getDirectCallee();
|
||||
if (!caleeDecl) {
|
||||
return false;
|
||||
}
|
||||
const clang::CallExpr* config = launchKernel->getConfig();
|
||||
if (!config) {
|
||||
return false;
|
||||
}
|
||||
bool HipifyAction::cudaLaunchKernel(const mat::MatchFinder::MatchResult &Result) {
|
||||
auto *launchKernel = Result.Nodes.getNodeAs<clang::CUDAKernelCallExpr>(sCudaLaunchKernel);
|
||||
if (!launchKernel) return false;
|
||||
auto *calleeExpr = launchKernel->getCallee();
|
||||
if (!calleeExpr) return false;
|
||||
auto *caleeDecl = launchKernel->getDirectCallee();
|
||||
if (!caleeDecl) return false;
|
||||
auto *config = launchKernel->getConfig();
|
||||
if (!config) return false;
|
||||
clang::SmallString<40> XStr;
|
||||
llvm::raw_svector_ostream OS(XStr);
|
||||
clang::LangOptions DefaultLangOptions;
|
||||
clang::SourceManager* SM = Result.SourceManager;
|
||||
OS << sHipLaunchKernelGGL;
|
||||
auto *SM = Result.SourceManager;
|
||||
OS << sHipLaunchKernelGGL << "(";
|
||||
if (caleeDecl->isTemplateInstantiation()) OS << sHIP_KERNEL_NAME << "(";
|
||||
OS << readSourceText(*SM, calleeExpr->getSourceRange());
|
||||
if (caleeDecl->isTemplateInstantiation()) OS << ")";
|
||||
@@ -381,29 +364,24 @@ 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_KERNEL_LAUNCH, ApiTypes::API_RUNTIME};
|
||||
Statistics::current().incrementCounter(counter, refName.str());
|
||||
hipCounter counter = {sHipLaunchKernelGGL, "", ConvTypes::CONV_KERNEL_LAUNCH, ApiTypes::API_RUNTIME};
|
||||
Statistics::current().incrementCounter(counter, sCudaLaunchKernel.str());
|
||||
return true;
|
||||
}
|
||||
|
||||
bool HipifyAction::cudaSharedIncompleteArrayVar(const clang::ast_matchers::MatchFinder::MatchResult& Result) {
|
||||
StringRef refName = "cudaSharedIncompleteArrayVar";
|
||||
auto* sharedVar = Result.Nodes.getNodeAs<clang::VarDecl>(refName);
|
||||
if (!sharedVar) {
|
||||
return false;
|
||||
}
|
||||
bool HipifyAction::cudaSharedIncompleteArrayVar(const mat::MatchFinder::MatchResult &Result) {
|
||||
auto *sharedVar = Result.Nodes.getNodeAs<clang::VarDecl>(sCudaSharedIncompleteArrayVar);
|
||||
if (!sharedVar) return false;
|
||||
// Example: extern __shared__ uint sRadix1[];
|
||||
if (!sharedVar->hasExternalFormalLinkage()) {
|
||||
return false;
|
||||
}
|
||||
if (!sharedVar->hasExternalFormalLinkage()) return false;
|
||||
clang::QualType QT = sharedVar->getType();
|
||||
std::string typeName;
|
||||
if (QT->isIncompleteArrayType()) {
|
||||
const clang::ArrayType* AT = QT.getTypePtr()->getAsArrayTypeUnsafe();
|
||||
const clang::ArrayType *AT = QT.getTypePtr()->getAsArrayTypeUnsafe();
|
||||
QT = AT->getElementType();
|
||||
if (QT.getTypePtr()->isBuiltinType()) {
|
||||
QT = QT.getCanonicalType();
|
||||
const auto* BT = clang::dyn_cast<clang::BuiltinType>(QT);
|
||||
auto *BT = clang::dyn_cast<clang::BuiltinType>(QT);
|
||||
if (BT) {
|
||||
clang::LangOptions LO;
|
||||
LO.CUDA = true;
|
||||
@@ -417,7 +395,7 @@ bool HipifyAction::cudaSharedIncompleteArrayVar(const clang::ast_matchers::Match
|
||||
if (!typeName.empty()) {
|
||||
clang::SourceLocation slStart = sharedVar->getOuterLocStart();
|
||||
clang::SourceLocation slEnd = llcompat::getEndLoc(sharedVar->getTypeSourceInfo()->getTypeLoc());
|
||||
clang::SourceManager* SM = Result.SourceManager;
|
||||
auto *SM = Result.SourceManager;
|
||||
size_t repLength = SM->getCharacterData(slEnd) - SM->getCharacterData(slStart) + 1;
|
||||
std::string varName = sharedVar->getNameAsString();
|
||||
std::string repName = sHIP_DYNAMIC_SHARED + "(" + typeName + ", " + varName + ")";
|
||||
@@ -425,52 +403,39 @@ bool HipifyAction::cudaSharedIncompleteArrayVar(const clang::ast_matchers::Match
|
||||
clang::FullSourceLoc fullSL(slStart, *SM);
|
||||
insertReplacement(Rep, fullSL);
|
||||
hipCounter counter = {sHIP_DYNAMIC_SHARED, "", ConvTypes::CONV_EXTERN_SHARED, ApiTypes::API_RUNTIME};
|
||||
Statistics::current().incrementCounter(counter, refName.str());
|
||||
Statistics::current().incrementCounter(counter, sCudaSharedIncompleteArrayVar.str());
|
||||
return true;
|
||||
}
|
||||
return false;
|
||||
}
|
||||
|
||||
bool HipifyAction::cudaDeviceFuncCall(const clang::ast_matchers::MatchFinder::MatchResult& Result) {
|
||||
if (const clang::CallExpr *call = Result.Nodes.getNodeAs<clang::CallExpr>("cudaDeviceFuncCall")) {
|
||||
const clang::FunctionDecl *funcDcl = call->getDirectCallee();
|
||||
if (!funcDcl) {
|
||||
return false;
|
||||
}
|
||||
bool HipifyAction::cudaDeviceFuncCall(const mat::MatchFinder::MatchResult &Result) {
|
||||
if (const clang::CallExpr *call = Result.Nodes.getNodeAs<clang::CallExpr>(sCudaDeviceFuncCall)) {
|
||||
auto *funcDcl = call->getDirectCallee();
|
||||
if (!funcDcl) return false;
|
||||
FindAndReplace(funcDcl->getDeclName().getAsString(), llcompat::getBeginLoc(call), CUDA_DEVICE_FUNC_MAP, false);
|
||||
return true;
|
||||
}
|
||||
return false;
|
||||
}
|
||||
|
||||
bool HipifyAction::cubNamespacePrefix(const clang::ast_matchers::MatchFinder::MatchResult& Result) {
|
||||
if (const clang::TypedefNameDecl *decl = Result.Nodes.getNodeAs<clang::TypedefNameDecl>("cubNamespacePrefix")) {
|
||||
if (!decl) {
|
||||
return false;
|
||||
}
|
||||
bool HipifyAction::cubNamespacePrefix(const mat::MatchFinder::MatchResult &Result) {
|
||||
if (auto *decl = Result.Nodes.getNodeAs<clang::TypedefNameDecl>(sCubNamespacePrefix)) {
|
||||
clang::QualType QT = decl->getUnderlyingType();
|
||||
const clang::Type* t = QT.getTypePtr();
|
||||
if (!t) {
|
||||
return false;
|
||||
}
|
||||
const clang::ElaboratedType* et = t->getAs<clang::ElaboratedType>();
|
||||
if (!et) {
|
||||
return false;
|
||||
}
|
||||
auto *t = QT.getTypePtr();
|
||||
if (!t) return false;
|
||||
const clang::ElaboratedType *et = t->getAs<clang::ElaboratedType>();
|
||||
if (!et) return false;
|
||||
const clang::NestedNameSpecifier *nns = et->getQualifier();
|
||||
if (!nns) {
|
||||
return false;
|
||||
}
|
||||
if (!nns) return false;
|
||||
const clang::NamespaceDecl *nsd = nns->getAsNamespace();
|
||||
if (!nsd) {
|
||||
return false;
|
||||
}
|
||||
if (!nsd) return false;
|
||||
const clang::TypeSourceInfo *si = decl->getTypeSourceInfo();
|
||||
const clang::TypeLoc tloc = si->getTypeLoc();
|
||||
const clang::SourceRange sr = tloc.getSourceRange();
|
||||
clang::SourceLocation sl(sr.getBegin());
|
||||
clang::SourceLocation end(sr.getEnd());
|
||||
clang::SourceManager& SM = getCompilerInstance().getSourceManager();
|
||||
auto &SM = getCompilerInstance().getSourceManager();
|
||||
size_t length = SM.getCharacterData(end) - SM.getCharacterData(sl);
|
||||
StringRef sfull = StringRef(SM.getCharacterData(sl), length);
|
||||
std::string name = nsd->getDeclName().getAsString();
|
||||
@@ -484,15 +449,11 @@ bool HipifyAction::cubNamespacePrefix(const clang::ast_matchers::MatchFinder::Ma
|
||||
return false;
|
||||
}
|
||||
|
||||
bool HipifyAction::cudaHostFuncCall(const clang::ast_matchers::MatchFinder::MatchResult& Result) {
|
||||
if (const clang::CallExpr * call = Result.Nodes.getNodeAs<clang::CallExpr>("cudaHostFuncCall")) {
|
||||
if (!call->getNumArgs()) {
|
||||
return false;
|
||||
}
|
||||
const clang::FunctionDecl* funcDcl = call->getDirectCallee();
|
||||
if (!funcDcl) {
|
||||
return false;
|
||||
}
|
||||
bool HipifyAction::cudaHostFuncCall(const mat::MatchFinder::MatchResult &Result) {
|
||||
if (auto *call = Result.Nodes.getNodeAs<clang::CallExpr>(sCudaHostFuncCall)) {
|
||||
if (!call->getNumArgs()) return false;
|
||||
auto *funcDcl = call->getDirectCallee();
|
||||
if (!funcDcl) return false;
|
||||
std::string sName = funcDcl->getDeclName().getAsString();
|
||||
unsigned int argNum = 0;
|
||||
bool b_reinterpret = (ReinterpretFunctions.find(sName) != ReinterpretFunctions.end()) ? true : false;
|
||||
@@ -506,7 +467,7 @@ bool HipifyAction::cudaHostFuncCall(const clang::ast_matchers::MatchFinder::Matc
|
||||
clang::SmallString<40> XStr;
|
||||
llvm::raw_svector_ostream OS(XStr);
|
||||
clang::SourceRange sr = call->getArg(argNum)->getSourceRange();
|
||||
clang::SourceManager* SM = Result.SourceManager;
|
||||
auto *SM = Result.SourceManager;
|
||||
OS << (b_reinterpret ? s_reinterpret_cast : sHIP_SYMBOL) << "(" << readSourceText(*SM, sr) << ")";
|
||||
clang::SourceRange replacementRange = getWriteRange(*SM, { sr.getBegin(), sr.getEnd() });
|
||||
clang::SourceLocation s = replacementRange.getBegin();
|
||||
@@ -521,7 +482,7 @@ bool HipifyAction::cudaHostFuncCall(const clang::ast_matchers::MatchFinder::Matc
|
||||
return false;
|
||||
}
|
||||
|
||||
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();
|
||||
@@ -530,10 +491,10 @@ void HipifyAction::insertReplacement(const ct::Replacement& rep, const clang::Fu
|
||||
}
|
||||
}
|
||||
|
||||
std::unique_ptr<clang::ASTConsumer> HipifyAction::CreateASTConsumer(clang::CompilerInstance& CI, llvm::StringRef) {
|
||||
Finder.reset(new clang::ast_matchers::MatchFinder);
|
||||
std::unique_ptr<clang::ASTConsumer> HipifyAction::CreateASTConsumer(clang::CompilerInstance &CI, StringRef) {
|
||||
Finder.reset(new mat::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(sCudaLaunchKernel), this);
|
||||
Finder->addMatcher(
|
||||
mat::varDecl(
|
||||
mat::isExpansionInMainFile(),
|
||||
@@ -541,7 +502,7 @@ std::unique_ptr<clang::ASTConsumer> HipifyAction::CreateASTConsumer(clang::Compi
|
||||
mat::hasAttr(clang::attr::CUDAShared),
|
||||
mat::hasType(mat::incompleteArrayType())
|
||||
)
|
||||
).bind("cudaSharedIncompleteArrayVar"),
|
||||
).bind(sCudaSharedIncompleteArrayVar),
|
||||
this
|
||||
);
|
||||
Finder->addMatcher(
|
||||
@@ -561,7 +522,7 @@ std::unique_ptr<clang::ASTConsumer> HipifyAction::CreateASTConsumer(clang::Compi
|
||||
)
|
||||
)
|
||||
)
|
||||
).bind("cudaHostFuncCall"),
|
||||
).bind(sCudaHostFuncCall),
|
||||
this
|
||||
);
|
||||
Finder->addMatcher(
|
||||
@@ -576,7 +537,7 @@ std::unique_ptr<clang::ASTConsumer> HipifyAction::CreateASTConsumer(clang::Compi
|
||||
mat::unless(mat::hasAttr(clang::attr::CUDAHost))
|
||||
)
|
||||
)
|
||||
).bind("cudaDeviceFuncCall"),
|
||||
).bind(sCudaDeviceFuncCall),
|
||||
this
|
||||
);
|
||||
Finder->addMatcher(
|
||||
@@ -586,12 +547,12 @@ std::unique_ptr<clang::ASTConsumer> HipifyAction::CreateASTConsumer(clang::Compi
|
||||
mat::elaboratedType(
|
||||
mat::hasQualifier(
|
||||
mat::specifiesNamespace(
|
||||
mat::hasName("cub")
|
||||
mat::hasName(sCub)
|
||||
)
|
||||
)
|
||||
)
|
||||
)
|
||||
).bind("cubNamespacePrefix"),
|
||||
).bind(sCubNamespacePrefix),
|
||||
this
|
||||
);
|
||||
// Ownership is transferred to the caller.
|
||||
@@ -599,10 +560,8 @@ std::unique_ptr<clang::ASTConsumer> HipifyAction::CreateASTConsumer(clang::Compi
|
||||
}
|
||||
|
||||
void HipifyAction::Ifndef(clang::SourceLocation Loc, const clang::Token &MacroNameTok, const clang::MacroDefinition &MD) {
|
||||
clang::SourceManager& SM = getCompilerInstance().getSourceManager();
|
||||
if (!SM.isWrittenInMainFile(Loc)) {
|
||||
return;
|
||||
}
|
||||
auto &SM = getCompilerInstance().getSourceManager();
|
||||
if (!SM.isWrittenInMainFile(Loc)) return;
|
||||
StringRef Text(SM.getCharacterData(MacroNameTok.getLocation()), MacroNameTok.getLength());
|
||||
Ifndefs.insert(std::make_pair(Text.str(), MacroNameTok.getEndLoc()));
|
||||
}
|
||||
@@ -615,12 +574,12 @@ void HipifyAction::EndSourceFileAction() {
|
||||
// one copy of the hip include into every file.
|
||||
bool placeForIncludeCalculated = false;
|
||||
clang::SourceLocation sl, controllingMacroLoc;
|
||||
clang::SourceManager& SM = getCompilerInstance().getSourceManager();
|
||||
clang::Preprocessor& PP = getCompilerInstance().getPreprocessor();
|
||||
clang::HeaderSearch& HS = PP.getHeaderSearchInfo();
|
||||
clang::ExternalPreprocessorSource* EPL = HS.getExternalLookup();
|
||||
const clang::FileEntry* FE = SM.getFileEntryForID(SM.getMainFileID());
|
||||
const clang::IdentifierInfo* controllingMacro = HS.getFileInfo(FE).getControllingMacro(EPL);
|
||||
auto &SM = getCompilerInstance().getSourceManager();
|
||||
clang::Preprocessor &PP = getCompilerInstance().getPreprocessor();
|
||||
clang::HeaderSearch &HS = PP.getHeaderSearchInfo();
|
||||
clang::ExternalPreprocessorSource *EPL = HS.getExternalLookup();
|
||||
const clang::FileEntry *FE = SM.getFileEntryForID(SM.getMainFileID());
|
||||
const clang::IdentifierInfo *controllingMacro = HS.getFileInfo(FE).getControllingMacro(EPL);
|
||||
if (controllingMacro) {
|
||||
auto found = Ifndefs.find(controllingMacro->getName().str());
|
||||
if (found != Ifndefs.end()) {
|
||||
@@ -629,19 +588,13 @@ void HipifyAction::EndSourceFileAction() {
|
||||
}
|
||||
}
|
||||
if (pragmaOnce) {
|
||||
if (placeForIncludeCalculated) {
|
||||
sl = pragmaOnceLoc < controllingMacroLoc ? pragmaOnceLoc : controllingMacroLoc;
|
||||
} else {
|
||||
sl = pragmaOnceLoc;
|
||||
}
|
||||
if (placeForIncludeCalculated) sl = pragmaOnceLoc < controllingMacroLoc ? pragmaOnceLoc : controllingMacroLoc;
|
||||
else sl = pragmaOnceLoc;
|
||||
placeForIncludeCalculated = true;
|
||||
}
|
||||
if (!placeForIncludeCalculated) {
|
||||
if (firstHeader) {
|
||||
sl = firstHeaderLoc;
|
||||
} else {
|
||||
sl = SM.getLocForStartOfFile(SM.getMainFileID());
|
||||
}
|
||||
if (firstHeader) sl = firstHeaderLoc;
|
||||
else sl = SM.getLocForStartOfFile(SM.getMainFileID());
|
||||
}
|
||||
clang::FullSourceLoc fullSL(sl, SM);
|
||||
ct::Replacement Rep(SM, sl, 0, "\n#include <hip/hip_runtime.h>\n");
|
||||
@@ -656,15 +609,15 @@ namespace {
|
||||
* A silly little class to proxy PPCallbacks back to the HipifyAction class.
|
||||
*/
|
||||
class PPCallbackProxy : public clang::PPCallbacks {
|
||||
HipifyAction& hipifyAction;
|
||||
HipifyAction &hipifyAction;
|
||||
|
||||
public:
|
||||
explicit PPCallbackProxy(HipifyAction& action): hipifyAction(action) {}
|
||||
explicit PPCallbackProxy(HipifyAction &action): hipifyAction(action) {}
|
||||
|
||||
void InclusionDirective(clang::SourceLocation hash_loc, const clang::Token& include_token,
|
||||
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
|
||||
const clang::FileEntry *file, StringRef search_path, StringRef relative_path,
|
||||
const clang::Module *imported
|
||||
#if LLVM_VERSION_MAJOR > 6
|
||||
, clang::SrcMgr::CharacteristicKind FileType
|
||||
#endif
|
||||
@@ -688,10 +641,10 @@ bool HipifyAction::BeginInvocation(clang::CompilerInstance &CI) {
|
||||
}
|
||||
|
||||
void HipifyAction::ExecuteAction() {
|
||||
clang::Preprocessor& PP = getCompilerInstance().getPreprocessor();
|
||||
clang::SourceManager& SM = getCompilerInstance().getSourceManager();
|
||||
clang::Preprocessor &PP = getCompilerInstance().getPreprocessor();
|
||||
auto &SM = getCompilerInstance().getSourceManager();
|
||||
// Start lexing the specified input file.
|
||||
const llvm::MemoryBuffer* FromFile = SM.getBuffer(SM.getMainFileID());
|
||||
const llvm::MemoryBuffer *FromFile = SM.getBuffer(SM.getMainFileID());
|
||||
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
|
||||
@@ -709,7 +662,7 @@ void HipifyAction::ExecuteAction() {
|
||||
clang::ASTFrontendAction::ExecuteAction();
|
||||
}
|
||||
|
||||
void HipifyAction::run(const clang::ast_matchers::MatchFinder::MatchResult& Result) {
|
||||
void HipifyAction::run(const mat::MatchFinder::MatchResult &Result) {
|
||||
if (cudaLaunchKernel(Result)) return;
|
||||
if (cudaSharedIncompleteArrayVar(Result)) return;
|
||||
if (cudaHostFuncCall(Result)) return;
|
||||
|
||||
@@ -31,17 +31,18 @@ THE SOFTWARE.
|
||||
#include "Statistics.h"
|
||||
|
||||
namespace ct = clang::tooling;
|
||||
namespace mat = clang::ast_matchers;
|
||||
using namespace llvm;
|
||||
|
||||
/**
|
||||
* A FrontendAction that hipifies CUDA programs.
|
||||
*/
|
||||
class HipifyAction : public clang::ASTFrontendAction,
|
||||
public clang::ast_matchers::MatchFinder::MatchCallback {
|
||||
public mat::MatchFinder::MatchCallback {
|
||||
private:
|
||||
ct::Replacements* replacements;
|
||||
ct::Replacements *replacements;
|
||||
std::map<std::string, clang::SourceLocation> Ifndefs;
|
||||
std::unique_ptr<clang::ast_matchers::MatchFinder> Finder;
|
||||
std::unique_ptr<mat::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.
|
||||
@@ -67,12 +68,11 @@ public:
|
||||
explicit HipifyAction(ct::Replacements *replacements): clang::ASTFrontendAction(),
|
||||
replacements(replacements) {}
|
||||
// MatchCallback listeners
|
||||
bool cudaBuiltin(const clang::ast_matchers::MatchFinder::MatchResult& Result);
|
||||
bool cudaLaunchKernel(const clang::ast_matchers::MatchFinder::MatchResult& Result);
|
||||
bool cudaSharedIncompleteArrayVar(const clang::ast_matchers::MatchFinder::MatchResult& Result);
|
||||
bool cudaDeviceFuncCall(const clang::ast_matchers::MatchFinder::MatchResult& Result);
|
||||
bool cudaHostFuncCall(const clang::ast_matchers::MatchFinder::MatchResult& Result);
|
||||
bool cubNamespacePrefix(const clang::ast_matchers::MatchFinder::MatchResult& Result);
|
||||
bool cudaLaunchKernel(const mat::MatchFinder::MatchResult &Result);
|
||||
bool cudaSharedIncompleteArrayVar(const mat::MatchFinder::MatchResult &Result);
|
||||
bool cudaDeviceFuncCall(const mat::MatchFinder::MatchResult &Result);
|
||||
bool cudaHostFuncCall(const mat::MatchFinder::MatchResult &Result);
|
||||
bool cubNamespacePrefix(const mat::MatchFinder::MatchResult &Result);
|
||||
// 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,
|
||||
@@ -91,7 +91,7 @@ public:
|
||||
|
||||
protected:
|
||||
// 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);
|
||||
void insertReplacement(const ct::Replacement &rep, const clang::FullSourceLoc &fullSL);
|
||||
// FrontendAction entry point.
|
||||
void ExecuteAction() override;
|
||||
// Callback before starting processing a single input; used by hipify-clang for setting Preprocessor options.
|
||||
@@ -99,8 +99,8 @@ protected:
|
||||
// Called at the start of each new file to process.
|
||||
void EndSourceFileAction() override;
|
||||
// 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;
|
||||
bool Exclude(const hipCounter & hipToken);
|
||||
void FindAndReplace(llvm::StringRef name, clang::SourceLocation sl, const std::map<llvm::StringRef, hipCounter>& repMap, bool bReplace = true);
|
||||
void run(const mat::MatchFinder::MatchResult &Result) override;
|
||||
std::unique_ptr<clang::ASTConsumer> CreateASTConsumer(clang::CompilerInstance &CI, StringRef InFile) override;
|
||||
bool Exclude(const hipCounter &hipToken);
|
||||
void FindAndReplace(StringRef name, clang::SourceLocation sl, const std::map<StringRef, hipCounter> &repMap, bool bReplace = true);
|
||||
};
|
||||
|
||||
Reference in New Issue
Block a user