[CUDA] When we emit an error that might have been deferred, also print a callstack.

[jlebar]

Summary

Previously, when you did something not allowed in a host+device function
and then caused it to be codegen'ed, we would print out an error telling
you that you did something bad, but we wouldn't tell you how we decided
that the function needed to be codegen'ed.

This change causes us to print out a callstack when emitting deferred
errors. This is immensely helpful when debugging highly-templated code,
where it's often unclear how a function became known-emitted.

We only print the callstack once per function, after we print the all
deferred errors.

This patch also switches all of our hashtables to using canonical
FunctionDecls instead of regular FunctionDecls. This prevents a number
of bugs, some of which are caught by tests added here, in which we
assume that two FDs for the same function have the same pointer value.

[rnk]

clang/include/clang/Sema/Sema.h

9292 ↗

(On Diff #74912)

Rather than having a custom key type, maybe it would be better to phrase this as a MapVector<CanonicalDeclPtr<FunctionDecl>, SourceLocation> ? Despite all the comments, I assumed FunctionDeclAndLoc was hashed by both elements for a long time.

9322 ↗

(On Diff #74912)

typo on "immediate"

[jlebar]

Address rnk's comments.

clang/include/clang/Sema/Sema.h

9292 ↗

(On Diff #74912)

That's much better; thank you.

[rnk]

lgtm

clang/include/clang/Basic/DiagnosticSemaKinds.td

6707 ↗

(On Diff #75178)

Do you think it's worth trying to indicate why the root of the "called by" notes must be emitted? I'm not suggesting we do it in this patch, just wondering.

clang/include/clang/Sema/Sema.h

9259 ↗

(On Diff #75178)

So, part of me wants to use FunctionDeclAndLoc here instead of std::pair, but then you'd have to bring back all the hashing machinery instead of getting it for free. Up to you, I guess.

9274 ↗

(On Diff #75178)

I guess there can be many callers, but we always record the first one that caused this function to be emitted.

[jlebar]

I'm going to submit this and send a patch to reuse FunctionDeclAndLoc. But I'm happy to add a comment about the note as well.

clang/include/clang/Basic/DiagnosticSemaKinds.td

6707 ↗

(On Diff #75178)

It seems just like e.g. note_template_class_instantiation_here, so I am not entirely sure what is the point of confusion. But if you can clarify that, I am happy to add a comment in a separate patch -- this stuff is already quite complicated, so I'm in favor of whatever we can do to make it clearer.

clang/include/clang/Sema/Sema.h

9259 ↗

(On Diff #75178)

Oh, and now make its hash function depend on both members instead of just the FD?

I actually like that change, but let me make it in a separate patch.

9274 ↗

(On Diff #75178)

Yes, exactly.

I was briefly worried about cycles here, but I think we're OK because ultimately we need to end up at an a priori known-emitted function, and that's a root in this tree.

[jlebar]

Thank you for the review, Reid.

cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td

 def err_nontemporal_builtin_must_be_pointer_intfltptr_or_vector : Error<
   "address argument to nontemporal builtin must be a pointer to integer, float, "
   "pointer, or a vector of such types (%0 invalid)">;
 
 def err_deleted_function_use : Error<"attempt to use a deleted function">;
 def err_deleted_inherited_ctor_use : Error<
   "constructor inherited by %0 from base class %1 is implicitly deleted">;
 
+def note_called_by : Note<"called by %0">;
 def err_kern_type_not_void_return : Error<
   "kernel function type %0 must have void return type">;
 def err_kern_is_nonstatic_method : Error<
   "kernel function %0 must be a free function or static member function">;
 def err_config_scalar_return : Error<
   "CUDA special function 'cudaConfigureCall' must have scalar return type">;
 def err_kern_call_not_global_function : Error<
   "kernel call to non-global function %0">;

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

public:
   /// Decrements our count of the number of times we've seen a pragma forcing
   /// functions to be __host__ __device__.  Returns false if the count is 0
   /// 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<const FunctionDecl *, std::vector<PartialDiagnosticAt>>
+  llvm::DenseMap<CanonicalDeclPtr<FunctionDecl>,
+                 std::vector<PartialDiagnosticAt>>
       CUDADeferredDiags;
 
   /// FunctionDecls plus raw encodings of 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<std::pair<FunctionDecl *, unsigned>> LocsWithCUDACallDiags;
+  llvm::DenseSet<std::pair<CanonicalDeclPtr<FunctionDecl>, unsigned>>
+      LocsWithCUDACallDiags;
 
-  /// The set of CUDA functions that we've discovered must be emitted by tracing
-  /// the call graph.  Functions that we can tell a priori must be emitted
-  /// aren't added to this set.
-  llvm::DenseSet<FunctionDecl *> CUDAKnownEmittedFns;
+  /// A pair of a canonical FunctionDecl and a SourceLocation.
+  struct FunctionDeclAndLoc {
+    CanonicalDeclPtr<FunctionDecl> FD;
+    SourceLocation Loc;
+  };
+
+  /// 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;
 
   /// A partial call graph maintained during CUDA compilation to support
-  /// deferred diagnostics.  Specifically, functions are only added here if, at
-  /// the time they're added, 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.
-  llvm::DenseMap<FunctionDecl *, llvm::SetVector<FunctionDecl *>> CUDACallGraph;
+  /// 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.
+  llvm::DenseMap</* Caller = */ CanonicalDeclPtr<FunctionDecl>,
+                 /* Callees = */ llvm::MapVector<CanonicalDeclPtr<FunctionDecl>,
+                                                 SourceLocation>>
+      CUDACallGraph;
 
   /// Diagnostic builder for CUDA 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 {
   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.
+      /// 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,
                     FunctionDecl *Fn, Sema &S);
+    ~CUDADiagBuilder();
 
     /// 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)
     ///     return ExprError();
     ///
     /// But see CUDADiagIfDeviceCode() and CUDADiagIfHostCode() -- you probably
     /// want to use these instead of creating a CUDADiagBuilder yourself.
-    operator bool() const { return ImmediateDiagBuilder.hasValue(); }
+    operator bool() const { return ImmediateDiag.hasValue(); }
 
     template <typename T>
     friend const CUDADiagBuilder &operator<<(const CUDADiagBuilder &Diag,
                                              const T &Value) {
-      if (Diag.ImmediateDiagBuilder.hasValue())
-        *Diag.ImmediateDiagBuilder << Value;
-      else if (Diag.PartialDiagInfo.hasValue())
-        Diag.PartialDiagInfo->PD << Value;
+      if (Diag.ImmediateDiag.hasValue())
+        *Diag.ImmediateDiag << Value;
+      else if (Diag.PartialDiag.hasValue())
+        *Diag.PartialDiag << Value;
       return Diag;
     }
 
   private:
-    struct PartialDiagnosticInfo {
-      PartialDiagnosticInfo(Sema &S, SourceLocation Loc, PartialDiagnostic PD,
-                            FunctionDecl *Fn)
-          : S(S), Loc(Loc), PD(std::move(PD)), Fn(Fn) {}
-
-      ~PartialDiagnosticInfo() {
-        S.CUDADeferredDiags[Fn].push_back({Loc, std::move(PD)});
-      }
-
-      Sema &S;
-      SourceLocation Loc;
-      PartialDiagnostic PD;
-      FunctionDecl *Fn;
-    };
+    unsigned DiagID;
+    bool ShowCallStack;
 
     // Invariant: At most one of these Optionals has a value.
     // FIXME: Switch these to a Variant once that exists.
-    llvm::Optional<Sema::SemaDiagnosticBuilder> ImmediateDiagBuilder;
-    llvm::Optional<PartialDiagnosticInfo> PartialDiagInfo;
+    llvm::Optional<SemaDiagnosticBuilder> ImmediateDiag;
+    llvm::Optional<PartialDiagnostic> PartialDiag;
   };
 
   /// Creates a CUDADiagBuilder 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.

cfe/trunk/lib/Sema/SemaCUDA.cpp

if (It != Previous.end()) {
     }
     return;
   }
 
   NewD->addAttr(CUDAHostAttr::CreateImplicit(Context));
   NewD->addAttr(CUDADeviceAttr::CreateImplicit(Context));
 }
 
-Sema::CUDADiagBuilder::CUDADiagBuilder(Kind K, SourceLocation Loc,
-                                       unsigned DiagID, FunctionDecl *Fn,
-                                       Sema &S) {
-  switch (K) {
-  case K_Nop:
-    break;
-  case K_Immediate:
-    ImmediateDiagBuilder.emplace(S.Diag(Loc, DiagID));
-    break;
-  case K_Deferred:
-    assert(Fn && "Must have a function to attach the deferred diag to.");
-    PartialDiagInfo.emplace(S, Loc, S.PDiag(DiagID), Fn);
-    break;
-  }
-}
-
 // 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.
CUDADiagBuilder::Kind DiagKind = [&] {
     case CFT_Device:
       return CUDADiagBuilder::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_Immediate
+                   ? CUDADiagBuilder::K_ImmediateWithCallStack
                    : CUDADiagBuilder::K_Deferred;
       }
       return CUDADiagBuilder::K_Nop;
 
     default:
       return CUDADiagBuilder::K_Nop;
     }
   }();
CUDADiagBuilder::Kind DiagKind = [&] {
     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 IsKnownEmitted(*this, dyn_cast<FunctionDecl>(CurContext))
-                 ? CUDADiagBuilder::K_Immediate
+                 ? CUDADiagBuilder::K_ImmediateWithCallStack
                  : CUDADiagBuilder::K_Deferred;
     default:
       return CUDADiagBuilder::K_Nop;
     }
   }();
   return CUDADiagBuilder(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 *FD) {
+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, FD)) {
-    assert(!S.CUDACallGraph.count(FD));
+  if (IsKnownEmitted(S, OrigCallee)) {
+    assert(!S.CUDACallGraph.count(OrigCallee));
     return;
   }
 
-  // We've just discovered that FD is known-emitted.  Walk our call graph to see
-  // what else we can now discover also must be emitted.
-  llvm::SmallVector<FunctionDecl *, 4> Worklist = {FD};
-  llvm::SmallSet<FunctionDecl *, 4> Seen;
-  Seen.insert(FD);
+  // 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()) {
-    FunctionDecl *Caller = Worklist.pop_back_val();
-    assert(!IsKnownEmitted(S, Caller) &&
+    CallInfo C = Worklist.pop_back_val();
+    assert(!IsKnownEmitted(S, C.Callee) &&
            "Worklist should not contain known-emitted functions.");
-    S.CUDAKnownEmittedFns.insert(Caller);
-    EmitDeferredDiags(S, Caller);
+    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 = Caller->getPrimaryTemplate()) {
+    if (auto *Templ = C.Callee->getPrimaryTemplate()) {
       FunctionDecl *TemplFD = Templ->getAsFunction();
       if (!Seen.count(TemplFD) && !S.CUDAKnownEmittedFns.count(TemplFD)) {
         Seen.insert(TemplFD);
-        Worklist.push_back(TemplFD);
+        Worklist.push_back(
+            {/* Caller = */ C.Caller, /* Callee = */ TemplFD, C.Loc});
       }
     }
 
-    // Add all functions called by Caller to our worklist.
-    auto CGIt = S.CUDACallGraph.find(Caller);
+    // Add all functions called by Callee to our worklist.
+    auto CGIt = S.CUDACallGraph.find(C.Callee);
     if (CGIt == S.CUDACallGraph.end())
       continue;
 
-    for (FunctionDecl *Callee : CGIt->second) {
-      if (Seen.count(Callee) || IsKnownEmitted(S, Callee))
+    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(Callee);
-      Worklist.push_back(Callee);
+      Seen.insert(NewCallee);
+      Worklist.push_back(
+          {/* Caller = */ C.Callee, /* Callee = */ NewCallee, CallLoc});
     }
 
-    // Caller is now known-emitted, so we no longer need to maintain its list of
-    // callees in CUDACallGraph.
+    // 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)
-    MarkKnownEmitted(*this, Callee);
+    MarkKnownEmitted(*this, Caller, Callee, Loc);
   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);
+      CUDACallGraph[Caller].insert({Callee, Loc});
   }
 
   CUDADiagBuilder::Kind DiagKind = [&] {
     switch (IdentifyCUDAPreference(Caller, Callee)) {
     case CFP_Never:
       return CUDADiagBuilder::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_Immediate
+      return CallerKnownEmitted ? CUDADiagBuilder::K_ImmediateWithCallStack
                                 : CUDADiagBuilder::K_Deferred;
     default:
       return CUDADiagBuilder::K_Nop;
     }
   }();
 
   if (DiagKind == CUDADiagBuilder::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.getRawEncoding()}).second)
     return true;
 
   CUDADiagBuilder(DiagKind, Loc, diag::err_ref_bad_target, Caller, *this)
       << IdentifyCUDATarget(Callee) << Callee << IdentifyCUDATarget(Caller);
   CUDADiagBuilder(DiagKind, Callee->getLocation(), diag::note_previous_decl,
                   Caller, *this)
       << Callee;
