[SEMA]Generalize deferred diagnostic interface, NFC.

[ABataev]

Summary

Deferred diagnostic interface is going to be used for OpenMP device
compilation. Generalized previously existed deferred diagnostic
interface for CUDA to be used with OpenMP and, possibly, other models.

[tra]

LGTM. I've added jlebar@ as he's originally written the code.

[rjmccall]

LGTM as well.

cfe/trunk/include/clang/Sema/Sema.h

public:
   /// before incrementing, so you can emit an error.
   bool PopForceCUDAHostDevice();
 
   /// Diagnostics that are emitted only if we discover that the given function
   /// must be codegen'ed.  Because handling these correctly adds overhead to
   /// compilation, this is currently only enabled for CUDA compilations.
   llvm::DenseMap<CanonicalDeclPtr<FunctionDecl>,
                  std::vector<PartialDiagnosticAt>>
-      CUDADeferredDiags;
+      DeviceDeferredDiags;
 
   /// A pair of a canonical FunctionDecl and a SourceLocation.  When used as the
   /// key in a hashtable, both the FD and location are hashed.
   struct FunctionDeclAndLoc {
     CanonicalDeclPtr<FunctionDecl> FD;
     SourceLocation Loc;
   };
 
   /// FunctionDecls and SourceLocations for which CheckCUDACall has emitted a
   /// (maybe deferred) "bad call" diagnostic.  We use this to avoid emitting the
   /// same deferred diag twice.
   llvm::DenseSet<FunctionDeclAndLoc> LocsWithCUDACallDiags;
 
   /// An inverse call graph, mapping known-emitted functions to one of their
   /// known-emitted callers (plus the location of the call).
   ///
   /// Functions that we can tell a priori must be emitted aren't added to this
   /// map.
   llvm::DenseMap</* Callee = */ CanonicalDeclPtr<FunctionDecl>,
                  /* Caller = */ FunctionDeclAndLoc>
-      CUDAKnownEmittedFns;
+      DeviceKnownEmittedFns;
 
-  /// A partial call graph maintained during CUDA compilation to support
-  /// deferred diagnostics.
+  /// 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 CUDAKnownEmittedFns.
+  /// set and add those functions to DeviceKnownEmittedFns.
   llvm::DenseMap</* Caller = */ CanonicalDeclPtr<FunctionDecl>,
                  /* Callees = */ llvm::MapVector<CanonicalDeclPtr<FunctionDecl>,
                                                  SourceLocation>>
-      CUDACallGraph;
+      DeviceCallGraph;
 
-  /// Diagnostic builder for CUDA errors which may or may not be deferred.
+  /// Diagnostic builder for CUDA/OpenMP devices errors which may or may not be
+  /// deferred.
   ///
   /// In CUDA, there exist constructs (e.g. variable-length arrays, try/catch)
   /// which are not allowed to appear inside __device__ functions and are
   /// allowed to appear in __host__ __device__ functions only if the host+device
   /// function is never codegen'ed.
   ///
   /// To handle this, we use the notion of "deferred diagnostics", where we
   /// attach a diagnostic to a FunctionDecl that's emitted iff it's codegen'ed.
   ///
   /// This class lets you emit either a regular diagnostic, a deferred
   /// diagnostic, or no diagnostic at all, according to an argument you pass to
   /// its constructor, thus simplifying the process of creating these "maybe
   /// deferred" diagnostics.
