[clang] 3453b69 - Revert "Recommit "[CUDA][HIP] Defer overloading resolution diagnostics for host device functions""

Reid Kleckner via cfe-commits cfe-commits at lists.llvm.org
Thu Sep 24 11:17:03 PDT 2020


Author: Reid Kleckner
Date: 2020-09-24T11:16:54-07:00
New Revision: 3453b6928da332bb67f902add71f5cd80f61c136

URL: https://github.com/llvm/llvm-project/commit/3453b6928da332bb67f902add71f5cd80f61c136
DIFF: https://github.com/llvm/llvm-project/commit/3453b6928da332bb67f902add71f5cd80f61c136.diff

LOG: Revert "Recommit "[CUDA][HIP] Defer overloading resolution diagnostics for host device functions""

This reverts commit e39da8ab6a286ac777d5fe7799f1eb782cf99938.

This depends on a change that needs additional design review and needs
to be reverted.

Added: 
    

Modified: 
    clang-tools-extra/clangd/Diagnostics.cpp
    clang/include/clang/Basic/Diagnostic.td
    clang/include/clang/Basic/DiagnosticAST.h
    clang/include/clang/Basic/DiagnosticAnalysis.h
    clang/include/clang/Basic/DiagnosticComment.h
    clang/include/clang/Basic/DiagnosticCrossTU.h
    clang/include/clang/Basic/DiagnosticDriver.h
    clang/include/clang/Basic/DiagnosticFrontend.h
    clang/include/clang/Basic/DiagnosticIDs.h
    clang/include/clang/Basic/DiagnosticLex.h
    clang/include/clang/Basic/DiagnosticParse.h
    clang/include/clang/Basic/DiagnosticRefactoring.h
    clang/include/clang/Basic/DiagnosticSema.h
    clang/include/clang/Basic/DiagnosticSemaKinds.td
    clang/include/clang/Basic/DiagnosticSerialization.h
    clang/include/clang/Basic/LangOptions.def
    clang/include/clang/Driver/Options.td
    clang/include/clang/Sema/Sema.h
    clang/lib/Basic/DiagnosticIDs.cpp
    clang/lib/Driver/ToolChains/Cuda.cpp
    clang/lib/Driver/ToolChains/HIP.cpp
    clang/lib/Frontend/CompilerInvocation.cpp
    clang/lib/Sema/AnalysisBasedWarnings.cpp
    clang/lib/Sema/Sema.cpp
    clang/lib/Sema/SemaAttr.cpp
    clang/lib/Sema/SemaCUDA.cpp
    clang/lib/Sema/SemaDecl.cpp
    clang/lib/Sema/SemaExprObjC.cpp
    clang/lib/Sema/SemaOpenMP.cpp
    clang/lib/Sema/SemaOverload.cpp
    clang/lib/Sema/SemaSYCL.cpp
    clang/lib/Sema/SemaStmt.cpp
    clang/lib/Sema/SemaStmtAsm.cpp
    clang/lib/Sema/SemaTemplateInstantiate.cpp
    clang/lib/Sema/SemaTemplateInstantiateDecl.cpp
    clang/lib/Sema/SemaTemplateVariadic.cpp
    clang/lib/Sema/SemaType.cpp
    clang/test/TableGen/DiagnosticBase.inc
    clang/tools/diagtool/DiagnosticNames.cpp
    clang/utils/TableGen/ClangDiagnosticsEmitter.cpp

Removed: 
    clang/test/SemaCUDA/deferred-oeverload.cu
    clang/test/TableGen/deferred-diag.td