-  return DiagKind != CUDADiagBuilder::K_Immediate;
+  return DiagKind != CUDADiagBuilder::K_Immediate &&
+         DiagKind != CUDADiagBuilder::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)

cfe/trunk/test/SemaCUDA/bad-calls-on-same-line.cu

inline __host__ __device__ void hd() {
   Selector<T>().f();
   // expected-error@-1 {{reference to __device__ function}}
   // expected-error@-2 {{reference to __device__ function}}
 }
 
 void host_fn() {
   hd<int>();
   hd<double>();  // expected-note {{function template specialization 'hd<double>'}}
+  // expected-note@-1 {{called by 'host_fn'}}
   hd<float>();  // expected-note {{function template specialization 'hd<float>'}}
+  // expected-note@-1 {{called by 'host_fn'}}
 }

cfe/trunk/test/SemaCUDA/call-device-fn-from-host.cu

-// RUN: %clang_cc1 %s --std=c++11 -triple x86_64-unknown-linux -emit-llvm -o - -verify
+// RUN: %clang_cc1 %s --std=c++11 -triple x86_64-unknown-linux -emit-llvm -o - \
+// RUN:   -verify -verify-ignore-unexpected=note
 
 // Note: This test won't work with -fsyntax-only, because some of these errors
 // are emitted during codegen.
 
 #include "Inputs/cuda.h"
 
 __device__ void device_fn() {}
 // expected-note@-1 {{'device_fn' declared here}}

