Revert "[CUDA][HIP][OpenMP] Emit deferred diagnostics by a post-parsing AST travese"

This reverts commit 1b978dd.

Author: yxsamliu

clang/include/clang/Sema/Sema.h

@@ -1464,12 +1464,6 @@ class Sema final {
 
   void emitAndClearUnusedLocalTypedefWarnings();
 
-  // Emit all deferred diagnostics.
-  void emitDeferredDiags();
-  // Emit any deferred diagnostics for FD and erase them from the map in which
-  // they're stored.
-  void emitDeferredDiags(FunctionDecl *FD, bool ShowCallStack);
-
   enum TUFragmentKind {
     /// The global module fragment, between 'module;' and a module-declaration.
     Global,
@@ -3689,8 +3683,7 @@ class Sema final {
     TemplateDiscarded, // Discarded due to uninstantiated templates
     Unknown,
   };
-  FunctionEmissionStatus getEmissionStatus(FunctionDecl *Decl,
-                                           bool Final = false);
+  FunctionEmissionStatus getEmissionStatus(FunctionDecl *Decl);
 
   // Whether the callee should be ignored in CUDA/HIP/OpenMP host/device check.
   bool shouldIgnoreInHostDeviceCheck(FunctionDecl *Callee);
@@ -9684,10 +9677,22 @@ class Sema final {
   /// Pop OpenMP function region for non-capturing function.
   void popOpenMPFunctionRegion(const sema::FunctionScopeInfo *OldFSI);
 
+  /// Check whether we're allowed to call Callee from the current function.
+  void checkOpenMPDeviceFunction(SourceLocation Loc, FunctionDecl *Callee,
+                                 bool CheckForDelayedContext = true);
+
+  /// Check whether we're allowed to call Callee from the current function.
+  void checkOpenMPHostFunction(SourceLocation Loc, FunctionDecl *Callee,
+                               bool CheckCaller = true);
+
   /// Check if the expression is allowed to be used in expressions for the
   /// OpenMP devices.
   void checkOpenMPDeviceExpr(const Expr *E);
 
+  /// Finishes analysis of the deferred functions calls that may be declared as
+  /// host/nohost during device/host compilation.
+  void finalizeOpenMPDelayedAnalysis();
+
   /// Checks if a type or a declaration is disabled due to the owning extension
   /// being disabled, and emits diagnostic messages if it is disabled.
   /// \param D type or declaration to be checked.
@@ -9870,11 +9875,6 @@ class Sema final {
   void
   checkDeclIsAllowedInOpenMPTarget(Expr *E, Decl *D,
                                    SourceLocation IdLoc = SourceLocation());
-  /// Finishes analysis of the deferred functions calls that may be declared as
-  /// host/nohost during device/host compilation.
-  void finalizeOpenMPDelayedAnalysis(const FunctionDecl *Caller,
-                                     const FunctionDecl *Callee,
-                                     SourceLocation Loc);
   /// Return true inside OpenMP declare target region.
   bool isInOpenMPDeclareTargetContext() const {
     return DeclareTargetNestingLevel > 0;
@@ -11223,6 +11223,18 @@ class Sema final {
                  /* Caller = */ FunctionDeclAndLoc>
       DeviceKnownEmittedFns;
 
+  /// A partial call graph maintained during CUDA/OpenMP device code compilation
+  /// to support deferred diagnostics.
+  ///
+  /// Functions are only added here if, at the time they're considered, they are
+  /// not known-emitted.  As soon as we discover that a function is
+  /// known-emitted, we remove it and everything it transitively calls from this
+  /// set and add those functions to DeviceKnownEmittedFns.
+  llvm::DenseMap</* Caller = */ CanonicalDeclPtr<FunctionDecl>,
+                 /* Callees = */ llvm::MapVector<CanonicalDeclPtr<FunctionDecl>,
+                                                 SourceLocation>>
+      DeviceCallGraph;
+
   /// Diagnostic builder for CUDA/OpenMP devices errors which may or may not be
   /// deferred.
   ///
@@ -11297,6 +11309,14 @@ class Sema final {
     llvm::Optional<unsigned> PartialDiagId;
   };
 
+  /// Indicate that this function (and thus everything it transtively calls)
+  /// will be codegen'ed, and emit any deferred diagnostics on this function and
+  /// its (transitive) callees.
+  void markKnownEmitted(
+      Sema &S, FunctionDecl *OrigCaller, FunctionDecl *OrigCallee,
+      SourceLocation OrigLoc,
+      const llvm::function_ref<bool(Sema &, FunctionDecl *)> IsKnownEmitted);
+
   /// Creates a DeviceDiagBuilder that emits the diagnostic if the current context
   /// is "used as device code".
   ///

clang/lib/Sema/Sema.cpp

@@ -11,7 +11,6 @@
 //
 //===----------------------------------------------------------------------===//
 
-#include "UsedDeclVisitor.h"
 #include "clang/AST/ASTContext.h"
 #include "clang/AST/ASTDiagnostic.h"
 #include "clang/AST/DeclCXX.h"
@@ -955,7 +954,9 @@ void Sema::ActOnEndOfTranslationUnitFragment(TUFragmentKind Kind) {
     PerformPendingInstantiations();
   }
 
-  emitDeferredDiags();
+  // Finalize analysis of OpenMP-specific constructs.
+  if (LangOpts.OpenMP)
+    finalizeOpenMPDelayedAnalysis();
 
   assert(LateParsedInstantiations.empty() &&
          "end of TU template instantiation should not create more "
@@ -1450,128 +1451,27 @@ static void emitCallStackNotes(Sema &S, FunctionDecl *FD) {
 
 // Emit any deferred diagnostics for FD and erase them from the map in which
 // they're stored.
-void Sema::emitDeferredDiags(FunctionDecl *FD, bool ShowCallStack) {
-  auto It = DeviceDeferredDiags.find(FD);
-  if (It == DeviceDeferredDiags.end())
+static void emitDeferredDiags(Sema &S, FunctionDecl *FD, bool ShowCallStack) {
+  auto It = S.DeviceDeferredDiags.find(FD);
+  if (It == S.DeviceDeferredDiags.end())
     return;
   bool HasWarningOrError = false;
   for (PartialDiagnosticAt &PDAt : It->second) {
     const SourceLocation &Loc = PDAt.first;
     const PartialDiagnostic &PD = PDAt.second;
-    HasWarningOrError |= getDiagnostics().getDiagnosticLevel(
+    HasWarningOrError |= S.getDiagnostics().getDiagnosticLevel(
                              PD.getDiagID(), Loc) >= DiagnosticsEngine::Warning;
-    DiagnosticBuilder Builder(Diags.Report(Loc, PD.getDiagID()));
+    DiagnosticBuilder Builder(S.Diags.Report(Loc, PD.getDiagID()));
     Builder.setForceEmit();
     PD.Emit(Builder);
   }
+  S.DeviceDeferredDiags.erase(It);
 
   // FIXME: Should this be called after every warning/error emitted in the loop
   // above, instead of just once per function?  That would be consistent with
   // how we handle immediate errors, but it also seems like a bit much.
   if (HasWarningOrError && ShowCallStack)
-    emitCallStackNotes(*this, FD);
-}
-
-namespace {
-/// Helper class that emits deferred diagnostic messages if an entity directly
-/// or indirectly using the function that causes the deferred diagnostic
-/// messages is known to be emitted.
-class DeferredDiagnosticsEmitter
-    : public UsedDeclVisitor<DeferredDiagnosticsEmitter> {
-public:
-  typedef UsedDeclVisitor<DeferredDiagnosticsEmitter> Inherited;
-  llvm::SmallSet<CanonicalDeclPtr<Decl>, 4> Visited;
-  llvm::SmallVector<CanonicalDeclPtr<FunctionDecl>, 4> UseStack;
-  bool ShouldEmit;
-  unsigned InOMPDeviceContext;
-
-  DeferredDiagnosticsEmitter(Sema &S)
-      : Inherited(S), ShouldEmit(false), InOMPDeviceContext(0) {}
-
-  void VisitDeclRefExpr(DeclRefExpr *E) {
-    if (FunctionDecl *FD = dyn_cast<FunctionDecl>(E->getDecl())) {
-      visitUsedDecl(E->getLocation(), FD);
-    }
-  }
-
-  void VisitMemberExpr(MemberExpr *E) {
-    if (FunctionDecl *FD = dyn_cast<FunctionDecl>(E->getMemberDecl()))
-      visitUsedDecl(E->getMemberLoc(), FD);
-  }
-
-  void VisitOMPTargetDirective(OMPTargetDirective *Node) {
-    ++InOMPDeviceContext;
-    Inherited::VisitOMPTargetDirective(Node);
-    --InOMPDeviceContext;
-  }
-
-  void VisitCapturedStmt(CapturedStmt *Node) {
-    visitUsedDecl(Node->getBeginLoc(), Node->getCapturedDecl());
-    Inherited::VisitCapturedStmt(Node);
-  }
-
-  void visitUsedDecl(SourceLocation Loc, Decl *D) {
-    if (auto *TD = dyn_cast<TranslationUnitDecl>(D)) {
-      for (auto *DD : TD->decls()) {
-        visitUsedDecl(Loc, DD);
-      }
-    } else if (auto *FTD = dyn_cast<FunctionTemplateDecl>(D)) {
-      for (auto *DD : FTD->specializations()) {
-        visitUsedDecl(Loc, DD);
-      }
-    } else if (auto *FD = dyn_cast<FunctionDecl>(D)) {
-      FunctionDecl *Caller = UseStack.empty() ? nullptr : UseStack.back();
-      auto IsKnownEmitted = S.getEmissionStatus(FD, /*Final=*/true) ==
-                            Sema::FunctionEmissionStatus::Emitted;
-      if (!Caller)
-        ShouldEmit = IsKnownEmitted;
-      if ((!ShouldEmit && !S.getLangOpts().OpenMP && !Caller) ||
-          S.shouldIgnoreInHostDeviceCheck(FD) || Visited.count(D))
-        return;
-      // Finalize analysis of OpenMP-specific constructs.
-      if (Caller && S.LangOpts.OpenMP && UseStack.size() == 1)
-        S.finalizeOpenMPDelayedAnalysis(Caller, FD, Loc);
-      if (Caller)
-        S.DeviceKnownEmittedFns[FD] = {Caller, Loc};
-      if (ShouldEmit || InOMPDeviceContext)
-        S.emitDeferredDiags(FD, Caller);
-      Visited.insert(D);
-      UseStack.push_back(FD);
-      if (auto *S = FD->getBody()) {
-        this->Visit(S);
-      }
-      UseStack.pop_back();
-      Visited.erase(D);
-    } else if (auto *RD = dyn_cast<RecordDecl>(D)) {
-      for (auto *DD : RD->decls()) {
-        visitUsedDecl(Loc, DD);
-      }
-    } else if (auto *CD = dyn_cast<CapturedDecl>(D)) {
-      if (auto *S = CD->getBody()) {
-        this->Visit(S);
-      }
-    } else if (auto *VD = dyn_cast<VarDecl>(D)) {
-      if (auto *Init = VD->getInit()) {
-        auto DevTy = OMPDeclareTargetDeclAttr::getDeviceType(VD);
-        bool IsDev = DevTy && (*DevTy == OMPDeclareTargetDeclAttr::DT_NoHost ||
-                               *DevTy == OMPDeclareTargetDeclAttr::DT_Any);
-        if (IsDev)
-          ++InOMPDeviceContext;
-        this->Visit(Init);
-        if (IsDev)
-          --InOMPDeviceContext;
-      }
-    }
-  }
-};
-} // namespace
-
-void Sema::emitDeferredDiags() {
-  if (DeviceDeferredDiags.empty() && !LangOpts.OpenMP)
-    return;
-
-  DeferredDiagnosticsEmitter(*this).visitUsedDecl(
-      SourceLocation(), Context.getTranslationUnitDecl());
+    emitCallStackNotes(S, FD);
 }
 
 // In CUDA, there are some constructs which may appear in semantically-valid
@@ -1644,6 +1544,71 @@ Sema::DeviceDiagBuilder::~DeviceDiagBuilder() {
   }
 }
 
+// Indicate that this function (and thus everything it transtively calls) will
+// be codegen'ed, and emit any deferred diagnostics on this function and its
+// (transitive) callees.
+void Sema::markKnownEmitted(
+    Sema &S, FunctionDecl *OrigCaller, FunctionDecl *OrigCallee,
+    SourceLocation OrigLoc,
+    const llvm::function_ref<bool(Sema &, FunctionDecl *)> IsKnownEmitted) {
+  // Nothing to do if we already know that FD is emitted.
+  if (IsKnownEmitted(S, OrigCallee)) {
+    assert(!S.DeviceCallGraph.count(OrigCallee));
+    return;
+  }
+
+  // We've just discovered that OrigCallee is known-emitted.  Walk our call
+  // graph to see what else we can now discover also must be emitted.
+
+  struct CallInfo {
+    FunctionDecl *Caller;
+    FunctionDecl *Callee;
+    SourceLocation Loc;
+  };
+  llvm::SmallVector<CallInfo, 4> Worklist = {{OrigCaller, OrigCallee, OrigLoc}};
+  llvm::SmallSet<CanonicalDeclPtr<FunctionDecl>, 4> Seen;
+  Seen.insert(OrigCallee);
+  while (!Worklist.empty()) {
+    CallInfo C = Worklist.pop_back_val();
+    assert(!IsKnownEmitted(S, C.Callee) &&
+           "Worklist should not contain known-emitted functions.");
+    S.DeviceKnownEmittedFns[C.Callee] = {C.Caller, C.Loc};
+    emitDeferredDiags(S, C.Callee, C.Caller);
+
+    // If this is a template instantiation, explore its callgraph as well:
+    // Non-dependent calls are part of the template's callgraph, while dependent
+    // calls are part of to the instantiation's call graph.
+    if (auto *Templ = C.Callee->getPrimaryTemplate()) {
+      FunctionDecl *TemplFD = Templ->getAsFunction();
+      if (!Seen.count(TemplFD) && !S.DeviceKnownEmittedFns.count(TemplFD)) {
+        Seen.insert(TemplFD);
+        Worklist.push_back(
+            {/* Caller = */ C.Caller, /* Callee = */ TemplFD, C.Loc});
+      }
+    }
+
+    // Add all functions called by Callee to our worklist.
+    auto CGIt = S.DeviceCallGraph.find(C.Callee);
+    if (CGIt == S.DeviceCallGraph.end())
+      continue;
+
+    for (std::pair<CanonicalDeclPtr<FunctionDecl>, SourceLocation> FDLoc :
+         CGIt->second) {
+      FunctionDecl *NewCallee = FDLoc.first;
+      SourceLocation CallLoc = FDLoc.second;
+      if (Seen.count(NewCallee) || IsKnownEmitted(S, NewCallee))
+        continue;
+      Seen.insert(NewCallee);
+      Worklist.push_back(
+          {/* Caller = */ C.Callee, /* Callee = */ NewCallee, CallLoc});
+    }
+
+    // C.Callee is now known-emitted, so we no longer need to maintain its list
+    // of callees in DeviceCallGraph.
+    S.DeviceCallGraph.erase(CGIt);
+  }
+}
+
 Sema::DeviceDiagBuilder Sema::targetDiag(SourceLocation Loc, unsigned DiagID) {
   if (LangOpts.OpenMP)
     return LangOpts.OpenMPIsDevice ? diagIfOpenMPDeviceCode(Loc, DiagID)

clang/lib/Sema/SemaCUDA.cpp

@@ -674,6 +674,25 @@ bool Sema::CheckCUDACall(SourceLocation Loc, FunctionDecl *Callee) {
   // Otherwise, mark the call in our call graph so we can traverse it later.
   bool CallerKnownEmitted =
       getEmissionStatus(Caller) == FunctionEmissionStatus::Emitted;
+  if (CallerKnownEmitted) {
+    // Host-side references to a __global__ function refer to the stub, so the
+    // function itself is never emitted and therefore should not be marked.
+    if (!shouldIgnoreInHostDeviceCheck(Callee))
+      markKnownEmitted(
+          *this, Caller, Callee, Loc, [](Sema &S, FunctionDecl *FD) {
+            return S.getEmissionStatus(FD) == FunctionEmissionStatus::Emitted;
+          });
+  } else {
+    // If we have
+    //   host fn calls kernel fn calls host+device,
+    // the HD function does not get instantiated on the host.  We model this by
+    // omitting at the call to the kernel from the callgraph.  This ensures
+    // that, when compiling for host, only HD functions actually called from the
+    // host get marked as known-emitted.
+    if (!shouldIgnoreInHostDeviceCheck(Callee))
+      DeviceCallGraph[Caller].insert({Callee, Loc});
+  }
+
   DeviceDiagBuilder::Kind DiagKind = [this, Caller, Callee,
                                       CallerKnownEmitted] {
     switch (IdentifyCUDAPreference(Caller, Callee)) {

clang/lib/Sema/SemaDecl.cpp

@@ -17929,8 +17929,7 @@ Decl *Sema::getObjCDeclContext() const {
   return (dyn_cast_or_null<ObjCContainerDecl>(CurContext));
 }
 
-Sema::FunctionEmissionStatus Sema::getEmissionStatus(FunctionDecl *FD,
-                                                     bool Final) {
+Sema::FunctionEmissionStatus Sema::getEmissionStatus(FunctionDecl *FD) {
   // Templates are emitted when they're instantiated.
   if (FD->isDependentContext())
     return FunctionEmissionStatus::TemplateDiscarded;
@@ -17942,10 +17941,8 @@ Sema::FunctionEmissionStatus Sema::getEmissionStatus(FunctionDecl *FD,
     if (DevTy.hasValue()) {
       if (*DevTy == OMPDeclareTargetDeclAttr::DT_Host)
         OMPES = FunctionEmissionStatus::OMPDiscarded;
-      else if (*DevTy == OMPDeclareTargetDeclAttr::DT_NoHost ||
-               *DevTy == OMPDeclareTargetDeclAttr::DT_Any) {
+      else if (DeviceKnownEmittedFns.count(FD) > 0)
         OMPES = FunctionEmissionStatus::Emitted;
-      }
     }
   } else if (LangOpts.OpenMP) {
     // In OpenMP 4.5 all the functions are host functions.
@@ -17961,11 +17958,10 @@ Sema::FunctionEmissionStatus Sema::getEmissionStatus(FunctionDecl *FD,
       if (DevTy.hasValue()) {
         if (*DevTy == OMPDeclareTargetDeclAttr::DT_NoHost) {
           OMPES = FunctionEmissionStatus::OMPDiscarded;
-        } else if (*DevTy == OMPDeclareTargetDeclAttr::DT_Host ||
-                   *DevTy == OMPDeclareTargetDeclAttr::DT_Any)
+        } else if (DeviceKnownEmittedFns.count(FD) > 0) {
           OMPES = FunctionEmissionStatus::Emitted;
-      } else if (Final)
-        OMPES = FunctionEmissionStatus::Emitted;
+        }
+      }
     }
   }
   if (OMPES == FunctionEmissionStatus::OMPDiscarded ||
@@ -18000,7 +17996,9 @@ Sema::FunctionEmissionStatus Sema::getEmissionStatus(FunctionDecl *FD,
 
   // Otherwise, the function is known-emitted if it's in our set of
   // known-emitted functions.
-  return FunctionEmissionStatus::Unknown;
+  return (DeviceKnownEmittedFns.count(FD) > 0)
+             ? FunctionEmissionStatus::Emitted
+             : FunctionEmissionStatus::Unknown;
 }
 
 bool Sema::shouldIgnoreInHostDeviceCheck(FunctionDecl *Callee) {