From 5d92a6c252a26d025e0121dfa5437dec2eb931ea Mon Sep 17 00:00:00 2001
From: Evgeny Mankov <Evgeniy.Mankov@amd.com>
Date: Wed, 6 Dec 2017 20:22:14 +0300
Subject: [PATCH] [HIPIFY] Remove cudaBuiltin matcher

---
 hipify-clang/src/HipifyAction.cpp | 61 -------------------------------
 1 file changed, 61 deletions(-)

diff --git a/hipify-clang/src/HipifyAction.cpp b/hipify-clang/src/HipifyAction.cpp
index f4c7fd7764..ee23387e1f 100644
--- a/hipify-clang/src/HipifyAction.cpp
+++ b/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<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,23 +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.
-    /* 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;
 }