cfe/trunk/test/SemaCUDA/call-host-fn-from-device.cu

 // RUN: %clang_cc1 %s --std=c++11 -triple nvptx-unknown-unknown -fcuda-is-device \
-// RUN:   -emit-llvm -o /dev/null -verify
+// RUN:   -emit-llvm -o /dev/null -verify -verify-ignore-unexpected=note
 
 // Note: This test won't work with -fsyntax-only, because some of these errors
 // are emitted during codegen.
 
 #include "Inputs/cuda.h"
 
 extern "C" void host_fn() {}
 // expected-note@-1 {{'host_fn' declared here}}

cfe/trunk/test/SemaCUDA/call-stack-for-deferred-err.cu

+// RUN: %clang_cc1 -fcuda-is-device -fsyntax-only -verify %s
+
+#include "Inputs/cuda.h"
+
+// We should emit an error for hd_fn's use of a VLA.  This would have been
+// legal if hd_fn were never codegen'ed on the device, so we should also print
+// out a callstack showing how we determine that hd_fn is known-emitted.
+//
+// Compare to no-call-stack-for-deferred-err.cu.
+
+inline __host__ __device__ void hd_fn(int n);
+inline __device__ void device_fn2() { hd_fn(42); } // expected-note {{called by 'device_fn2'}}
+
+__global__ void kernel() { device_fn2(); } // expected-note {{called by 'kernel'}}
+
+inline __host__ __device__ void hd_fn(int n) {
+  int vla[n]; // expected-error {{variable-length array}}
+}