-  class CUDADiagBuilder {
+  class DeviceDiagBuilder {
   public:
     enum Kind {
       /// Emit no diagnostics.
       K_Nop,
       /// Emit the diagnostic immediately (i.e., behave like Sema::Diag()).
       K_Immediate,
       /// Emit the diagnostic immediately, and, if it's a warning or error, also
       /// emit a call stack showing how this function can be reached by an a
       /// priori known-emitted function.
       K_ImmediateWithCallStack,
       /// Create a deferred diagnostic, which is emitted only if the function
       /// it's attached to is codegen'ed.  Also emit a call stack as with
       /// K_ImmediateWithCallStack.
       K_Deferred
     };
 
-    CUDADiagBuilder(Kind K, SourceLocation Loc, unsigned DiagID,
+    DeviceDiagBuilder(Kind K, SourceLocation Loc, unsigned DiagID,
                       FunctionDecl *Fn, Sema &S);
-    ~CUDADiagBuilder();
+    ~DeviceDiagBuilder();
 
     /// Convertible to bool: True if we immediately emitted an error, false if
     /// we didn't emit an error or we created a deferred error.
     ///
     /// Example usage:
     ///
-    ///   if (CUDADiagBuilder(...) << foo << bar)
+    ///   if (DeviceDiagBuilder(...) << foo << bar)
     ///     return ExprError();
     ///
     /// But see CUDADiagIfDeviceCode() and CUDADiagIfHostCode() -- you probably
-    /// want to use these instead of creating a CUDADiagBuilder yourself.
+    /// want to use these instead of creating a DeviceDiagBuilder yourself.
     operator bool() const { return ImmediateDiag.hasValue(); }
 
     template <typename T>
-    friend const CUDADiagBuilder &operator<<(const CUDADiagBuilder &Diag,
+    friend const DeviceDiagBuilder &operator<<(const DeviceDiagBuilder &Diag,
                                                const T &Value) {
       if (Diag.ImmediateDiag.hasValue())
         *Diag.ImmediateDiag << Value;
       else if (Diag.PartialDiag.hasValue())
         *Diag.PartialDiag << Value;
       return Diag;
     }
 
   private:
     Sema &S;
     SourceLocation Loc;
     unsigned DiagID;
     FunctionDecl *Fn;
     bool ShowCallStack;
 
     // Invariant: At most one of these Optionals has a value.
     // FIXME: Switch these to a Variant once that exists.
     llvm::Optional<SemaDiagnosticBuilder> ImmediateDiag;
     llvm::Optional<PartialDiagnostic> PartialDiag;
   };
 
-  /// Creates a CUDADiagBuilder that emits the diagnostic if the current context
+  /// 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".
   ///
   /// - If CurContext is a __host__ function, does not emit any diagnostics.
   /// - If CurContext is a __device__ or __global__ function, emits the
   ///   diagnostics immediately.
   /// - If CurContext is a __host__ __device__ function and we are compiling for
   ///   the device, creates a diagnostic which is emitted if and when we realize
   ///   that the function will be codegen'ed.
   ///
   /// Example usage:
   ///
   ///  // Variable-length arrays are not allowed in CUDA device code.
   ///  if (CUDADiagIfDeviceCode(Loc, diag::err_cuda_vla) << CurrentCUDATarget())
   ///    return ExprError();
   ///  // Otherwise, continue parsing as normal.
-  CUDADiagBuilder CUDADiagIfDeviceCode(SourceLocation Loc, unsigned DiagID);
+  DeviceDiagBuilder CUDADiagIfDeviceCode(SourceLocation Loc, unsigned DiagID);
 
-  /// Creates a CUDADiagBuilder that emits the diagnostic if the current context
+  /// Creates a DeviceDiagBuilder that emits the diagnostic if the current context
   /// is "used as host code".
   ///
   /// Same as CUDADiagIfDeviceCode, with "host" and "device" switched.
-  CUDADiagBuilder CUDADiagIfHostCode(SourceLocation Loc, unsigned DiagID);
+  DeviceDiagBuilder CUDADiagIfHostCode(SourceLocation Loc, unsigned DiagID);
 
   enum CUDAFunctionTarget {
     CFT_Device,
     CFT_Global,
     CFT_Host,
     CFT_HostDevice,
     CFT_InvalidTarget
   };

cfe/trunk/lib/Sema/Sema.cpp

 Sema::SemaDiagnosticBuilder
 Sema::Diag(SourceLocation Loc, const PartialDiagnostic& PD) {
   SemaDiagnosticBuilder Builder(Diag(Loc, PD.getDiagID()));
   PD.Emit(Builder);
 
   return Builder;
 }
 
+// Print notes showing how we can reach FD starting from an a priori
+// known-callable function.
+static void emitCallStackNotes(Sema &S, FunctionDecl *FD) {
+  auto FnIt = S.DeviceKnownEmittedFns.find(FD);
+  while (FnIt != S.DeviceKnownEmittedFns.end()) {
+    DiagnosticBuilder Builder(
+        S.Diags.Report(FnIt->second.Loc, diag::note_called_by));
+    Builder << FnIt->second.FD;
+    Builder.setForceEmit();
+
+    FnIt = S.DeviceKnownEmittedFns.find(FnIt->second.FD);
+  }
+}
+
+// Emit any deferred diagnostics for FD and erase them from the map in which
+// they're stored.
+static void emitDeferredDiags(Sema &S, FunctionDecl *FD) {
+  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 |= S.getDiagnostics().getDiagnosticLevel(
+                             PD.getDiagID(), Loc) >= DiagnosticsEngine::Warning;
+    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)
+    emitCallStackNotes(S, FD);
+}
+
+// In CUDA, there are some constructs which may appear in semantically-valid
+// code, but trigger errors if we ever generate code for the function in which
+// they appear.  Essentially every construct you're not allowed to use on the
+// device falls into this category, because you are allowed to use these
+// constructs in a __host__ __device__ function, but only if that function is
+// never codegen'ed on the device.
+//
+// To handle semantic checking for these constructs, we keep track of the set of
+// functions we know will be emitted, either because we could tell a priori that
+// they would be emitted, or because they were transitively called by a
+// known-emitted function.
+//
+// We also keep a partial call graph of which not-known-emitted functions call
+// which other not-known-emitted functions.
+//
+// When we see something which is illegal if the current function is emitted
+// (usually by way of CUDADiagIfDeviceCode, CUDADiagIfHostCode, or
+// CheckCUDACall), we first check if the current function is known-emitted.  If
+// so, we immediately output the diagnostic.
+//
+// Otherwise, we "defer" the diagnostic.  It sits in Sema::DeviceDeferredDiags
+// until we discover that the function is known-emitted, at which point we take
+// it out of this map and emit the diagnostic.
+
+Sema::DeviceDiagBuilder::DeviceDiagBuilder(Kind K, SourceLocation Loc,
+                                           unsigned DiagID, FunctionDecl *Fn,
+                                           Sema &S)
+    : S(S), Loc(Loc), DiagID(DiagID), Fn(Fn),
+      ShowCallStack(K == K_ImmediateWithCallStack || K == K_Deferred) {
+  switch (K) {
+  case K_Nop:
+    break;
+  case K_Immediate:
+  case K_ImmediateWithCallStack:
+    ImmediateDiag.emplace(S.Diag(Loc, DiagID));
+    break;
+  case K_Deferred:
+    assert(Fn && "Must have a function to attach the deferred diag to.");
+    PartialDiag.emplace(S.PDiag(DiagID));
+    break;
+  }
+}
+
+Sema::DeviceDiagBuilder::~DeviceDiagBuilder() {
+  if (ImmediateDiag) {
+    // Emit our diagnostic and, if it was a warning or error, output a callstack
+    // if Fn isn't a priori known-emitted.
+    bool IsWarningOrError = S.getDiagnostics().getDiagnosticLevel(
+                                DiagID, Loc) >= DiagnosticsEngine::Warning;
+    ImmediateDiag.reset(); // Emit the immediate diag.
+    if (IsWarningOrError && ShowCallStack)
+      emitCallStackNotes(S, Fn);
+  } else if (PartialDiag) {
+    assert(ShowCallStack && "Must always show call stack for deferred diags.");
+    S.DeviceDeferredDiags[Fn].push_back({Loc, std::move(*PartialDiag)});
+  }
+}
+
+// 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);
+
+    // 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);
+  }
+}
+
 /// Looks through the macro-expansion chain for the given
 /// location, looking for a macro expansion with the given name.
 /// If one is found, returns true and sets the location to that
 /// expansion loc.
 bool Sema::findMacroSpelling(SourceLocation &locref, StringRef name) {
   SourceLocation loc = locref;
   if (!loc.isMacroID()) return false;
 

cfe/trunk/lib/Sema/SemaCUDA.cpp

if (It != Previous.end()) {
     }
     return;
   }
 
   NewD->addAttr(CUDAHostAttr::CreateImplicit(Context));
   NewD->addAttr(CUDADeviceAttr::CreateImplicit(Context));
 }
 
