diff --git a/hipamd/hipify-clang/src/CUDA2HipMap.cpp b/hipamd/hipify-clang/src/CUDA2HipMap.cpp index 8e76b5fdde..1893f6ce5b 100644 --- a/hipamd/hipify-clang/src/CUDA2HipMap.cpp +++ b/hipamd/hipify-clang/src/CUDA2HipMap.cpp @@ -1461,35 +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}}, - - {"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..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,21 +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. - 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; }