[clang-tools-extra] 772bd8a - Revert "[CUDA][HIP] Defer overloading resolution diagnostics for host device functions"

Roman Lebedev via cfe-commits cfe-commits at lists.llvm.org
Thu Sep 17 11:14:38 PDT 2020


Same comment about commit message, thanks

On Thu, Sep 17, 2020 at 8:56 PM Yaxun Liu via cfe-commits
<cfe-commits at lists.llvm.org> wrote:
>
>
> Author: Yaxun (Sam) Liu
> Date: 2020-09-17T13:55:31-04:00
> New Revision: 772bd8a7d99b8db899f594d393986e4b6cd85aa1
>
> URL: https://github.com/llvm/llvm-project/commit/772bd8a7d99b8db899f594d393986e4b6cd85aa1
> DIFF: https://github.com/llvm/llvm-project/commit/772bd8a7d99b8db899f594d393986e4b6cd85aa1.diff
>
> LOG: Revert "[CUDA][HIP] Defer overloading resolution diagnostics for host device functions"
>
> This reverts commit 7f1f89ec8d9944559042bb6d3b1132eabe3409de.
>
> This reverts commit 40df06cdafc010002fc9cfe1dda73d689b7d27a6.
>
> 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 20a5105fca4b..2e265e114191 100644
> --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
> +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td
> @@ -4060,8 +4060,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<
> @@ -4375,8 +4373,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 670bd8983265..4a22580a22ff 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.
> @@ -11799,11 +11673,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
> @@ -11816,16 +11763,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
> @@ -11840,10 +11786,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
> @@ -11856,14 +11801,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.
> @@ -12636,7 +12576,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
> @@ -12654,8 +12594,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 07e56fbcd611..8c7e63e06301 100644
> --- a/clang/lib/Basic/DiagnosticIDs.cpp
> +++ b/clang/lib/Basic/DiagnosticIDs.cpp
> @@ -42,7 +42,6 @@ struct StaticDiagInfoRec {
>    unsigned SFINAE : 2;
>    unsigned WarnNoWerror : 1;
>    unsigned WarnShowInSystemHeader : 1;
> -  unsigned Deferrable : 1;
>    unsigned Category : 6;
>
>    uint16_t OptionGroupIndex;
> @@ -97,10 +96,12 @@ VALIDATE_DIAG_SIZE(REFACTORING)
>
>  static const StaticDiagInfoRec StaticDiagInfo[] = {
>  #define DIAG(ENUM, CLASS, DEFAULT_SEVERITY, DESC, GROUP, SFINAE, NOWERROR,     \
> -             SHOWINSYSHEADER, DEFERRABLE, CATEGORY)                            \
> -  {diag::ENUM, DEFAULT_SEVERITY,         CLASS,      DiagnosticIDs::SFINAE,    \
> -   NOWERROR,   SHOWINSYSHEADER,          DEFERRABLE, CATEGORY,                 \
> -   GROUP,      STR_SIZE(DESC, uint16_t), DESC},
> +             SHOWINSYSHEADER, CATEGORY)                                        \
> +  {                                                                            \
> +    diag::ENUM, DEFAULT_SEVERITY, CLASS, DiagnosticIDs::SFINAE, NOWERROR,      \
> +        SHOWINSYSHEADER, CATEGORY, GROUP, STR_SIZE(DESC, uint16_t), DESC       \
> +  }                                                                            \
> +  ,
>  #include "clang/Basic/DiagnosticCommonKinds.inc"
>  #include "clang/Basic/DiagnosticDriverKinds.inc"
>  #include "clang/Basic/DiagnosticFrontendKinds.inc"
> @@ -252,12 +253,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 c62147f1dbac..20fb5a4d27e7 100644
> --- a/clang/lib/Sema/SemaDecl.cpp
> +++ b/clang/lib/Sema/SemaDecl.cpp
> @@ -14563,11 +14563,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.
> @@ -14619,7 +14619,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 c5072b5563e4..92f6141b6d38 100644
> --- a/clang/lib/Sema/SemaOpenMP.cpp
> +++ b/clang/lib/Sema/SemaOpenMP.cpp
> @@ -1888,27 +1888,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");
> @@ -1916,30 +1916,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 71c31fd7b836..95d110e754f4 100644
> --- a/clang/lib/Sema/SemaOverload.cpp
> +++ b/clang/lib/Sema/SemaOverload.cpp
> @@ -11522,18 +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.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 9c908878ab95..921d94036a2c 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 41e6d767c796..623d808b59f6 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";
>
>
>
> _______________________________________________
> cfe-commits mailing list
> cfe-commits at lists.llvm.org
> https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits


More information about the cfe-commits mailing list