cfe/trunk/test/SemaCUDA/exceptions.cu

inline __host__ __device__ void hd3() {
   try {} catch(void*) {}
 #ifdef __CUDA_ARCH__
   // expected-error@-3 {{cannot use 'throw' in __host__ __device__ function}}
   // expected-error@-3 {{cannot use 'try' in __host__ __device__ function}}
 #endif
 }
 
 __device__ void call_hd3() { hd3(); }
+#ifdef __CUDA_ARCH__
+// expected-note@-2 {{called by 'call_hd3'}}
+#endif

cfe/trunk/test/SemaCUDA/no-call-stack-for-immediate-errs.cu

+// RUN: %clang_cc1 -fcuda-is-device -fsyntax-only -verify %s
+
+#include "Inputs/cuda.h"
+
+// Here we should dump an error about the VLA in device_fn, but we should not
+// print a callstack indicating how device_fn becomes known-emitted, because
+// it's an error to use a VLA in any __device__ function, even one that doesn't
+// get emitted.
+
+inline __device__ void device_fn(int n);
+inline __device__ void device_fn2() { device_fn(42); }
+
+__global__ void kernel() { device_fn2(); }
+
+inline __device__ void device_fn(int n) {
+  int vla[n]; // expected-error {{variable-length array}}
+}

cfe/trunk/test/SemaCUDA/trace-through-global.cu

inline __host__ __device__ void hd4(double) {
   device_fn(0);  // No error; this function is never called.
 }
 
 __global__ void kernel(int) { hd2(); }
 
 template <typename T>
 void launch_kernel() {
   kernel<<<0, 0>>>(T());
-  hd1();
-  hd3(T());
+
+  // Notice that these two diagnostics are different: Because the call to hd1
+  // is not dependent on T, the call to hd1 comes from 'launch_kernel', while
+  // the call to hd3, being dependent, comes from 'launch_kernel<int>'.
+  hd1(); // expected-note {{called by 'launch_kernel'}}
+  hd3(T()); // expected-note {{called by 'launch_kernel<int>'}}
 }
 
 void host_fn() {
   launch_kernel<int>();
+  // expected-note@-1 {{called by 'host_fn'}}
+  // expected-note@-2 {{called by 'host_fn'}}
 }