AI/rocm-systems
2
0
Fork 0
Código Incidencias Pull Requests Acciones Paquetes Proyectos Lanzamientos Wiki Actividad

Merge pull request #283 from emankov/master

Explorar el Código 
[HIPIFY] Disable cudaBuiltin matcher.
Este commit está contenido en:
Ben Sander
2017-12-07 12:38:02 -08:00
cometido por GitHub
padre faef491778 45befb73d3
commit f245fe8a6c
Se han modificado 2 ficheros con 0 adiciones y 88 borrados
Descargar archivo de parche Descargar archivo de diferencias Expandir todos los archivos Contraer todos los archivos
hipamd/hipify-clang/src/CUDA2HipMap.cpp
-29
Ver fichero
@@ -1461,35 +1461,6 @@ const std::map<llvm::StringRef, hipCounter> CUDA_IDENTIFIER_MAP{
{"cudaHostRegisterMapped", {"hipHostRegisterMapped", CONV_MEM, API_RUNTIME}},
{"cudaHostRegisterIoMemory", {"hipHostRegisterIoMemory", CONV_MEM, API_RUNTIME}},
// Coordinate Indexing and Dimensions
{"threadIdx.x", {"hipThreadIdx_x", CONV_COORD_FUNC, API_RUNTIME}},
{"threadIdx.y", {"hipThreadIdx_y", CONV_COORD_FUNC, API_RUNTIME}},
{"threadIdx.z", {"hipThreadIdx_z", CONV_COORD_FUNC, API_RUNTIME}},
{"blockIdx.x", {"hipBlockIdx_x", CONV_COORD_FUNC, API_RUNTIME}},
{"blockIdx.y", {"hipBlockIdx_y", CONV_COORD_FUNC, API_RUNTIME}},
{"blockIdx.z", {"hipBlockIdx_z", CONV_COORD_FUNC, API_RUNTIME}},
{"blockDim.x", {"hipBlockDim_x", CONV_COORD_FUNC, API_RUNTIME}},
{"blockDim.y", {"hipBlockDim_y", CONV_COORD_FUNC, API_RUNTIME}},
{"blockDim.z", {"hipBlockDim_z", CONV_COORD_FUNC, API_RUNTIME}},
{"gridDim.x", {"hipGridDim_x", CONV_COORD_FUNC, API_RUNTIME}},
{"gridDim.y", {"hipGridDim_y", CONV_COORD_FUNC, API_RUNTIME}},
{"gridDim.z", {"hipGridDim_z", CONV_COORD_FUNC, API_RUNTIME}},
{"blockIdx.x", {"hipBlockIdx_x", CONV_COORD_FUNC, API_RUNTIME}},
{"blockIdx.y", {"hipBlockIdx_y", CONV_COORD_FUNC, API_RUNTIME}},
{"blockIdx.z", {"hipBlockIdx_z", CONV_COORD_FUNC, API_RUNTIME}},
{"blockDim.x", {"hipBlockDim_x", CONV_COORD_FUNC, API_RUNTIME}},
{"blockDim.y", {"hipBlockDim_y", CONV_COORD_FUNC, API_RUNTIME}},
{"blockDim.z", {"hipBlockDim_z", CONV_COORD_FUNC, API_RUNTIME}},
{"gridDim.x", {"hipGridDim_x", CONV_COORD_FUNC, API_RUNTIME}},
{"gridDim.y", {"hipGridDim_y", CONV_COORD_FUNC, API_RUNTIME}},
{"gridDim.z", {"hipGridDim_z", CONV_COORD_FUNC, API_RUNTIME}},
{"warpSize", {"hipWarpSize", CONV_SPECIAL_FUNC, API_RUNTIME}},
// Events
hipamd/hipify-clang/src/HipifyAction.cpp
-59
Ver fichero
@@ -251,49 +251,6 @@ bool HipifyAction::cudaLaunchKernel(const clang::ast_matchers::MatchFinder::Matc
return true;
}
bool HipifyAction::cudaBuiltin(const clang::ast_matchers::MatchFinder::MatchResult& Result) {
const clang::MemberExpr* threadIdx = Result.Nodes.getNodeAs<clang::MemberExpr>("cudaBuiltin");
if (!threadIdx) {
return false;
}
const clang::OpaqueValueExpr* refBase = clang::dyn_cast<clang::OpaqueValueExpr>(threadIdx->getBase());
if (!refBase) {
return false;
}
const clang::DeclRefExpr* declRef = clang::dyn_cast<clang::DeclRefExpr>(refBase->getSourceExpr());
if (!declRef) {
return false;
}
clang::SourceLocation sl = threadIdx->getLocStart();
clang::SourceManager* SM = Result.SourceManager;
StringRef name = declRef->getDecl()->getName();
StringRef memberName = threadIdx->getMemberDecl()->getName();
size_t pos = memberName.find_first_not_of("__fetch_builtin_");
memberName = memberName.slice(pos, memberName.size());
clang::SmallString<128> tmpData;
name = clang::Twine(name + "." + memberName).toStringRef(tmpData);
const auto found = CUDA_IDENTIFIER_MAP.find(name);
if (found != CUDA_IDENTIFIER_MAP.end()) {
Statistics::current().incrementCounter(found->second, name.str());
if (!found->second.unsupported) {
StringRef repName = found->second.hipName;
ct::Replacement Rep(*SM, sl, name.size(), repName);
clang::FullSourceLoc fullSL(sl, *SM);
insertReplacement(Rep, fullSL);
}
} else {
clang::DiagnosticsEngine& DE = getCompilerInstance().getDiagnostics();
const auto ID = DE.getCustomDiagID(clang::DiagnosticsEngine::Warning, "Unknown CUDA builtin");
DE.Report(sl, ID);
}
return true;
}
bool HipifyAction::cudaSharedIncompleteArrayVar(const clang::ast_matchers::MatchFinder::MatchResult& Result) {
StringRef refName = "cudaSharedIncompleteArrayVar";
auto* sharedVar = Result.Nodes.getNodeAs<clang::VarDecl>(refName);
@@ -357,21 +314,6 @@ std::unique_ptr<clang::ASTConsumer> HipifyAction::CreateASTConsumer(clang::Compi
// Replace the <<<...>>> language extension with a hip kernel launch
Finder->addMatcher(mat::cudaKernelCallExpr(mat::isExpansionInMainFile()).bind("cudaLaunchKernel"), this);
// Replace cuda builtins.
Finder->addMatcher(
mat::memberExpr(
mat::isExpansionInMainFile(),
mat::hasObjectExpression(
mat::hasType(
mat::cxxRecordDecl(
mat::matchesName("__cuda_builtin_")
)
)
)
).bind("cudaBuiltin"),
this
);
Finder->addMatcher(
mat::varDecl(
mat::isExpansionInMainFile(),
@@ -454,7 +396,6 @@ void HipifyAction::ExecuteAction() {
}
void HipifyAction::run(const clang::ast_matchers::MatchFinder::MatchResult& Result) {
if (cudaBuiltin(Result)) return;
if (cudaLaunchKernel(Result)) return;
if (cudaSharedIncompleteArrayVar(Result)) return;
}