From 43f0a14005f2fe3debb6af1da501249eb41ba591 Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Tue, 5 Dec 2017 19:46:53 +0300 Subject: [PATCH 1/3] [HIPIFY] remove duplicates from CUDA_IDENTIFIER_MAP --- hipamd/hipify-clang/src/CUDA2HipMap.cpp | 12 ------------ 1 file changed, 12 deletions(-) diff --git a/hipamd/hipify-clang/src/CUDA2HipMap.cpp b/hipamd/hipify-clang/src/CUDA2HipMap.cpp index 8e76b5fdde..0413d76022 100644 --- a/hipamd/hipify-clang/src/CUDA2HipMap.cpp +++ b/hipamd/hipify-clang/src/CUDA2HipMap.cpp @@ -1478,18 +1478,6 @@ const std::map CUDA_IDENTIFIER_MAP{ {"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 From ccc959fdc9d4b37f870c09390043632137ef9d26 Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Tue, 5 Dec 2017 20:28:51 +0300 Subject: [PATCH 2/3] [HIPIFY] Disable cudaBuiltin matcher. As HIP has started to support vanilla CUDA syntax for threadIdx, blockIdx, blockDim and gridDim. Other CUDA builtins are not tracked for now. --- hipamd/hipify-clang/src/CUDA2HipMap.cpp | 17 ----------------- hipamd/hipify-clang/src/HipifyAction.cpp | 6 ++++-- 2 files changed, 4 insertions(+), 19 deletions(-) diff --git a/hipamd/hipify-clang/src/CUDA2HipMap.cpp b/hipamd/hipify-clang/src/CUDA2HipMap.cpp index 0413d76022..1893f6ce5b 100644 --- a/hipamd/hipify-clang/src/CUDA2HipMap.cpp +++ b/hipamd/hipify-clang/src/CUDA2HipMap.cpp @@ -1461,23 +1461,6 @@ const std::map 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}}, - {"warpSize", {"hipWarpSize", CONV_SPECIAL_FUNC, API_RUNTIME}}, // Events diff --git a/hipamd/hipify-clang/src/HipifyAction.cpp b/hipamd/hipify-clang/src/HipifyAction.cpp index 192dd00949..f4c7fd7764 100644 --- a/hipamd/hipify-clang/src/HipifyAction.cpp +++ b/hipamd/hipify-clang/src/HipifyAction.cpp @@ -358,6 +358,8 @@ std::unique_ptr HipifyAction::CreateASTConsumer(clang::Compi Finder->addMatcher(mat::cudaKernelCallExpr(mat::isExpansionInMainFile()).bind("cudaLaunchKernel"), this); // Replace cuda builtins. + /* Disable as HIP has started to support vanilla CUDA syntax for threadIdx, blockIdx, blockDim and gridDim. */ + /* Other CUDA builtins are not tracked for now. Finder->addMatcher( mat::memberExpr( mat::isExpansionInMainFile(), @@ -371,7 +373,7 @@ std::unique_ptr HipifyAction::CreateASTConsumer(clang::Compi ).bind("cudaBuiltin"), this ); - + */ Finder->addMatcher( mat::varDecl( mat::isExpansionInMainFile(), @@ -454,7 +456,7 @@ void HipifyAction::ExecuteAction() { } void HipifyAction::run(const clang::ast_matchers::MatchFinder::MatchResult& Result) { - if (cudaBuiltin(Result)) return; +// if (cudaBuiltin(Result)) return; if (cudaLaunchKernel(Result)) return; if (cudaSharedIncompleteArrayVar(Result)) return; } From 45befb73d36a4bcb44925475dfac671a53d031de Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Wed, 6 Dec 2017 20:22:14 +0300 Subject: [PATCH 3/3] [HIPIFY] Remove cudaBuiltin matcher --- hipamd/hipify-clang/src/HipifyAction.cpp | 61 ------------------------ 1 file changed, 61 deletions(-) diff --git a/hipamd/hipify-clang/src/HipifyAction.cpp b/hipamd/hipify-clang/src/HipifyAction.cpp index f4c7fd7764..ee23387e1f 100644 --- a/hipamd/hipify-clang/src/HipifyAction.cpp +++ b/hipamd/hipify-clang/src/HipifyAction.cpp @@ -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("cudaBuiltin"); - if (!threadIdx) { - return false; - } - - const clang::OpaqueValueExpr* refBase = clang::dyn_cast(threadIdx->getBase()); - if (!refBase) { - return false; - } - - const clang::DeclRefExpr* declRef = clang::dyn_cast(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(refName); @@ -357,23 +314,6 @@ std::unique_ptr 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. - /* Disable as HIP has started to support vanilla CUDA syntax for threadIdx, blockIdx, blockDim and gridDim. */ - /* Other CUDA builtins are not tracked for now. - 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(), @@ -456,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; }