-// In CUDA, there are some constructs which may appear in semantically-valid
-// code, but trigger errors if we ever generate code for the function in which
-// they appear.  Essentially every construct you're not allowed to use on the
-// device falls into this category, because you are allowed to use these
-// constructs in a __host__ __device__ function, but only if that function is
-// never codegen'ed on the device.
-//
-// To handle semantic checking for these constructs, we keep track of the set of
-// functions we know will be emitted, either because we could tell a priori that
-// they would be emitted, or because they were transitively called by a
-// known-emitted function.
-//
-// We also keep a partial call graph of which not-known-emitted functions call
-// which other not-known-emitted functions.
-//
-// When we see something which is illegal if the current function is emitted
-// (usually by way of CUDADiagIfDeviceCode, CUDADiagIfHostCode, or
-// CheckCUDACall), we first check if the current function is known-emitted.  If
-// so, we immediately output the diagnostic.
-//
-// Otherwise, we "defer" the diagnostic.  It sits in Sema::CUDADeferredDiags
-// until we discover that the function is known-emitted, at which point we take
-// it out of this map and emit the diagnostic.
-
-Sema::CUDADiagBuilder::CUDADiagBuilder(Kind K, SourceLocation Loc,
-                                       unsigned DiagID, FunctionDecl *Fn,
-                                       Sema &S)
-    : S(S), Loc(Loc), DiagID(DiagID), Fn(Fn),
-      ShowCallStack(K == K_ImmediateWithCallStack || K == K_Deferred) {
-  switch (K) {
-  case K_Nop:
-    break;
-  case K_Immediate:
-  case K_ImmediateWithCallStack:
-    ImmediateDiag.emplace(S.Diag(Loc, DiagID));
-    break;
-  case K_Deferred:
-    assert(Fn && "Must have a function to attach the deferred diag to.");
-    PartialDiag.emplace(S.PDiag(DiagID));
-    break;
-  }
-}
-
-// Print notes showing how we can reach FD starting from an a priori
-// known-callable function.
-static void EmitCallStackNotes(Sema &S, FunctionDecl *FD) {
-  auto FnIt = S.CUDAKnownEmittedFns.find(FD);
-  while (FnIt != S.CUDAKnownEmittedFns.end()) {
-    DiagnosticBuilder Builder(
-        S.Diags.Report(FnIt->second.Loc, diag::note_called_by));
-    Builder << FnIt->second.FD;
-    Builder.setForceEmit();
-
-    FnIt = S.CUDAKnownEmittedFns.find(FnIt->second.FD);
-  }
-}
-
-Sema::CUDADiagBuilder::~CUDADiagBuilder() {
-  if (ImmediateDiag) {
-    // Emit our diagnostic and, if it was a warning or error, output a callstack
-    // if Fn isn't a priori known-emitted.
-    bool IsWarningOrError = S.getDiagnostics().getDiagnosticLevel(
-                                DiagID, Loc) >= DiagnosticsEngine::Warning;
-    ImmediateDiag.reset(); // Emit the immediate diag.
-    if (IsWarningOrError && ShowCallStack)
-      EmitCallStackNotes(S, Fn);
-  } else if (PartialDiag) {
-    assert(ShowCallStack && "Must always show call stack for deferred diags.");
-    S.CUDADeferredDiags[Fn].push_back({Loc, std::move(*PartialDiag)});
-  }
-}
-
 // Do we know that we will eventually codegen the given function?
 static bool IsKnownEmitted(Sema &S, FunctionDecl *FD) {
   // Templates are emitted when they're instantiated.
   if (FD->isDependentContext())
     return false;
 
   // When compiling for device, host functions are never emitted.  Similarly,
   // when compiling for host, device and global functions are never emitted.
static bool IsKnownEmitted(Sema &S, FunctionDecl *FD) {
   FunctionDecl *Def = FD->getDefinition();
 
   if (Def &&
       !isDiscardableGVALinkage(S.getASTContext().GetGVALinkageForFunction(Def)))
     return true;
 
   // Otherwise, the function is known-emitted if it's in our set of
   // known-emitted functions.
-  return S.CUDAKnownEmittedFns.count(FD) > 0;
+  return S.DeviceKnownEmittedFns.count(FD) > 0;
 }
 
-Sema::CUDADiagBuilder Sema::CUDADiagIfDeviceCode(SourceLocation Loc,
+Sema::DeviceDiagBuilder Sema::CUDADiagIfDeviceCode(SourceLocation Loc,
                                                    unsigned DiagID) {
   assert(getLangOpts().CUDA && "Should only be called during CUDA compilation");
-  CUDADiagBuilder::Kind DiagKind = [&] {
+  DeviceDiagBuilder::Kind DiagKind = [this] {
     switch (CurrentCUDATarget()) {
     case CFT_Global:
     case CFT_Device:
-      return CUDADiagBuilder::K_Immediate;
+      return DeviceDiagBuilder::K_Immediate;
     case CFT_HostDevice:
       // An HD function counts as host code if we're compiling for host, and
       // device code if we're compiling for device.  Defer any errors in device
       // mode until the function is known-emitted.
       if (getLangOpts().CUDAIsDevice) {
         return IsKnownEmitted(*this, dyn_cast<FunctionDecl>(CurContext))
-                   ? CUDADiagBuilder::K_ImmediateWithCallStack
-                   : CUDADiagBuilder::K_Deferred;
+                   ? DeviceDiagBuilder::K_ImmediateWithCallStack
+                   : DeviceDiagBuilder::K_Deferred;
       }
-      return CUDADiagBuilder::K_Nop;
+      return DeviceDiagBuilder::K_Nop;
 
     default:
-      return CUDADiagBuilder::K_Nop;
+      return DeviceDiagBuilder::K_Nop;
     }
   }();
-  return CUDADiagBuilder(DiagKind, Loc, DiagID,
+  return DeviceDiagBuilder(DiagKind, Loc, DiagID,
                            dyn_cast<FunctionDecl>(CurContext), *this);
 }
 
-Sema::CUDADiagBuilder Sema::CUDADiagIfHostCode(SourceLocation Loc,
+Sema::DeviceDiagBuilder Sema::CUDADiagIfHostCode(SourceLocation Loc,
                                                  unsigned DiagID) {
   assert(getLangOpts().CUDA && "Should only be called during CUDA compilation");
-  CUDADiagBuilder::Kind DiagKind = [&] {
+  DeviceDiagBuilder::Kind DiagKind = [this] {
     switch (CurrentCUDATarget()) {
     case CFT_Host:
-      return CUDADiagBuilder::K_Immediate;
+      return DeviceDiagBuilder::K_Immediate;
     case CFT_HostDevice:
       // An HD function counts as host code if we're compiling for host, and
       // device code if we're compiling for device.  Defer any errors in device
       // mode until the function is known-emitted.
       if (getLangOpts().CUDAIsDevice)
-        return CUDADiagBuilder::K_Nop;
+        return DeviceDiagBuilder::K_Nop;
 
       return IsKnownEmitted(*this, dyn_cast<FunctionDecl>(CurContext))
-                 ? CUDADiagBuilder::K_ImmediateWithCallStack
-                 : CUDADiagBuilder::K_Deferred;
+                 ? DeviceDiagBuilder::K_ImmediateWithCallStack
+                 : DeviceDiagBuilder::K_Deferred;
     default:
-      return CUDADiagBuilder::K_Nop;
+      return DeviceDiagBuilder::K_Nop;
     }
   }();
-  return CUDADiagBuilder(DiagKind, Loc, DiagID,
+  return DeviceDiagBuilder(DiagKind, Loc, DiagID,
                            dyn_cast<FunctionDecl>(CurContext), *this);
 }
 
-// Emit any deferred diagnostics for FD and erase them from the map in which
-// they're stored.
-static void EmitDeferredDiags(Sema &S, FunctionDecl *FD) {
-  auto It = S.CUDADeferredDiags.find(FD);
-  if (It == S.CUDADeferredDiags.end())
-    return;
-  bool HasWarningOrError = false;
-  for (PartialDiagnosticAt &PDAt : It->second) {
-    const SourceLocation &Loc = PDAt.first;
-    const PartialDiagnostic &PD = PDAt.second;
-    HasWarningOrError |= S.getDiagnostics().getDiagnosticLevel(
-                             PD.getDiagID(), Loc) >= DiagnosticsEngine::Warning;
-    DiagnosticBuilder Builder(S.Diags.Report(Loc, PD.getDiagID()));
-    Builder.setForceEmit();
-    PD.Emit(Builder);
-  }
-  S.CUDADeferredDiags.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)
-    EmitCallStackNotes(S, FD);
-}
-
-// 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.
-static void MarkKnownEmitted(Sema &S, FunctionDecl *OrigCaller,
-                             FunctionDecl *OrigCallee, SourceLocation OrigLoc) {
-  // Nothing to do if we already know that FD is emitted.
-  if (IsKnownEmitted(S, OrigCallee)) {
-    assert(!S.CUDACallGraph.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.CUDAKnownEmittedFns[C.Callee] = {C.Caller, C.Loc};
-    EmitDeferredDiags(S, C.Callee);
-
-    // 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.CUDAKnownEmittedFns.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.CUDACallGraph.find(C.Callee);
-    if (CGIt == S.CUDACallGraph.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 CUDACallGraph.
-    S.CUDACallGraph.erase(CGIt);
-  }
-}
-
 bool Sema::CheckCUDACall(SourceLocation Loc, FunctionDecl *Callee) {
   assert(getLangOpts().CUDA && "Should only be called during CUDA compilation");
   assert(Callee && "Callee may not be null.");
   // FIXME: Is bailing out early correct here?  Should we instead assume that
   // the caller is a global initializer?
   FunctionDecl *Caller = dyn_cast<FunctionDecl>(CurContext);
   if (!Caller)
     return true;
 
   // If the caller is known-emitted, mark the callee as known-emitted.
   // Otherwise, mark the call in our call graph so we can traverse it later.
   bool CallerKnownEmitted = IsKnownEmitted(*this, Caller);
   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 (getLangOpts().CUDAIsDevice || IdentifyCUDATarget(Callee) != CFT_Global)
-      MarkKnownEmitted(*this, Caller, Callee, Loc);
+      markKnownEmitted(*this, Caller, Callee, Loc, IsKnownEmitted);
   } 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 (getLangOpts().CUDAIsDevice || IdentifyCUDATarget(Callee) != CFT_Global)
-      CUDACallGraph[Caller].insert({Callee, Loc});
+      DeviceCallGraph[Caller].insert({Callee, Loc});
   }
 
-  CUDADiagBuilder::Kind DiagKind = [&] {
+  DeviceDiagBuilder::Kind DiagKind = [this, Caller, Callee,
+                                      CallerKnownEmitted] {
     switch (IdentifyCUDAPreference(Caller, Callee)) {
     case CFP_Never:
-      return CUDADiagBuilder::K_Immediate;
+      return DeviceDiagBuilder::K_Immediate;
     case CFP_WrongSide:
       assert(Caller && "WrongSide calls require a non-null caller");
       // If we know the caller will be emitted, we know this wrong-side call
       // will be emitted, so it's an immediate error.  Otherwise, defer the
       // error until we know the caller is emitted.
-      return CallerKnownEmitted ? CUDADiagBuilder::K_ImmediateWithCallStack
-                                : CUDADiagBuilder::K_Deferred;
+      return CallerKnownEmitted ? DeviceDiagBuilder::K_ImmediateWithCallStack
+                                : DeviceDiagBuilder::K_Deferred;
     default:
-      return CUDADiagBuilder::K_Nop;
+      return DeviceDiagBuilder::K_Nop;
     }
   }();
 
-  if (DiagKind == CUDADiagBuilder::K_Nop)
+  if (DiagKind == DeviceDiagBuilder::K_Nop)
     return true;
 
   // Avoid emitting this error twice for the same location.  Using a hashtable
   // like this is unfortunate, but because we must continue parsing as normal
   // after encountering a deferred error, it's otherwise very tricky for us to
   // ensure that we only emit this deferred error once.
   if (!LocsWithCUDACallDiags.insert({Caller, Loc}).second)
     return true;
 
-  CUDADiagBuilder(DiagKind, Loc, diag::err_ref_bad_target, Caller, *this)
+  DeviceDiagBuilder(DiagKind, Loc, diag::err_ref_bad_target, Caller, *this)
       << IdentifyCUDATarget(Callee) << Callee << IdentifyCUDATarget(Caller);
-  CUDADiagBuilder(DiagKind, Callee->getLocation(), diag::note_previous_decl,
+  DeviceDiagBuilder(DiagKind, Callee->getLocation(), diag::note_previous_decl,
                     Caller, *this)
       << Callee;
-  return DiagKind != CUDADiagBuilder::K_Immediate &&
-         DiagKind != CUDADiagBuilder::K_ImmediateWithCallStack;
+  return DiagKind != DeviceDiagBuilder::K_Immediate &&
+         DiagKind != DeviceDiagBuilder::K_ImmediateWithCallStack;
 }
 
 void Sema::CUDASetLambdaAttrs(CXXMethodDecl *Method) {
   assert(getLangOpts().CUDA && "Should only be called during CUDA compilation");
   if (Method->hasAttr<CUDAHostAttr>() || Method->hasAttr<CUDADeviceAttr>())
     return;
   FunctionDecl *CurFn = dyn_cast<FunctionDecl>(CurContext);
   if (!CurFn)