################################################################################
diff  --git a/clang-tools-extra/clangd/Diagnostics.cpp b/clang-tools-extra/clangd/Diagnostics.cpp
index 18ff96202e0a..afa72f9d4051 100644
--- a/clang-tools-extra/clangd/Diagnostics.cpp
+++ b/clang-tools-extra/clangd/Diagnostics.cpp
@@ -43,7 +43,7 @@ namespace {
 const char *getDiagnosticCode(unsigned ID) {
   switch (ID) {
 #define DIAG(ENUM, CLASS, DEFAULT_MAPPING, DESC, GROPU, SFINAE, NOWERROR,      \
-             SHOWINSYSHEADER, DEFERRABLE, CATEGORY)                            \
+             SHOWINSYSHEADER, CATEGORY)                                        \
   case clang::diag::ENUM:                                                      \
     return #ENUM;
 #include "clang/Basic/DiagnosticASTKinds.inc"

diff  --git a/clang/include/clang/Basic/Diagnostic.td b/clang/include/clang/Basic/Diagnostic.td
index ab2c738a2ace..48ba8c0f469f 100644
--- a/clang/include/clang/Basic/Diagnostic.td
+++ b/clang/include/clang/Basic/Diagnostic.td
@@ -45,7 +45,6 @@ class TextSubstitution<string Text> {
   // diagnostics
   string Component = "";
   string CategoryName = "";
-  bit Deferrable = 0;
 }
 
 // Diagnostic Categories.  These can be applied to groups or individual
@@ -84,7 +83,6 @@ class Diagnostic<string text, DiagClass DC, Severity defaultmapping> {
   bit            AccessControl = 0;
   bit            WarningNoWerror = 0;
   bit            ShowInSystemHeader = 0;
-  bit            Deferrable = 0;
   Severity       DefaultSeverity = defaultmapping;
   DiagGroup      Group;
   string         CategoryName = "";
@@ -108,14 +106,6 @@ class SuppressInSystemHeader {
   bit ShowInSystemHeader = 0;
 }
 
-class Deferrable {
-  bit Deferrable = 1;
-}
-
-class NonDeferrable {
-  bit Deferrable = 0;
-}
-
 // FIXME: ExtWarn and Extension should also be SFINAEFailure by default.
 class Error<string str>     : Diagnostic<str, CLASS_ERROR, SEV_Error>, SFINAEFailure {
   bit ShowInSystemHeader = 1;

diff  --git a/clang/include/clang/Basic/DiagnosticAST.h b/clang/include/clang/Basic/DiagnosticAST.h
index 76c31ad9508e..afe5f62e2012 100644
--- a/clang/include/clang/Basic/DiagnosticAST.h
+++ b/clang/include/clang/Basic/DiagnosticAST.h
@@ -15,7 +15,7 @@ namespace clang {
 namespace diag {
 enum {
 #define DIAG(ENUM, FLAGS, DEFAULT_MAPPING, DESC, GROUP, SFINAE, NOWERROR,      \
-             SHOWINSYSHEADER, DEFERRABLE, CATEGORY)                            \
+             SHOWINSYSHEADER, CATEGORY)                                        \
   ENUM,
 #define ASTSTART
 #include "clang/Basic/DiagnosticASTKinds.inc"

diff  --git a/clang/include/clang/Basic/DiagnosticAnalysis.h b/clang/include/clang/Basic/DiagnosticAnalysis.h
index f9037cc8d75a..eea35a4d616e 100644
--- a/clang/include/clang/Basic/DiagnosticAnalysis.h
+++ b/clang/include/clang/Basic/DiagnosticAnalysis.h
@@ -15,7 +15,7 @@ namespace clang {
 namespace diag {
 enum {
 #define DIAG(ENUM, FLAGS, DEFAULT_MAPPING, DESC, GROUP, SFINAE, NOWERROR,      \
-             SHOWINSYSHEADER, DEFERRABLE, CATEGORY)                            \
+             SHOWINSYSHEADER, CATEGORY)                                        \
   ENUM,
 #define ANALYSISSTART
 #include "clang/Basic/DiagnosticAnalysisKinds.inc"

diff  --git a/clang/include/clang/Basic/DiagnosticComment.h b/clang/include/clang/Basic/DiagnosticComment.h
index 6e011bfcebab..a87bafa8b3a5 100644
--- a/clang/include/clang/Basic/DiagnosticComment.h
+++ b/clang/include/clang/Basic/DiagnosticComment.h
@@ -15,7 +15,7 @@ namespace clang {
 namespace diag {
 enum {
 #define DIAG(ENUM, FLAGS, DEFAULT_MAPPING, DESC, GROUP, SFINAE, NOWERROR,      \
-             SHOWINSYSHEADER, DEFERRABLE, CATEGORY)                            \
+             SHOWINSYSHEADER, CATEGORY)                                        \
   ENUM,
 #define COMMENTSTART
 #include "clang/Basic/DiagnosticCommentKinds.inc"

diff  --git a/clang/include/clang/Basic/DiagnosticCrossTU.h b/clang/include/clang/Basic/DiagnosticCrossTU.h
index ded85ec3f840..c1c582bd6ee4 100644
--- a/clang/include/clang/Basic/DiagnosticCrossTU.h
+++ b/clang/include/clang/Basic/DiagnosticCrossTU.h
@@ -15,7 +15,7 @@ namespace clang {
 namespace diag {
 enum {
 #define DIAG(ENUM, FLAGS, DEFAULT_MAPPING, DESC, GROUP, SFINAE, NOWERROR,      \
-             SHOWINSYSHEADER, DEFERRABLE, CATEGORY)                            \
+             SHOWINSYSHEADER, CATEGORY)                                        \
   ENUM,
 #define CROSSTUSTART
 #include "clang/Basic/DiagnosticCrossTUKinds.inc"

diff  --git a/clang/include/clang/Basic/DiagnosticDriver.h b/clang/include/clang/Basic/DiagnosticDriver.h
index cecd8fd6b4d5..63913df4523b 100644
--- a/clang/include/clang/Basic/DiagnosticDriver.h
+++ b/clang/include/clang/Basic/DiagnosticDriver.h
@@ -15,7 +15,7 @@ namespace clang {
 namespace diag {
 enum {
 #define DIAG(ENUM, FLAGS, DEFAULT_MAPPING, DESC, GROUP, SFINAE, NOWERROR,      \
-             SHOWINSYSHEADER, DEFERRABLE, CATEGORY)                            \
+             SHOWINSYSHEADER, CATEGORY)                                        \
   ENUM,
 #define DRIVERSTART
 #include "clang/Basic/DiagnosticDriverKinds.inc"

diff  --git a/clang/include/clang/Basic/DiagnosticFrontend.h b/clang/include/clang/Basic/DiagnosticFrontend.h
index f57c587fb469..57f00e73abb4 100644
--- a/clang/include/clang/Basic/DiagnosticFrontend.h
+++ b/clang/include/clang/Basic/DiagnosticFrontend.h
@@ -15,7 +15,7 @@ namespace clang {
 namespace diag {
 enum {
 #define DIAG(ENUM, FLAGS, DEFAULT_MAPPING, DESC, GROUP, SFINAE, NOWERROR,      \
-             SHOWINSYSHEADER, DEFERRABLE, CATEGORY)                            \
+             SHOWINSYSHEADER, CATEGORY)                                        \
   ENUM,
 #define FRONTENDSTART
 #include "clang/Basic/DiagnosticFrontendKinds.inc"

diff  --git a/clang/include/clang/Basic/DiagnosticIDs.h b/clang/include/clang/Basic/DiagnosticIDs.h
index 7fd107c4add7..00c939650e54 100644
--- a/clang/include/clang/Basic/DiagnosticIDs.h
+++ b/clang/include/clang/Basic/DiagnosticIDs.h
@@ -64,9 +64,8 @@ namespace clang {
 
     // Get typedefs for common diagnostics.
     enum {
-#define DIAG(ENUM, FLAGS, DEFAULT_MAPPING, DESC, GROUP, SFINAE, CATEGORY,      \
-             NOWERROR, SHOWINSYSHEADER, DEFFERABLE)                            \
-  ENUM,
+#define DIAG(ENUM,FLAGS,DEFAULT_MAPPING,DESC,GROUP,\
+             SFINAE,CATEGORY,NOWERROR,SHOWINSYSHEADER) ENUM,
 #define COMMONSTART
 #include "clang/Basic/DiagnosticCommonKinds.inc"
       NUM_BUILTIN_COMMON_DIAGNOSTICS
@@ -281,13 +280,6 @@ class DiagnosticIDs : public RefCountedBase<DiagnosticIDs> {
   /// are not SFINAE errors.
   static SFINAEResponse getDiagnosticSFINAEResponse(unsigned DiagID);
 
-  /// Whether the diagnostic message can be deferred.
-  ///
-  /// For single source offloading languages, a diagnostic message occurred
-  /// in a device host function may be deferred until the function is sure
-  /// to be emitted.
-  static bool isDeferrable(unsigned DiagID);
-
   /// Get the string of all diagnostic flags.
   ///
   /// \returns A list of all diagnostics flags as they would be written in a

diff  --git a/clang/include/clang/Basic/DiagnosticLex.h b/clang/include/clang/Basic/DiagnosticLex.h
index 7a3128de3b82..33789051b286 100644
--- a/clang/include/clang/Basic/DiagnosticLex.h
+++ b/clang/include/clang/Basic/DiagnosticLex.h
@@ -15,7 +15,7 @@ namespace clang {
 namespace diag {
 enum {
 #define DIAG(ENUM, FLAGS, DEFAULT_MAPPING, DESC, GROUP, SFINAE, NOWERROR,      \
-             SHOWINSYSHEADER, DEFERRABLE, CATEGORY)                            \
+             SHOWINSYSHEADER, CATEGORY)                                        \
   ENUM,
 #define LEXSTART
 #include "clang/Basic/DiagnosticLexKinds.inc"

diff  --git a/clang/include/clang/Basic/DiagnosticParse.h b/clang/include/clang/Basic/DiagnosticParse.h
index d066d3f71a25..0c21ff93c5fa 100644
--- a/clang/include/clang/Basic/DiagnosticParse.h
+++ b/clang/include/clang/Basic/DiagnosticParse.h
@@ -15,7 +15,7 @@ namespace clang {
 namespace diag {
 enum {
 #define DIAG(ENUM, FLAGS, DEFAULT_MAPPING, DESC, GROUP, SFINAE, NOWERROR,      \
-             SHOWINSYSHEADER, DEFERRABLE, CATEGORY)                            \
+             SHOWINSYSHEADER, CATEGORY)                                        \
   ENUM,
 #define PARSESTART
 #include "clang/Basic/DiagnosticParseKinds.inc"

diff  --git a/clang/include/clang/Basic/DiagnosticRefactoring.h b/clang/include/clang/Basic/DiagnosticRefactoring.h
index fc7564047a24..aded0162ab33 100644
--- a/clang/include/clang/Basic/DiagnosticRefactoring.h
+++ b/clang/include/clang/Basic/DiagnosticRefactoring.h
@@ -15,7 +15,7 @@ namespace clang {
 namespace diag {
 enum {
 #define DIAG(ENUM, FLAGS, DEFAULT_MAPPING, DESC, GROUP, SFINAE, NOWERROR,      \
-             SHOWINSYSHEADER, DEFERRABLE, CATEGORY)                            \
+             SHOWINSYSHEADER, CATEGORY)                                        \
   ENUM,
 #define REFACTORINGSTART
 #include "clang/Basic/DiagnosticRefactoringKinds.inc"

diff  --git a/clang/include/clang/Basic/DiagnosticSema.h b/clang/include/clang/Basic/DiagnosticSema.h
index 7323167aeee8..72a6b9753893 100644
--- a/clang/include/clang/Basic/DiagnosticSema.h
+++ b/clang/include/clang/Basic/DiagnosticSema.h
@@ -15,7 +15,7 @@ namespace clang {
 namespace diag {
 enum {
 #define DIAG(ENUM, FLAGS, DEFAULT_MAPPING, DESC, GROUP, SFINAE, NOWERROR,      \
-             SHOWINSYSHEADER, DEFERRABLE, CATEGORY)                            \
+             SHOWINSYSHEADER, CATEGORY)                                        \
   ENUM,
 #define SEMASTART
 #include "clang/Basic/DiagnosticSemaKinds.inc"

diff  --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td
index df61aacca64c..0cb817df9db3 100644
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -4100,8 +4100,6 @@ def err_ovl_static_nonstatic_member : Error<
   "static and non-static member functions with the same parameter types "
   "cannot be overloaded">;
 
-let Deferrable = 1 in {
-
 def err_ovl_no_viable_function_in_call : Error<
   "no matching function for call to %0">;
 def err_ovl_no_viable_member_function_in_call : Error<
@@ -4415,8 +4413,6 @@ def err_addr_ovl_not_func_ptrref : Error<
 def err_addr_ovl_no_qualifier : Error<
   "cannot form member pointer of type %0 without '&' and class name">;
 
-} // let Deferrable
-
 // C++11 Literal Operators
 def err_ovl_no_viable_literal_operator : Error<
   "no matching literal operator for call to %0"

diff  --git a/clang/include/clang/Basic/DiagnosticSerialization.h b/clang/include/clang/Basic/DiagnosticSerialization.h
index b3d99fb3feaa..7e46a36a7fd3 100644
--- a/clang/include/clang/Basic/DiagnosticSerialization.h
+++ b/clang/include/clang/Basic/DiagnosticSerialization.h
@@ -15,7 +15,7 @@ namespace clang {
 namespace diag {
 enum {
 #define DIAG(ENUM, FLAGS, DEFAULT_MAPPING, DESC, GROUP, SFINAE, NOWERROR,      \
-             SHOWINSYSHEADER, DEFERRABLE, CATEGORY)                            \
+             SHOWINSYSHEADER, CATEGORY)                                        \
   ENUM,
 #define SERIALIZATIONSTART
 #include "clang/Basic/DiagnosticSerializationKinds.inc"

diff  --git a/clang/include/clang/Basic/LangOptions.def b/clang/include/clang/Basic/LangOptions.def
index d711d66784a4..9846809763f8 100644
--- a/clang/include/clang/Basic/LangOptions.def
+++ b/clang/include/clang/Basic/LangOptions.def
@@ -241,7 +241,6 @@ LANGOPT(CUDADeviceApproxTranscendentals, 1, 0, "using approximate transcendental
 LANGOPT(GPURelocatableDeviceCode, 1, 0, "generate relocatable device code")
 LANGOPT(GPUAllowDeviceInit, 1, 0, "allowing device side global init functions for HIP")
 LANGOPT(GPUMaxThreadsPerBlock, 32, 256, "default max threads per block for kernel launch bounds for HIP")
-LANGOPT(GPUDeferDiag, 1, 0, "defer host/device related diagnostic messages for CUDA/HIP")
 
 LANGOPT(SYCL              , 1, 0, "SYCL")
 LANGOPT(SYCLIsDevice      , 1, 0, "Generate code for SYCL device")

diff  --git a/clang/include/clang/Driver/Options.td b/clang/include/clang/Driver/Options.td
index f7261babd16a..d7c2496b8a5d 100644
--- a/clang/include/clang/Driver/Options.td
+++ b/clang/include/clang/Driver/Options.td
@@ -669,9 +669,6 @@ defm hip_new_launch_api : OptInFFlag<"hip-new-launch-api",
   "Use", "Don't use", " new kernel launching API for HIP">;
 defm gpu_allow_device_init : OptInFFlag<"gpu-allow-device-init",
   "Allow", "Don't allow", " device side init function in HIP">;
-defm gpu_defer_diag : OptInFFlag<"gpu-defer-diag",
-  "Defer", "Don't defer", " host/device related diagnostic messages"
-  " for CUDA/HIP">;
 def gpu_max_threads_per_block_EQ : Joined<["--"], "gpu-max-threads-per-block=">,
   Flags<[CC1Option]>,
   HelpText<"Default max threads per block for kernel launch bounds for HIP">;

diff  --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h
index d6b9b0c6166f..9fa201df1802 100644
--- a/clang/include/clang/Sema/Sema.h
+++ b/clang/include/clang/Sema/Sema.h
@@ -1462,30 +1462,28 @@ class Sema final {
   /// template instantiation stacks.
   ///
   /// This class provides a wrapper around the basic DiagnosticBuilder
-  /// class that emits diagnostics. ImmediateDiagBuilder is
+  /// class that emits diagnostics. SemaDiagnosticBuilder is
   /// responsible for emitting the diagnostic (as DiagnosticBuilder
   /// does) and, if the diagnostic comes from inside a template
   /// instantiation, printing the template instantiation stack as
   /// well.
-  class ImmediateDiagBuilder : public DiagnosticBuilder {
+  class SemaDiagnosticBuilder : public DiagnosticBuilder {
     Sema &SemaRef;
     unsigned DiagID;
 
   public:
-    ImmediateDiagBuilder(DiagnosticBuilder &DB, Sema &SemaRef, unsigned DiagID)
-        : DiagnosticBuilder(DB), SemaRef(SemaRef), DiagID(DiagID) {}
-    ImmediateDiagBuilder(DiagnosticBuilder &&DB, Sema &SemaRef, unsigned DiagID)
-        : DiagnosticBuilder(DB), SemaRef(SemaRef), DiagID(DiagID) {}
+    SemaDiagnosticBuilder(DiagnosticBuilder &DB, Sema &SemaRef, unsigned DiagID)
+      : DiagnosticBuilder(DB), SemaRef(SemaRef), DiagID(DiagID) { }
 
     // This is a cunning lie. DiagnosticBuilder actually performs move
     // construction in its copy constructor (but due to varied uses, it's not
     // possible to conveniently express this as actual move construction). So
     // the default copy ctor here is fine, because the base class disables the
-    // source anyway, so the user-defined ~ImmediateDiagBuilder is a safe no-op
+    // source anyway, so the user-defined ~SemaDiagnosticBuilder is a safe no-op
     // in that case anwyay.
-    ImmediateDiagBuilder(const ImmediateDiagBuilder &) = default;
+    SemaDiagnosticBuilder(const SemaDiagnosticBuilder&) = default;
 
-    ~ImmediateDiagBuilder() {
+    ~SemaDiagnosticBuilder() {
       // If we aren't active, there is nothing to do.
       if (!isActive()) return;
 
@@ -1506,162 +1504,38 @@ class Sema final {
     }
 
     /// Teach operator<< to produce an object of the correct type.
-    template <typename T>
-    friend const ImmediateDiagBuilder &
-    operator<<(const ImmediateDiagBuilder &Diag, const T &Value) {
+    template<typename T>
+    friend const SemaDiagnosticBuilder &operator<<(
+        const SemaDiagnosticBuilder &Diag, const T &Value) {
       const DiagnosticBuilder &BaseDiag = Diag;
       BaseDiag << Value;
       return Diag;
     }
 
-    // It is necessary to limit this to rvalue reference to avoid calling this
-    // function with a bitfield lvalue argument since non-const reference to
-    // bitfield is not allowed.
-    template <typename T, typename = typename std::enable_if<
-                              !std::is_lvalue_reference<T>::value>::type>
-    const ImmediateDiagBuilder &operator<<(T &&V) const {
-      const DiagnosticBuilder &BaseDiag = *this;
-      BaseDiag << std::move(V);
-      return *this;
-    }
-  };
-
-  /// A generic diagnostic builder for 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 SemaDiagnosticBuilder {
-  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
-    };
-
-    SemaDiagnosticBuilder(Kind K, SourceLocation Loc, unsigned DiagID,
-                          FunctionDecl *Fn, Sema &S);
-    SemaDiagnosticBuilder(SemaDiagnosticBuilder &&D);
-    SemaDiagnosticBuilder(const SemaDiagnosticBuilder &) = default;
-    ~SemaDiagnosticBuilder();
-
-    bool isImmediate() const { return ImmediateDiag.hasValue(); }
-
-    /// 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 (SemaDiagnosticBuilder(...) << foo << bar)
-    ///     return ExprError();
-    ///
-    /// But see CUDADiagIfDeviceCode() and CUDADiagIfHostCode() -- you probably
-    /// want to use these instead of creating a SemaDiagnosticBuilder yourself.
-    operator bool() const { return isImmediate(); }
-
-    template <typename T>
-    friend const SemaDiagnosticBuilder &
-    operator<<(const SemaDiagnosticBuilder &Diag, const T &Value) {
-      if (Diag.ImmediateDiag.hasValue())
-        *Diag.ImmediateDiag << Value;
-      else if (Diag.PartialDiagId.hasValue())
-        Diag.S.DeviceDeferredDiags[Diag.Fn][*Diag.PartialDiagId].second
-            << Value;
-      return Diag;
-    }
-
     // It is necessary to limit this to rvalue reference to avoid calling this
     // function with a bitfield lvalue argument since non-const reference to
     // bitfield is not allowed.
     template <typename T, typename = typename std::enable_if<
                               !std::is_lvalue_reference<T>::value>::type>
     const SemaDiagnosticBuilder &operator<<(T &&V) const {
-      if (ImmediateDiag.hasValue())
-        *ImmediateDiag << std::move(V);
-      else if (PartialDiagId.hasValue())
-        S.DeviceDeferredDiags[Fn][*PartialDiagId].second << std::move(V);
+      const StreamableDiagnosticBase &DB = *this;
+      DB << std::move(V);
       return *this;
     }
-
-    friend const SemaDiagnosticBuilder &
-    operator<<(const SemaDiagnosticBuilder &Diag, const PartialDiagnostic &PD) {
-      if (Diag.ImmediateDiag.hasValue())
-        PD.Emit(*Diag.ImmediateDiag);
-      else if (Diag.PartialDiagId.hasValue())
-        Diag.S.DeviceDeferredDiags[Diag.Fn][*Diag.PartialDiagId].second = PD;
-      return Diag;
-    }
-
-    void AddFixItHint(const FixItHint &Hint) const {
-      if (ImmediateDiag.hasValue())
-        ImmediateDiag->AddFixItHint(Hint);
-      else if (PartialDiagId.hasValue())
-        S.DeviceDeferredDiags[Fn][*PartialDiagId].second.AddFixItHint(Hint);
-    }
-
-    friend ExprResult ExprError(const SemaDiagnosticBuilder &) {
-      return ExprError();
-    }
-    friend StmtResult StmtError(const SemaDiagnosticBuilder &) {
-      return StmtError();
-    }
-    operator ExprResult() const { return ExprError(); }
-    operator StmtResult() const { return StmtError(); }
-    operator TypeResult() const { return TypeError(); }
-    operator DeclResult() const { return DeclResult(true); }
-    operator MemInitResult() const { return MemInitResult(true); }
-
-  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<ImmediateDiagBuilder> ImmediateDiag;
-    llvm::Optional<unsigned> PartialDiagId;
   };
-  using DiagBuilderT = SemaDiagnosticBuilder;
-
-  /// Is the last error level diagnostic immediate. This is used to determined
-  /// whether the next info diagnostic should be immediate.
-  bool IsLastErrorImmediate = true;
 
   /// Emit a diagnostic.
-  SemaDiagnosticBuilder Diag(SourceLocation Loc, unsigned DiagID,
-                             bool DeferHint = false);
+  SemaDiagnosticBuilder Diag(SourceLocation Loc, unsigned DiagID) {
+    DiagnosticBuilder DB = Diags.Report(Loc, DiagID);
+    return SemaDiagnosticBuilder(DB, *this, DiagID);
+  }
 
   /// Emit a partial diagnostic.
-  SemaDiagnosticBuilder Diag(SourceLocation Loc, const PartialDiagnostic &PD,
-                             bool DeferHint = false);
+  SemaDiagnosticBuilder Diag(SourceLocation Loc, const PartialDiagnostic& PD);
 
   /// Build a partial diagnostic.
   PartialDiagnostic PDiag(unsigned DiagID = 0); // in SemaInternal.h
 
-  /// Whether uncompilable error has occurred. This includes error happens
-  /// in deferred diagnostics.
-  bool hasUncompilableErrorOccurred() const;
-
   bool findMacroSpelling(SourceLocation &loc, StringRef name);
 
   /// Get a string to suggest for zero-initialization of a type.
@@ -11814,11 +11688,84 @@ class Sema final {
                  /* Caller = */ FunctionDeclAndLoc>
       DeviceKnownEmittedFns;
 
-  /// Creates a SemaDiagnosticBuilder that emits the diagnostic if the current
-  /// context is "used as device code".
+  /// 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.
   ///
-  /// - If CurContext is a __host__ function, does not emit any diagnostics
-  ///   unless \p EmitOnBothSides is true.
+  /// 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 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
+    };
+
+    DeviceDiagBuilder(Kind K, SourceLocation Loc, unsigned DiagID,
+                      FunctionDecl *Fn, Sema &S);
+    DeviceDiagBuilder(DeviceDiagBuilder &&D);
+    DeviceDiagBuilder(const DeviceDiagBuilder &) = default;
+    ~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 (DeviceDiagBuilder(...) << foo << bar)
+    ///     return ExprError();
+    ///
+    /// But see CUDADiagIfDeviceCode() and CUDADiagIfHostCode() -- you probably
+    /// want to use these instead of creating a DeviceDiagBuilder yourself.
+    operator bool() const { return ImmediateDiag.hasValue(); }
+
+    template <typename T>
+    friend const DeviceDiagBuilder &operator<<(const DeviceDiagBuilder &Diag,
+                                               const T &Value) {
+      if (Diag.ImmediateDiag.hasValue())
+        *Diag.ImmediateDiag << Value;
+      else if (Diag.PartialDiagId.hasValue())
+        Diag.S.DeviceDeferredDiags[Diag.Fn][*Diag.PartialDiagId].second
+            << 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<unsigned> PartialDiagId;
+  };
+
+  /// 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
@@ -11831,16 +11778,15 @@ class Sema final {
   ///  if (CUDADiagIfDeviceCode(Loc, diag::err_cuda_vla) << CurrentCUDATarget())
   ///    return ExprError();
   ///  // Otherwise, continue parsing as normal.
-  SemaDiagnosticBuilder CUDADiagIfDeviceCode(SourceLocation Loc,
-                                             unsigned DiagID);
+  DeviceDiagBuilder CUDADiagIfDeviceCode(SourceLocation Loc, unsigned DiagID);
 
-  /// Creates a SemaDiagnosticBuilder that emits the diagnostic if the current
-  /// context is "used as host code".
+  /// Creates a DeviceDiagBuilder that emits the diagnostic if the current context
+  /// is "used as host code".
   ///
   /// Same as CUDADiagIfDeviceCode, with "host" and "device" switched.
-  SemaDiagnosticBuilder CUDADiagIfHostCode(SourceLocation Loc, unsigned DiagID);
+  DeviceDiagBuilder CUDADiagIfHostCode(SourceLocation Loc, unsigned DiagID);
 
-  /// Creates a SemaDiagnosticBuilder that emits the diagnostic if the current
+  /// Creates a DeviceDiagBuilder that emits the diagnostic if the current
   /// context is "used as device code".
   ///
   /// - If CurContext is a `declare target` function or it is known that the
@@ -11855,10 +11801,9 @@ class Sema final {
   ///  if (diagIfOpenMPDeviceCode(Loc, diag::err_vla_unsupported))
   ///    return ExprError();
   ///  // Otherwise, continue parsing as normal.
-  SemaDiagnosticBuilder diagIfOpenMPDeviceCode(SourceLocation Loc,
-                                               unsigned DiagID);
+  DeviceDiagBuilder diagIfOpenMPDeviceCode(SourceLocation Loc, unsigned DiagID);
 
-  /// Creates a SemaDiagnosticBuilder that emits the diagnostic if the current
+  /// Creates a DeviceDiagBuilder that emits the diagnostic if the current
   /// context is "used as host code".
   ///
   /// - If CurContext is a `declare target` function or it is known that the
@@ -11871,14 +11816,9 @@ class Sema final {
   ///  if (diagIfOpenMPHostode(Loc, diag::err_vla_unsupported))
   ///    return ExprError();
   ///  // Otherwise, continue parsing as normal.
-  SemaDiagnosticBuilder diagIfOpenMPHostCode(SourceLocation Loc,
-                                             unsigned DiagID);
+  DeviceDiagBuilder diagIfOpenMPHostCode(SourceLocation Loc, unsigned DiagID);
 
-  SemaDiagnosticBuilder targetDiag(SourceLocation Loc, unsigned DiagID);
-  SemaDiagnosticBuilder targetDiag(SourceLocation Loc,
-                                   const PartialDiagnostic &PD) {
-    return targetDiag(Loc, PD.getDiagID()) << PD;
-  }
+  DeviceDiagBuilder targetDiag(SourceLocation Loc, unsigned DiagID);
 
   /// Check if the expression is allowed to be used in expressions for the
   /// offloading devices.
@@ -12651,7 +12591,7 @@ class Sema final {
     ConstructorDestructor,
     BuiltinFunction
   };
-  /// Creates a SemaDiagnosticBuilder that emits the diagnostic if the current
+  /// Creates a DeviceDiagBuilder that emits the diagnostic if the current
   /// context is "used as device code".
   ///
   /// - If CurLexicalContext is a kernel function or it is known that the
@@ -12669,8 +12609,7 @@ class Sema final {
   /// if (!S.Context.getTargetInfo().hasFloat128Type() &&
   ///     S.getLangOpts().SYCLIsDevice)
   ///   SYCLDiagIfDeviceCode(Loc, diag::err_type_unsupported) << "__float128";
-  SemaDiagnosticBuilder SYCLDiagIfDeviceCode(SourceLocation Loc,
-                                             unsigned DiagID);
+  DeviceDiagBuilder SYCLDiagIfDeviceCode(SourceLocation Loc, unsigned DiagID);
 
   /// Check whether we're allowed to call Callee from the current context.
   ///

diff  --git a/clang/lib/Basic/DiagnosticIDs.cpp b/clang/lib/Basic/DiagnosticIDs.cpp
index b641077f1ac2..ef615036082e 100644
--- a/clang/lib/Basic/DiagnosticIDs.cpp
+++ b/clang/lib/Basic/DiagnosticIDs.cpp
@@ -33,7 +33,7 @@ struct StaticDiagInfoRec;
 // platforms. See "How To Write Shared Libraries" by Ulrich Drepper.
 struct StaticDiagInfoDescriptionStringTable {
 #define DIAG(ENUM, CLASS, DEFAULT_SEVERITY, DESC, GROUP, SFINAE, NOWERROR,     \
-             SHOWINSYSHEADER, DEFERRABLE, CATEGORY)                            \
+             SHOWINSYSHEADER, CATEGORY)                                        \
   char ENUM##_desc[sizeof(DESC)];
   // clang-format off
 #include "clang/Basic/DiagnosticCommonKinds.inc"
@@ -54,9 +54,9 @@ struct StaticDiagInfoDescriptionStringTable {
 
 const StaticDiagInfoDescriptionStringTable StaticDiagInfoDescriptions = {
 #define DIAG(ENUM, CLASS, DEFAULT_SEVERITY, DESC, GROUP, SFINAE, NOWERROR,     \
-             SHOWINSYSHEADER, DEFERRABLE, CATEGORY)                            \
+             SHOWINSYSHEADER, CATEGORY)                                        \
   DESC,
-  // clang-format off
+// clang-format off
 #include "clang/Basic/DiagnosticCommonKinds.inc"
 #include "clang/Basic/DiagnosticDriverKinds.inc"
 #include "clang/Basic/DiagnosticFrontendKinds.inc"
@@ -79,9 +79,9 @@ extern const StaticDiagInfoRec StaticDiagInfo[];
 // StaticDiagInfoRec would have extra padding on 64-bit platforms.
 const uint32_t StaticDiagInfoDescriptionOffsets[] = {
 #define DIAG(ENUM, CLASS, DEFAULT_SEVERITY, DESC, GROUP, SFINAE, NOWERROR,     \
-             SHOWINSYSHEADER, DEFERRABLE, CATEGORY)                            \
+             SHOWINSYSHEADER, CATEGORY)                                        \
   offsetof(StaticDiagInfoDescriptionStringTable, ENUM##_desc),
-  // clang-format off
+// clang-format off
 #include "clang/Basic/DiagnosticCommonKinds.inc"
 #include "clang/Basic/DiagnosticDriverKinds.inc"
 #include "clang/Basic/DiagnosticFrontendKinds.inc"
@@ -114,7 +114,6 @@ struct StaticDiagInfoRec {
   unsigned SFINAE : 2;
   unsigned WarnNoWerror : 1;
   unsigned WarnShowInSystemHeader : 1;
-  unsigned Deferrable : 1;
   unsigned Category : 6;
 
   uint16_t OptionGroupIndex;
@@ -169,7 +168,7 @@ VALIDATE_DIAG_SIZE(REFACTORING)
 
 const StaticDiagInfoRec StaticDiagInfo[] = {
 #define DIAG(ENUM, CLASS, DEFAULT_SEVERITY, DESC, GROUP, SFINAE, NOWERROR,     \
-             SHOWINSYSHEADER, DEFERRABLE, CATEGORY)                            \
+             SHOWINSYSHEADER, CATEGORY)                                        \
   {                                                                            \
       diag::ENUM,                                                              \
       DEFAULT_SEVERITY,                                                        \
@@ -177,11 +176,10 @@ const StaticDiagInfoRec StaticDiagInfo[] = {
       DiagnosticIDs::SFINAE,                                                   \
       NOWERROR,                                                                \
       SHOWINSYSHEADER,                                                         \
-      DEFERRABLE,                                                              \
       CATEGORY,                                                                \
       GROUP,                                                                   \
       STR_SIZE(DESC, uint16_t)},
-  // clang-format off
+// clang-format off
 #include "clang/Basic/DiagnosticCommonKinds.inc"
 #include "clang/Basic/DiagnosticDriverKinds.inc"
 #include "clang/Basic/DiagnosticFrontendKinds.inc"
@@ -336,12 +334,6 @@ DiagnosticIDs::getDiagnosticSFINAEResponse(unsigned DiagID) {
   return SFINAE_Report;
 }
 
-bool DiagnosticIDs::isDeferrable(unsigned DiagID) {
-  if (const StaticDiagInfoRec *Info = GetDiagInfo(DiagID))
-    return Info->Deferrable;
-  return false;
-}
-
 /// getBuiltinDiagClass - Return the class field of the diagnostic.
 ///
 static unsigned getBuiltinDiagClass(unsigned DiagID) {

diff  --git a/clang/lib/Driver/ToolChains/Cuda.cpp b/clang/lib/Driver/ToolChains/Cuda.cpp
index f8af765f600f..d7933534a5d3 100644
--- a/clang/lib/Driver/ToolChains/Cuda.cpp
+++ b/clang/lib/Driver/ToolChains/Cuda.cpp
@@ -634,10 +634,6 @@ void CudaToolChain::addClangTargetOptions(
     if (DriverArgs.hasFlag(options::OPT_fgpu_rdc, options::OPT_fno_gpu_rdc,
                            false))
       CC1Args.push_back("-fgpu-rdc");
-
-    if (DriverArgs.hasFlag(options::OPT_fgpu_defer_diag,
-                           options::OPT_fno_gpu_defer_diag, false))
-      CC1Args.push_back("-fgpu-defer-diag");
   }
 
   if (DriverArgs.hasArg(options::OPT_nogpulib))

diff  --git a/clang/lib/Driver/ToolChains/HIP.cpp b/clang/lib/Driver/ToolChains/HIP.cpp
index 13bd59f926f5..43e557c98050 100644
--- a/clang/lib/Driver/ToolChains/HIP.cpp
+++ b/clang/lib/Driver/ToolChains/HIP.cpp
@@ -268,10 +268,6 @@ void HIPToolChain::addClangTargetOptions(
                          options::OPT_fno_gpu_allow_device_init, false))
     CC1Args.push_back("-fgpu-allow-device-init");
 
-  if (DriverArgs.hasFlag(options::OPT_fgpu_defer_diag,
-                         options::OPT_fno_gpu_defer_diag, false))
-    CC1Args.push_back("-fgpu-defer-diag");
-
   CC1Args.push_back("-fcuda-allow-variadic-functions");
 
   // Default to "hidden" visibility, as object level linking will not be

diff  --git a/clang/lib/Frontend/CompilerInvocation.cpp b/clang/lib/Frontend/CompilerInvocation.cpp
index 488a9dd0f8eb..a88a91182307 100644
--- a/clang/lib/Frontend/CompilerInvocation.cpp
+++ b/clang/lib/Frontend/CompilerInvocation.cpp
@@ -2632,9 +2632,6 @@ static void ParseLangArgs(LangOptions &Opts, ArgList &Args, InputKind IK,
   if (Args.hasArg(OPT_fno_cuda_host_device_constexpr))
     Opts.CUDAHostDeviceConstexpr = 0;
 
-  if (Args.hasArg(OPT_fgpu_defer_diag))
-    Opts.GPUDeferDiag = 1;
-
   if (Opts.CUDAIsDevice && Args.hasArg(OPT_fcuda_approx_transcendentals))
     Opts.CUDADeviceApproxTranscendentals = 1;
 

diff  --git a/clang/lib/Sema/AnalysisBasedWarnings.cpp b/clang/lib/Sema/AnalysisBasedWarnings.cpp
index 2850162141c9..37fd26d7c22d 100644
--- a/clang/lib/Sema/AnalysisBasedWarnings.cpp
+++ b/clang/lib/Sema/AnalysisBasedWarnings.cpp
@@ -2096,7 +2096,7 @@ AnalysisBasedWarnings::IssueWarnings(sema::AnalysisBasedWarnings::Policy P,
   if (cast<DeclContext>(D)->isDependentContext())
     return;
 
-  if (S.hasUncompilableErrorOccurred()) {
+  if (Diags.hasUncompilableErrorOccurred()) {
     // Flush out any possibly unreachable diagnostics.
     flushDiagnostics(S, fscope);
     return;

diff  --git a/clang/lib/Sema/Sema.cpp b/clang/lib/Sema/Sema.cpp
index 53ff2b62c437..375fe3b28dec 100644
--- a/clang/lib/Sema/Sema.cpp
+++ b/clang/lib/Sema/Sema.cpp
@@ -1436,24 +1436,11 @@ void Sema::EmitCurrentDiagnostic(unsigned DiagID) {
 }
 
 Sema::SemaDiagnosticBuilder
-Sema::Diag(SourceLocation Loc, const PartialDiagnostic &PD, bool DeferHint) {
-  return Diag(Loc, PD.getDiagID(), DeferHint) << PD;
-}
+Sema::Diag(SourceLocation Loc, const PartialDiagnostic& PD) {
+  SemaDiagnosticBuilder Builder(Diag(Loc, PD.getDiagID()));
+  PD.Emit(Builder);
 
-bool Sema::hasUncompilableErrorOccurred() const {
-  if (getDiagnostics().hasUncompilableErrorOccurred())
-    return true;
-  auto *FD = dyn_cast<FunctionDecl>(CurContext);
-  if (!FD)
-    return false;
-  auto Loc = DeviceDeferredDiags.find(FD);
-  if (Loc == DeviceDeferredDiags.end())
-    return false;
-  for (auto PDAt : Loc->second) {
-    if (DiagnosticIDs::isDefaultMappingAsError(PDAt.second.getDiagID()))
-      return true;
-  }
-  return false;
+  return Builder;
 }
 
 // Print notes showing how we can reach FD starting from an a priori
@@ -1666,9 +1653,9 @@ void Sema::emitDeferredDiags() {
 // until we discover that the function is known-emitted, at which point we take
 // it out of this map and emit the diagnostic.
 
-Sema::SemaDiagnosticBuilder::SemaDiagnosticBuilder(Kind K, SourceLocation Loc,
-                                                   unsigned DiagID,
-                                                   FunctionDecl *Fn, Sema &S)
+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) {
@@ -1676,8 +1663,7 @@ Sema::SemaDiagnosticBuilder::SemaDiagnosticBuilder(Kind K, SourceLocation Loc,
     break;
   case K_Immediate:
   case K_ImmediateWithCallStack:
-    ImmediateDiag.emplace(
-        ImmediateDiagBuilder(S.Diags.Report(Loc, DiagID), S, DiagID));
+    ImmediateDiag.emplace(S.Diag(Loc, DiagID));
     break;
   case K_Deferred:
     assert(Fn && "Must have a function to attach the deferred diag to.");
@@ -1688,7 +1674,7 @@ Sema::SemaDiagnosticBuilder::SemaDiagnosticBuilder(Kind K, SourceLocation Loc,
   }
 }
 
-Sema::SemaDiagnosticBuilder::SemaDiagnosticBuilder(SemaDiagnosticBuilder &&D)
+Sema::DeviceDiagBuilder::DeviceDiagBuilder(DeviceDiagBuilder &&D)
     : S(D.S), Loc(D.Loc), DiagID(D.DiagID), Fn(D.Fn),
       ShowCallStack(D.ShowCallStack), ImmediateDiag(D.ImmediateDiag),
       PartialDiagId(D.PartialDiagId) {
@@ -1698,7 +1684,7 @@ Sema::SemaDiagnosticBuilder::SemaDiagnosticBuilder(SemaDiagnosticBuilder &&D)
   D.PartialDiagId.reset();
 }
 
-Sema::SemaDiagnosticBuilder::~SemaDiagnosticBuilder() {
+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.
@@ -1713,8 +1699,7 @@ Sema::SemaDiagnosticBuilder::~SemaDiagnosticBuilder() {
   }
 }
 
-Sema::SemaDiagnosticBuilder Sema::targetDiag(SourceLocation Loc,
-                                             unsigned DiagID) {
+Sema::DeviceDiagBuilder Sema::targetDiag(SourceLocation Loc, unsigned DiagID) {
   if (LangOpts.OpenMP)
     return LangOpts.OpenMPIsDevice ? diagIfOpenMPDeviceCode(Loc, DiagID)
                                    : diagIfOpenMPHostCode(Loc, DiagID);
@@ -1725,32 +1710,8 @@ Sema::SemaDiagnosticBuilder Sema::targetDiag(SourceLocation Loc,
   if (getLangOpts().SYCLIsDevice)
     return SYCLDiagIfDeviceCode(Loc, DiagID);
 
-  return SemaDiagnosticBuilder(SemaDiagnosticBuilder::K_Immediate, Loc, DiagID,
-                               getCurFunctionDecl(), *this);
-}
-
-Sema::SemaDiagnosticBuilder Sema::Diag(SourceLocation Loc, unsigned DiagID,
-                                       bool DeferHint) {
-  bool IsError = Diags.getDiagnosticIDs()->isDefaultMappingAsError(DiagID);
-  bool ShouldDefer = getLangOpts().CUDA && LangOpts.GPUDeferDiag &&
-                     DiagnosticIDs::isDeferrable(DiagID) &&
-                     (DeferHint || !IsError);
-  auto SetIsLastErrorImmediate = [&](bool Flag) {
-    if (IsError)
-      IsLastErrorImmediate = Flag;
-  };
-  if (!ShouldDefer) {
-    SetIsLastErrorImmediate(true);
-    return SemaDiagnosticBuilder(SemaDiagnosticBuilder::K_Immediate, Loc,
-                                 DiagID, getCurFunctionDecl(), *this);
-  }
-
-  SemaDiagnosticBuilder DB =
-      getLangOpts().CUDAIsDevice
-          ? CUDADiagIfDeviceCode(Loc, DiagID)
-          : CUDADiagIfHostCode(Loc, DiagID);
-  SetIsLastErrorImmediate(DB.isImmediate());
-  return DB;
+  return DeviceDiagBuilder(DeviceDiagBuilder::K_Immediate, Loc, DiagID,
+                           getCurFunctionDecl(), *this);
 }
 
 void Sema::checkDeviceDecl(const ValueDecl *D, SourceLocation Loc) {

diff  --git a/clang/lib/Sema/SemaAttr.cpp b/clang/lib/Sema/SemaAttr.cpp
index 1e58627e94a3..bd5fc586b6af 100644
--- a/clang/lib/Sema/SemaAttr.cpp
+++ b/clang/lib/Sema/SemaAttr.cpp
@@ -380,8 +380,8 @@ void Sema::DiagnoseUnterminatedPragmaPack() {
     // The user might have already reset the alignment, so suggest replacing
     // the reset with a pop.
     if (IsInnermost && PackStack.CurrentValue == PackStack.DefaultValue) {
-      auto DB = Diag(PackStack.CurrentPragmaLocation,
-                     diag::note_pragma_pack_pop_instead_reset);
+      DiagnosticBuilder DB = Diag(PackStack.CurrentPragmaLocation,
+                                  diag::note_pragma_pack_pop_instead_reset);
       SourceLocation FixItLoc = Lexer::findLocationAfterToken(
           PackStack.CurrentPragmaLocation, tok::l_paren, SourceMgr, LangOpts,
           /*SkipTrailing=*/false);

diff  --git a/clang/lib/Sema/SemaCUDA.cpp b/clang/lib/Sema/SemaCUDA.cpp
index 13c735678583..6203edea7112 100644
--- a/clang/lib/Sema/SemaCUDA.cpp
+++ b/clang/lib/Sema/SemaCUDA.cpp
@@ -639,63 +639,58 @@ void Sema::MaybeAddCUDAConstantAttr(VarDecl *VD) {
   }
 }
 
-Sema::SemaDiagnosticBuilder Sema::CUDADiagIfDeviceCode(SourceLocation Loc,
-                                                       unsigned DiagID) {
+Sema::DeviceDiagBuilder Sema::CUDADiagIfDeviceCode(SourceLocation Loc,
+                                                   unsigned DiagID) {
   assert(getLangOpts().CUDA && "Should only be called during CUDA compilation");
-  SemaDiagnosticBuilder::Kind DiagKind = [&] {
-    if (!isa<FunctionDecl>(CurContext))
-      return SemaDiagnosticBuilder::K_Immediate;
+  DeviceDiagBuilder::Kind DiagKind = [this] {
     switch (CurrentCUDATarget()) {
     case CFT_Global:
     case CFT_Device:
-      return SemaDiagnosticBuilder::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 SemaDiagnosticBuilder::K_Nop;
-      if (IsLastErrorImmediate && Diags.getDiagnosticIDs()->isBuiltinNote(DiagID))
-        return SemaDiagnosticBuilder::K_Immediate;
-      return (getEmissionStatus(cast<FunctionDecl>(CurContext)) ==
-              FunctionEmissionStatus::Emitted)
-                 ? SemaDiagnosticBuilder::K_ImmediateWithCallStack
-                 : SemaDiagnosticBuilder::K_Deferred;
+      if (getLangOpts().CUDAIsDevice) {
+        return (getEmissionStatus(cast<FunctionDecl>(CurContext)) ==
+                FunctionEmissionStatus::Emitted)
+                   ? DeviceDiagBuilder::K_ImmediateWithCallStack
+                   : DeviceDiagBuilder::K_Deferred;
+      }
+      return DeviceDiagBuilder::K_Nop;
+
     default:
-      return SemaDiagnosticBuilder::K_Nop;
+      return DeviceDiagBuilder::K_Nop;
     }
   }();
-  return SemaDiagnosticBuilder(DiagKind, Loc, DiagID,
-                               dyn_cast<FunctionDecl>(CurContext), *this);
+  return DeviceDiagBuilder(DiagKind, Loc, DiagID,
+                           dyn_cast<FunctionDecl>(CurContext), *this);
 }
 
-Sema::SemaDiagnosticBuilder Sema::CUDADiagIfHostCode(SourceLocation Loc,
-                                                     unsigned DiagID) {
+Sema::DeviceDiagBuilder Sema::CUDADiagIfHostCode(SourceLocation Loc,
+                                                 unsigned DiagID) {
   assert(getLangOpts().CUDA && "Should only be called during CUDA compilation");
-  SemaDiagnosticBuilder::Kind DiagKind = [&] {
-    if (!isa<FunctionDecl>(CurContext))
-      return SemaDiagnosticBuilder::K_Immediate;
+  DeviceDiagBuilder::Kind DiagKind = [this] {
     switch (CurrentCUDATarget()) {
     case CFT_Host:
-      return SemaDiagnosticBuilder::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 SemaDiagnosticBuilder::K_Nop;
-      if (IsLastErrorImmediate && Diags.getDiagnosticIDs()->isBuiltinNote(DiagID))
-        return SemaDiagnosticBuilder::K_Immediate;
+        return DeviceDiagBuilder::K_Nop;
+
       return (getEmissionStatus(cast<FunctionDecl>(CurContext)) ==
               FunctionEmissionStatus::Emitted)
-                 ? SemaDiagnosticBuilder::K_ImmediateWithCallStack
-                 : SemaDiagnosticBuilder::K_Deferred;
+                 ? DeviceDiagBuilder::K_ImmediateWithCallStack
+                 : DeviceDiagBuilder::K_Deferred;
     default:
-      return SemaDiagnosticBuilder::K_Nop;
+      return DeviceDiagBuilder::K_Nop;
     }
   }();
-  return SemaDiagnosticBuilder(DiagKind, Loc, DiagID,
-                               dyn_cast<FunctionDecl>(CurContext), *this);
+  return DeviceDiagBuilder(DiagKind, Loc, DiagID,
+                           dyn_cast<FunctionDecl>(CurContext), *this);
 }
 
 bool Sema::CheckCUDACall(SourceLocation Loc, FunctionDecl *Callee) {
@@ -716,8 +711,8 @@ 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;
-  SemaDiagnosticBuilder::Kind DiagKind = [this, Caller, Callee,
-                                          CallerKnownEmitted] {
+  DeviceDiagBuilder::Kind DiagKind = [this, Caller, Callee,
+                                      CallerKnownEmitted] {
     switch (IdentifyCUDAPreference(Caller, Callee)) {
     case CFP_Never:
     case CFP_WrongSide:
@@ -725,15 +720,14 @@ bool Sema::CheckCUDACall(SourceLocation Loc, FunctionDecl *Callee) {
       // 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
-                 ? SemaDiagnosticBuilder::K_ImmediateWithCallStack
-                 : SemaDiagnosticBuilder::K_Deferred;
+      return CallerKnownEmitted ? DeviceDiagBuilder::K_ImmediateWithCallStack
+                                : DeviceDiagBuilder::K_Deferred;
     default:
-      return SemaDiagnosticBuilder::K_Nop;
+      return DeviceDiagBuilder::K_Nop;
     }
   }();
 
-  if (DiagKind == SemaDiagnosticBuilder::K_Nop)
+  if (DiagKind == DeviceDiagBuilder::K_Nop)
     return true;
 
   // Avoid emitting this error twice for the same location.  Using a hashtable
@@ -743,14 +737,14 @@ bool Sema::CheckCUDACall(SourceLocation Loc, FunctionDecl *Callee) {
   if (!LocsWithCUDACallDiags.insert({Caller, Loc}).second)
     return true;
 
-  SemaDiagnosticBuilder(DiagKind, Loc, diag::err_ref_bad_target, Caller, *this)
+  DeviceDiagBuilder(DiagKind, Loc, diag::err_ref_bad_target, Caller, *this)
       << IdentifyCUDATarget(Callee) << Callee << IdentifyCUDATarget(Caller);
   if (!Callee->getBuiltinID())
-    SemaDiagnosticBuilder(DiagKind, Callee->getLocation(),
-                          diag::note_previous_decl, Caller, *this)
+    DeviceDiagBuilder(DiagKind, Callee->getLocation(), diag::note_previous_decl,
+                      Caller, *this)
         << Callee;
-  return DiagKind != SemaDiagnosticBuilder::K_Immediate &&
-         DiagKind != SemaDiagnosticBuilder::K_ImmediateWithCallStack;
+  return DiagKind != DeviceDiagBuilder::K_Immediate &&
+         DiagKind != DeviceDiagBuilder::K_ImmediateWithCallStack;
 }
 
 // Check the wrong-sided reference capture of lambda for CUDA/HIP.
@@ -787,14 +781,14 @@ void Sema::CUDACheckLambdaCapture(CXXMethodDecl *Callee,
   bool ShouldCheck = CalleeIsDevice && CallerIsHost;
   if (!ShouldCheck || !Capture.isReferenceCapture())
     return;
-  auto DiagKind = SemaDiagnosticBuilder::K_Deferred;
+  auto DiagKind = DeviceDiagBuilder::K_Deferred;
   if (Capture.isVariableCapture()) {
-    SemaDiagnosticBuilder(DiagKind, Capture.getLocation(),
-                          diag::err_capture_bad_target, Callee, *this)
+    DeviceDiagBuilder(DiagKind, Capture.getLocation(),
+                      diag::err_capture_bad_target, Callee, *this)
         << Capture.getVariable();
   } else if (Capture.isThisCapture()) {
-    SemaDiagnosticBuilder(DiagKind, Capture.getLocation(),
-                          diag::err_capture_bad_target_this_ptr, Callee, *this);
+    DeviceDiagBuilder(DiagKind, Capture.getLocation(),
+                      diag::err_capture_bad_target_this_ptr, Callee, *this);
   }
   return;
 }

diff  --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp
index cf0413669345..434f13e38e3c 100644
--- a/clang/lib/Sema/SemaDecl.cpp
+++ b/clang/lib/Sema/SemaDecl.cpp
@@ -14549,11 +14549,11 @@ Decl *Sema::ActOnFinishFunctionBody(Decl *dcl, Stmt *Body,
     // If any errors have occurred, clear out any temporaries that may have
     // been leftover. This ensures that these temporaries won't be picked up for
     // deletion in some later function.
-    if (hasUncompilableErrorOccurred() ||
+    if (getDiagnostics().hasUncompilableErrorOccurred() ||
         getDiagnostics().getSuppressAllDiagnostics()) {
       DiscardCleanupsInEvaluationContext();
     }
-    if (!hasUncompilableErrorOccurred() &&
+    if (!getDiagnostics().hasUncompilableErrorOccurred() &&
         !isa<FunctionTemplateDecl>(dcl)) {
       // Since the body is valid, issue any analysis-based warnings that are
       // enabled.
@@ -14605,7 +14605,7 @@ Decl *Sema::ActOnFinishFunctionBody(Decl *dcl, Stmt *Body,
   // If any errors have occurred, clear out any temporaries that may have
   // been leftover. This ensures that these temporaries won't be picked up for
   // deletion in some later function.
-  if (hasUncompilableErrorOccurred()) {
+  if (getDiagnostics().hasUncompilableErrorOccurred()) {
     DiscardCleanupsInEvaluationContext();
   }
 

diff  --git a/clang/lib/Sema/SemaExprObjC.cpp b/clang/lib/Sema/SemaExprObjC.cpp
index 60587db0cc69..2c088c8b15a3 100644
--- a/clang/lib/Sema/SemaExprObjC.cpp
+++ b/clang/lib/Sema/SemaExprObjC.cpp
@@ -2445,8 +2445,8 @@ static void applyCocoaAPICheck(Sema &S, const ObjCMessageExpr *Msg,
   SourceManager &SM = S.SourceMgr;
   edit::Commit ECommit(SM, S.LangOpts);
   if (refactor(Msg,*S.NSAPIObj, ECommit)) {
-    auto Builder = S.Diag(MsgLoc, DiagID)
-                   << Msg->getSelector() << Msg->getSourceRange();
+    DiagnosticBuilder Builder = S.Diag(MsgLoc, DiagID)
+                        << Msg->getSelector() << Msg->getSourceRange();
     // FIXME: Don't emit diagnostic at all if fixits are non-commitable.
     if (!ECommit.isCommitable())
       return;
@@ -3139,8 +3139,9 @@ ExprResult Sema::BuildInstanceMessage(Expr *Receiver,
     if (ReceiverType->isObjCClassType() && !isImplicit &&
         !(Receiver->isObjCSelfExpr() && getLangOpts().ObjCAutoRefCount)) {
       {
-        auto Builder = Diag(Receiver->getExprLoc(),
-                            diag::err_messaging_class_with_direct_method);
+        DiagnosticBuilder Builder =
+            Diag(Receiver->getExprLoc(),
+                 diag::err_messaging_class_with_direct_method);
         if (Receiver->isObjCSelfExpr()) {
           Builder.AddFixItHint(FixItHint::CreateReplacement(
               RecRange, Method->getClassInterface()->getName()));
@@ -3152,7 +3153,7 @@ ExprResult Sema::BuildInstanceMessage(Expr *Receiver,
 
     if (SuperLoc.isValid()) {
       {
-        auto Builder =
+        DiagnosticBuilder Builder =
             Diag(SuperLoc, diag::err_messaging_super_with_direct_method);
         if (ReceiverType->isObjCClassType()) {
           Builder.AddFixItHint(FixItHint::CreateReplacement(
@@ -3735,11 +3736,15 @@ bool Sema::isKnownName(StringRef name) {
   return LookupName(R, TUScope, false);
 }
 
-template <typename DiagBuilderT>
-static void addFixitForObjCARCConversion(
-    Sema &S, DiagBuilderT &DiagB, Sema::CheckedConversionKind CCK,
-    SourceLocation afterLParen, QualType castType, Expr *castExpr,
-    Expr *realCast, const char *bridgeKeyword, const char *CFBridgeName) {
+static void addFixitForObjCARCConversion(Sema &S,
+                                         DiagnosticBuilder &DiagB,
+                                         Sema::CheckedConversionKind CCK,
+                                         SourceLocation afterLParen,
+                                         QualType castType,
+                                         Expr *castExpr,
+                                         Expr *realCast,
+                                         const char *bridgeKeyword,
+                                         const char *CFBridgeName) {
   // We handle C-style and implicit casts here.
   switch (CCK) {
   case Sema::CCK_ImplicitConversion:
@@ -3916,9 +3921,9 @@ diagnoseObjCARCConversion(Sema &S, SourceRange castRange,
     assert(CreateRule != ACC_bottom && "This cast should already be accepted.");
     if (CreateRule != ACC_plusOne)
     {
-      auto DiagB = (CCK != Sema::CCK_OtherCast)
-                       ? S.Diag(noteLoc, diag::note_arc_bridge)
-                       : S.Diag(noteLoc, diag::note_arc_cstyle_bridge);
+      DiagnosticBuilder DiagB =
+        (CCK != Sema::CCK_OtherCast) ? S.Diag(noteLoc, diag::note_arc_bridge)
+                              : S.Diag(noteLoc, diag::note_arc_cstyle_bridge);
 
       addFixitForObjCARCConversion(S, DiagB, CCK, afterLParen,
                                    castType, castExpr, realCast, "__bridge ",
@@ -3926,12 +3931,12 @@ diagnoseObjCARCConversion(Sema &S, SourceRange castRange,
     }
     if (CreateRule != ACC_plusZero)
     {
-      auto DiagB = (CCK == Sema::CCK_OtherCast && !br)
-                       ? S.Diag(noteLoc, diag::note_arc_cstyle_bridge_transfer)
-                             << castExprType
-                       : S.Diag(br ? castExpr->getExprLoc() : noteLoc,
-                                diag::note_arc_bridge_transfer)
-                             << castExprType << br;
+      DiagnosticBuilder DiagB =
+        (CCK == Sema::CCK_OtherCast && !br) ?
+          S.Diag(noteLoc, diag::note_arc_cstyle_bridge_transfer) << castExprType :
+          S.Diag(br ? castExpr->getExprLoc() : noteLoc,
+                 diag::note_arc_bridge_transfer)
+            << castExprType << br;
 
       addFixitForObjCARCConversion(S, DiagB, CCK, afterLParen,
                                    castType, castExpr, realCast, "__bridge_transfer ",
@@ -3957,21 +3962,21 @@ diagnoseObjCARCConversion(Sema &S, SourceRange castRange,
     assert(CreateRule != ACC_bottom && "This cast should already be accepted.");
     if (CreateRule != ACC_plusOne)
     {
-      auto DiagB = (CCK != Sema::CCK_OtherCast)
-                       ? S.Diag(noteLoc, diag::note_arc_bridge)
-                       : S.Diag(noteLoc, diag::note_arc_cstyle_bridge);
+      DiagnosticBuilder DiagB =
+      (CCK != Sema::CCK_OtherCast) ? S.Diag(noteLoc, diag::note_arc_bridge)
+                               : S.Diag(noteLoc, diag::note_arc_cstyle_bridge);
       addFixitForObjCARCConversion(S, DiagB, CCK, afterLParen,
                                    castType, castExpr, realCast, "__bridge ",
                                    nullptr);
     }
     if (CreateRule != ACC_plusZero)
     {
-      auto DiagB = (CCK == Sema::CCK_OtherCast && !br)
-                       ? S.Diag(noteLoc, diag::note_arc_cstyle_bridge_retained)
-                             << castType
-                       : S.Diag(br ? castExpr->getExprLoc() : noteLoc,
-                                diag::note_arc_bridge_retained)
-                             << castType << br;
+      DiagnosticBuilder DiagB =
+        (CCK == Sema::CCK_OtherCast && !br) ?
+          S.Diag(noteLoc, diag::note_arc_cstyle_bridge_retained) << castType :
+          S.Diag(br ? castExpr->getExprLoc() : noteLoc,
+                 diag::note_arc_bridge_retained)
+            << castType << br;
 
       addFixitForObjCARCConversion(S, DiagB, CCK, afterLParen,
                                    castType, castExpr, realCast, "__bridge_retained ",

diff  --git a/clang/lib/Sema/SemaOpenMP.cpp b/clang/lib/Sema/SemaOpenMP.cpp
index 5094724c3118..91897cbe25b4 100644
--- a/clang/lib/Sema/SemaOpenMP.cpp
+++ b/clang/lib/Sema/SemaOpenMP.cpp
@@ -1896,27 +1896,27 @@ enum class FunctionEmissionStatus {
 };
 } // anonymous namespace
 
-Sema::SemaDiagnosticBuilder Sema::diagIfOpenMPDeviceCode(SourceLocation Loc,
-                                                         unsigned DiagID) {
+Sema::DeviceDiagBuilder Sema::diagIfOpenMPDeviceCode(SourceLocation Loc,
+                                                     unsigned DiagID) {
   assert(LangOpts.OpenMP && LangOpts.OpenMPIsDevice &&
          "Expected OpenMP device compilation.");
 
   FunctionDecl *FD = getCurFunctionDecl();
-  SemaDiagnosticBuilder::Kind Kind = SemaDiagnosticBuilder::K_Nop;
+  DeviceDiagBuilder::Kind Kind = DeviceDiagBuilder::K_Nop;
   if (FD) {
     FunctionEmissionStatus FES = getEmissionStatus(FD);
     switch (FES) {
     case FunctionEmissionStatus::Emitted:
-      Kind = SemaDiagnosticBuilder::K_Immediate;
+      Kind = DeviceDiagBuilder::K_Immediate;
       break;
     case FunctionEmissionStatus::Unknown:
       Kind = isOpenMPDeviceDelayedContext(*this)
-                 ? SemaDiagnosticBuilder::K_Deferred
-                 : SemaDiagnosticBuilder::K_Immediate;
+                 ? DeviceDiagBuilder::K_Deferred
+                 : DeviceDiagBuilder::K_Immediate;
       break;
     case FunctionEmissionStatus::TemplateDiscarded:
     case FunctionEmissionStatus::OMPDiscarded:
-      Kind = SemaDiagnosticBuilder::K_Nop;
+      Kind = DeviceDiagBuilder::K_Nop;
       break;
     case FunctionEmissionStatus::CUDADiscarded:
       llvm_unreachable("CUDADiscarded unexpected in OpenMP device compilation");
@@ -1924,30 +1924,30 @@ Sema::SemaDiagnosticBuilder Sema::diagIfOpenMPDeviceCode(SourceLocation Loc,
     }
   }
 
-  return SemaDiagnosticBuilder(Kind, Loc, DiagID, getCurFunctionDecl(), *this);
+  return DeviceDiagBuilder(Kind, Loc, DiagID, getCurFunctionDecl(), *this);
 }
 
-Sema::SemaDiagnosticBuilder Sema::diagIfOpenMPHostCode(SourceLocation Loc,
-                                                       unsigned DiagID) {
+Sema::DeviceDiagBuilder Sema::diagIfOpenMPHostCode(SourceLocation Loc,
+                                                   unsigned DiagID) {
   assert(LangOpts.OpenMP && !LangOpts.OpenMPIsDevice &&
          "Expected OpenMP host compilation.");
   FunctionEmissionStatus FES = getEmissionStatus(getCurFunctionDecl());
-  SemaDiagnosticBuilder::Kind Kind = SemaDiagnosticBuilder::K_Nop;
+  DeviceDiagBuilder::Kind Kind = DeviceDiagBuilder::K_Nop;
   switch (FES) {
   case FunctionEmissionStatus::Emitted:
-    Kind = SemaDiagnosticBuilder::K_Immediate;
+    Kind = DeviceDiagBuilder::K_Immediate;
     break;
   case FunctionEmissionStatus::Unknown:
-    Kind = SemaDiagnosticBuilder::K_Deferred;
+    Kind = DeviceDiagBuilder::K_Deferred;
     break;
   case FunctionEmissionStatus::TemplateDiscarded:
   case FunctionEmissionStatus::OMPDiscarded:
   case FunctionEmissionStatus::CUDADiscarded:
-    Kind = SemaDiagnosticBuilder::K_Nop;
+    Kind = DeviceDiagBuilder::K_Nop;
     break;
   }
 
-  return SemaDiagnosticBuilder(Kind, Loc, DiagID, getCurFunctionDecl(), *this);
+  return DeviceDiagBuilder(Kind, Loc, DiagID, getCurFunctionDecl(), *this);
 }
 
 static OpenMPDefaultmapClauseKind

diff  --git a/clang/lib/Sema/SemaOverload.cpp b/clang/lib/Sema/SemaOverload.cpp
index 1252937ae65d..95d110e754f4 100644
--- a/clang/lib/Sema/SemaOverload.cpp
+++ b/clang/lib/Sema/SemaOverload.cpp
@@ -11522,19 +11522,9 @@ void OverloadCandidateSet::NoteCandidates(PartialDiagnosticAt PD,
     StringRef Opc, SourceLocation OpLoc,
     llvm::function_ref<bool(OverloadCandidate &)> Filter) {
 
-  bool DeferHint = false;
-  if (S.getLangOpts().CUDA && S.getLangOpts().GPUDeferDiag) {
-    // Defer diagnostic for CUDA/HIP if there are wrong-sided candidates.
-    auto WrongSidedCands =
-        CompleteCandidates(S, OCD_AllCandidates, Args, OpLoc, [](auto &Cand) {
-          return Cand.Viable == false &&
-                 Cand.FailureKind == ovl_fail_bad_target;
-        });
-    DeferHint = WrongSidedCands.size();
-  }
   auto Cands = CompleteCandidates(S, OCD, Args, OpLoc, Filter);
 
-  S.Diag(PD.first, PD.second, DeferHint);
+  S.Diag(PD.first, PD.second);
 
   NoteCandidates(S, Args, Cands, Opc, OpLoc);
 

diff  --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index af35052ee1e3..db7603b42f7b 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -17,19 +17,19 @@ using namespace clang;
 // SYCL device specific diagnostics implementation
 // -----------------------------------------------------------------------------
 
-Sema::SemaDiagnosticBuilder Sema::SYCLDiagIfDeviceCode(SourceLocation Loc,
-                                                       unsigned DiagID) {
+Sema::DeviceDiagBuilder Sema::SYCLDiagIfDeviceCode(SourceLocation Loc,
+                                                   unsigned DiagID) {
   assert(getLangOpts().SYCLIsDevice &&
          "Should only be called during SYCL compilation");
   FunctionDecl *FD = dyn_cast<FunctionDecl>(getCurLexicalContext());
-  SemaDiagnosticBuilder::Kind DiagKind = [this, FD] {
+  DeviceDiagBuilder::Kind DiagKind = [this, FD] {
     if (!FD)
-      return SemaDiagnosticBuilder::K_Nop;
+      return DeviceDiagBuilder::K_Nop;
     if (getEmissionStatus(FD) == Sema::FunctionEmissionStatus::Emitted)
-      return SemaDiagnosticBuilder::K_ImmediateWithCallStack;
-    return SemaDiagnosticBuilder::K_Deferred;
+      return DeviceDiagBuilder::K_ImmediateWithCallStack;
+    return DeviceDiagBuilder::K_Deferred;
   }();
-  return SemaDiagnosticBuilder(DiagKind, Loc, DiagID, FD, *this);
+  return DeviceDiagBuilder(DiagKind, Loc, DiagID, FD, *this);
 }
 
 bool Sema::checkSYCLDeviceFunction(SourceLocation Loc, FunctionDecl *Callee) {
@@ -42,8 +42,8 @@ bool Sema::checkSYCLDeviceFunction(SourceLocation Loc, FunctionDecl *Callee) {
   if (isUnevaluatedContext() || isConstantEvaluated())
     return true;
 
-  SemaDiagnosticBuilder::Kind DiagKind = SemaDiagnosticBuilder::K_Nop;
+  DeviceDiagBuilder::Kind DiagKind = DeviceDiagBuilder::K_Nop;
 
-  return DiagKind != SemaDiagnosticBuilder::K_Immediate &&
-         DiagKind != SemaDiagnosticBuilder::K_ImmediateWithCallStack;
+  return DiagKind != DeviceDiagBuilder::K_Immediate &&
+         DiagKind != DeviceDiagBuilder::K_ImmediateWithCallStack;
 }

diff  --git a/clang/lib/Sema/SemaStmt.cpp b/clang/lib/Sema/SemaStmt.cpp
index 0e860a663a7a..5b4aaa678974 100644
--- a/clang/lib/Sema/SemaStmt.cpp
+++ b/clang/lib/Sema/SemaStmt.cpp
@@ -1261,10 +1261,10 @@ Sema::ActOnFinishSwitchStmt(SourceLocation SwitchLoc, Stmt *Switch,
 
       // Produce a nice diagnostic if multiple values aren't handled.
       if (!UnhandledNames.empty()) {
-        auto DB = Diag(CondExpr->getExprLoc(), TheDefaultStmt
-                                                   ? diag::warn_def_missing_case
+        DiagnosticBuilder DB = Diag(CondExpr->getExprLoc(),
+                                    TheDefaultStmt ? diag::warn_def_missing_case
                                                    : diag::warn_missing_case)
-                  << (int)UnhandledNames.size();
+                               << (int)UnhandledNames.size();
 
         for (size_t I = 0, E = std::min(UnhandledNames.size(), (size_t)3);
              I != E; ++I)

diff  --git a/clang/lib/Sema/SemaStmtAsm.cpp b/clang/lib/Sema/SemaStmtAsm.cpp
index 3b631bf747c6..10fa24682f9c 100644
--- a/clang/lib/Sema/SemaStmtAsm.cpp
+++ b/clang/lib/Sema/SemaStmtAsm.cpp
@@ -448,9 +448,9 @@ StmtResult Sema::ActOnGCCAsmStmt(SourceLocation AsmLoc, bool IsSimple,
     unsigned Size = Context.getTypeSize(Ty);
     if (!Context.getTargetInfo().validateInputSize(FeatureMap,
                                                    Literal->getString(), Size))
-      return targetDiag(InputExpr->getBeginLoc(),
-                        diag::err_asm_invalid_input_size)
-             << Info.getConstraintStr();
+      return StmtResult(
+          targetDiag(InputExpr->getBeginLoc(), diag::err_asm_invalid_input_size)
+          << Info.getConstraintStr());
   }
 
   // Check that the clobbers are valid.

diff  --git a/clang/lib/Sema/SemaTemplateInstantiate.cpp b/clang/lib/Sema/SemaTemplateInstantiate.cpp
index 54049d177ac0..11e03c517d01 100644
--- a/clang/lib/Sema/SemaTemplateInstantiate.cpp
+++ b/clang/lib/Sema/SemaTemplateInstantiate.cpp
@@ -237,7 +237,7 @@ Sema::InstantiatingTemplate::InstantiatingTemplate(
   // error have occurred. Any diagnostics we might have raised will not be
   // visible, and we do not need to construct a correct AST.
   if (SemaRef.Diags.hasFatalErrorOccurred() &&
-      SemaRef.hasUncompilableErrorOccurred()) {
+      SemaRef.Diags.hasUncompilableErrorOccurred()) {
     Invalid = true;
     return;
   }

diff  --git a/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp b/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp
index 488c9fb5a213..cbbb44b82adc 100644
--- a/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp
+++ b/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp
@@ -6008,7 +6008,7 @@ NamedDecl *Sema::FindInstantiatedDecl(SourceLocation Loc, NamedDecl *D,
     if (!Result) {
       if (isa<UsingShadowDecl>(D)) {
         // UsingShadowDecls can instantiate to nothing because of using hiding.
-      } else if (hasUncompilableErrorOccurred()) {
+      } else if (Diags.hasUncompilableErrorOccurred()) {
         // We've already complained about some ill-formed code, so most likely
         // this declaration failed to instantiate. There's no point in
         // complaining further, since this is normal in invalid code.

diff  --git a/clang/lib/Sema/SemaTemplateVariadic.cpp b/clang/lib/Sema/SemaTemplateVariadic.cpp
index 1951aec3d17d..3c927f6b749f 100644
--- a/clang/lib/Sema/SemaTemplateVariadic.cpp
+++ b/clang/lib/Sema/SemaTemplateVariadic.cpp
@@ -368,8 +368,8 @@ Sema::DiagnoseUnexpandedParameterPacks(SourceLocation Loc,
       Locations.push_back(Unexpanded[I].second);
   }
 
-  auto DB = Diag(Loc, diag::err_unexpanded_parameter_pack)
-            << (int)UPPC << (int)Names.size();
+  DiagnosticBuilder DB = Diag(Loc, diag::err_unexpanded_parameter_pack)
+                         << (int)UPPC << (int)Names.size();
   for (size_t I = 0, E = std::min(Names.size(), (size_t)2); I != E; ++I)
     DB << Names[I];
 

diff  --git a/clang/lib/Sema/SemaType.cpp b/clang/lib/Sema/SemaType.cpp
index d9ff7c155ef9..d8ea9c037259 100644
--- a/clang/lib/Sema/SemaType.cpp
+++ b/clang/lib/Sema/SemaType.cpp
@@ -4133,8 +4133,7 @@ static FileID getNullabilityCompletenessCheckFileID(Sema &S,
 
 /// Creates a fix-it to insert a C-style nullability keyword at \p pointerLoc,
 /// taking into account whitespace before and after.
-template <typename DiagBuilderT>
-static void fixItNullability(Sema &S, DiagBuilderT &Diag,
+static void fixItNullability(Sema &S, DiagnosticBuilder &Diag,
                              SourceLocation PointerLoc,
                              NullabilityKind Nullability) {
   assert(PointerLoc.isValid());

diff  --git a/clang/test/SemaCUDA/deferred-oeverload.cu b/clang/test/SemaCUDA/deferred-oeverload.cu
deleted file mode 100644
index f89732b581ff..000000000000
--- a/clang/test/SemaCUDA/deferred-oeverload.cu
+++ /dev/null
@@ -1,78 +0,0 @@
-// RUN: %clang_cc1 -fcuda-is-device -fsyntax-only -verify=dev,com %s \
-// RUN:   -std=c++11 -fgpu-defer-diag
-// RUN: %clang_cc1 -fsyntax-only -verify=host,com %s \
-// RUN:   -std=c++11 -fgpu-defer-diag
-
-#include "Inputs/cuda.h"
-
-// When callee is called by a host function with integer arguments, there is an error for ambiguity.
-// It should be deferred since it involves wrong-sided candidates.
-__device__ void callee(int);
-__host__ void callee(float); // host-note {{candidate function}}
-__host__ void callee(double); // host-note {{candidate function}}
-
-// When callee2 is called by a device function without arguments, there is an error for 'no matching function'.
-// It should be deferred since it involves wrong-sided candidates.
-__host__ void callee2(); // dev-note{{candidate function not viable: call to __host__ function from __device__ function}}
-
-// When callee3 is called by a device function without arguments, there is an error for 'no matching function'.
-// It should be deferred since it involves wrong-sided candidates.
-__host__ void callee3(); // dev-note{{candidate function not viable: call to __host__ function from __device__ function}}
-__device__ void callee3(int); // dev-note{{candidate function not viable: requires 1 argument, but 0 were provided}}
-
-// When callee4 is called by a host or device function without arguments, there is an error for 'no matching function'.
-// It should be immediate since it involves no wrong-sided candidates (it is not a viable candiate due to signature).
-__host__ void callee4(int); // com-note 2{{candidate function not viable: requires 1 argument, but 0 were provided}}
-
-// When callee5 is called by a host function with integer arguments, there is an error for ambiguity.
-// It should be immediate since it involves no wrong-sided candidates.
-__host__ void callee5(float); // com-note {{candidate function}}
-__host__ void callee5(double); // com-note {{candidate function}}
-
-__host__ void hf() {
- callee(1); // host-error {{call to 'callee' is ambiguous}}
- callee2();
- callee3();
- callee4(); // com-error {{no matching function for call to 'callee4'}}
- callee5(1); // com-error {{call to 'callee5' is ambiguous}}
- undeclared_func(); // com-error {{use of undeclared identifier 'undeclared_func'}}
-}
-
-__device__ void df() {
- callee(1);
- callee2(); // dev-error {{no matching function for call to 'callee2'}}
- callee3(); // dev-error {{no matching function for call to 'callee3'}}
- callee4(); // com-error {{no matching function for call to 'callee4'}}
-}
-
-struct A { int x; typedef int isA; };
-struct B { int x; };
-
-// This function is invalid for A and B by SFINAE.
-// This fails to substitue for A but no diagnostic
-// should be emitted.
-template<typename T, typename T::foo* = nullptr>
-__host__ __device__ void sfinae(T t) { // com-note {{candidate template ignored: substitution failure [with T = B]}}
-  t.x = 1;
-}
-
-// This function is defined for A only by SFINAE.
-// Calling it with A should succeed, with B should fail.
-// The error should not be deferred since it happens in
-// file scope.
-
-template<typename T, typename T::isA* = nullptr>
-__host__ __device__ void sfinae(T t) { // com-note {{candidate template ignored: substitution failure [with T = B]}}
-  t.x = 1;
-}
-
-void test_sfinae() {
-  sfinae(A());
-  sfinae(B()); // com-error{{no matching function for call to 'sfinae'}}
-}
-
-// If a syntax error causes a function not declared, it cannot
-// be deferred.
-
-inline __host__ __device__ void bad_func() { // com-note {{to match this '{'}}
-// com-error {{expected '}'}}

diff  --git a/clang/test/TableGen/DiagnosticBase.inc b/clang/test/TableGen/DiagnosticBase.inc
index 291850e35364..6f5bd818aa9c 100644
--- a/clang/test/TableGen/DiagnosticBase.inc
+++ b/clang/test/TableGen/DiagnosticBase.inc
@@ -45,7 +45,6 @@ class TextSubstitution<string Text> {
   // diagnostics
   string Component = "";
   string CategoryName = "";
-  bit Deferrable = 0;
 }
 
 // Diagnostic Categories.  These can be applied to groups or individual
@@ -76,7 +75,6 @@ class Diagnostic<string text, DiagClass DC, Severity defaultmapping> {
   bit            AccessControl = 0;
   bit            WarningNoWerror = 0;
   bit            ShowInSystemHeader = 0;
-  bit            Deferrable = 0;
   Severity       DefaultSeverity = defaultmapping;
   DiagGroup      Group;
   string         CategoryName = "";
@@ -100,14 +98,6 @@ class SuppressInSystemHeader {
   bit ShowInSystemHeader = 0;
 }
 
-class Deferrable {
-  bit Deferrable = 1;
-}
-
-class NonDeferrable {
-  bit Deferrable = 0;
-}
-
 // FIXME: ExtWarn and Extension should also be SFINAEFailure by default.
 class Error<string str>     : Diagnostic<str, CLASS_ERROR, SEV_Error>, SFINAEFailure {
   bit ShowInSystemHeader = 1;

diff  --git a/clang/test/TableGen/deferred-diag.td b/clang/test/TableGen/deferred-diag.td
deleted file mode 100644
index bf95af31f587..000000000000
--- a/clang/test/TableGen/deferred-diag.td
+++ /dev/null
@@ -1,27 +0,0 @@
-// RUN: clang-tblgen -gen-clang-diags-defs -I%S %s -o - 2>&1 | \
-// RUN:    FileCheck --strict-whitespace %s
-include "DiagnosticBase.inc"
-
-// Test usage of Deferrable and NonDeferrable in diagnostics.
-
-def test_default : Error<"This error is non-deferrable by default">;
-// CHECK-DAG: DIAG(test_default, {{.*}}SFINAE_SubstitutionFailure, false, true, false, 0)
-
-def test_deferrable : Error<"This error is deferrable">, Deferrable;
-// CHECK-DAG: DIAG(test_deferrable, {{.*}} SFINAE_SubstitutionFailure, false, true, true, 0)
-
-def test_non_deferrable : Error<"This error is non-deferrable">, NonDeferrable;
-// CHECK-DAG: DIAG(test_non_deferrable, {{.*}} SFINAE_SubstitutionFailure, false, true, false, 0)
-
-let Deferrable = 1 in {
-
-def test_let : Error<"This error is deferrable by let">;
-// CHECK-DAG: DIAG(test_let, {{.*}} SFINAE_SubstitutionFailure, false, true, true, 0)
-
-// Make sure TextSubstitution is allowed in the let Deferrable block.
-def textsub : TextSubstitution<"%select{text1|text2}0">;
-
-def test_let2 : Error<"This error is deferrable by let %sub{textsub}0">;
-// CHECK-DAG: DIAG(test_let2, {{.*}} SFINAE_SubstitutionFailure, false, true, true, 0)
-
-}
\ No newline at end of file

diff  --git a/clang/tools/diagtool/DiagnosticNames.cpp b/clang/tools/diagtool/DiagnosticNames.cpp
index c54f81481a26..eddb99d1f57d 100644
--- a/clang/tools/diagtool/DiagnosticNames.cpp
+++ b/clang/tools/diagtool/DiagnosticNames.cpp
@@ -28,7 +28,7 @@ llvm::ArrayRef<DiagnosticRecord> diagtool::getBuiltinDiagnosticsByName() {
 // out of sync easily?
 static const DiagnosticRecord BuiltinDiagnosticsByID[] = {
 #define DIAG(ENUM,CLASS,DEFAULT_MAPPING,DESC,GROUP,               \
-             SFINAE,NOWERROR,SHOWINSYSHEADER,DEFER,CATEGORY)            \
+             SFINAE,NOWERROR,SHOWINSYSHEADER,CATEGORY)            \
   { #ENUM, diag::ENUM, STR_SIZE(#ENUM, uint8_t) },
 #include "clang/Basic/DiagnosticCommonKinds.inc"
 #include "clang/Basic/DiagnosticCrossTUKinds.inc"

diff  --git a/clang/utils/TableGen/ClangDiagnosticsEmitter.cpp b/clang/utils/TableGen/ClangDiagnosticsEmitter.cpp
index 430895d8425f..76d412203009 100644
--- a/clang/utils/TableGen/ClangDiagnosticsEmitter.cpp
+++ b/clang/utils/TableGen/ClangDiagnosticsEmitter.cpp
@@ -1294,11 +1294,6 @@ void clang::EmitClangDiagsDefs(RecordKeeper &Records, raw_ostream &OS,
     else
       OS << ", false";
 
-    if (R.getValueAsBit("Deferrable"))
-      OS << ", true";
-    else
-      OS << ", false";
-
     // Category number.
     OS << ", " << CategoryIDs.getID(getDiagnosticCategory(&R, DGParentMap));
     OS << ")\n";


        


More information about the cfe-commits mailing list