Index: clang/include/clang/Basic/Diagnostic.td =================================================================== --- clang/include/clang/Basic/Diagnostic.td +++ clang/include/clang/Basic/Diagnostic.td @@ -45,6 +45,7 @@ // diagnostics string Component = ""; string CategoryName = ""; + bit Deferrable = 0; } // Diagnostic Categories. These can be applied to groups or individual @@ -83,6 +84,7 @@ bit AccessControl = 0; bit WarningNoWerror = 0; bit ShowInSystemHeader = 0; + bit Deferrable = 0; Severity DefaultSeverity = defaultmapping; DiagGroup Group; string CategoryName = ""; @@ -106,6 +108,14 @@ 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 : Diagnostic, SFINAEFailure { bit ShowInSystemHeader = 1; Index: clang/include/clang/Basic/DiagnosticAST.h =================================================================== --- clang/include/clang/Basic/DiagnosticAST.h +++ clang/include/clang/Basic/DiagnosticAST.h @@ -15,7 +15,7 @@ namespace diag { enum { #define DIAG(ENUM, FLAGS, DEFAULT_MAPPING, DESC, GROUP, SFINAE, NOWERROR, \ - SHOWINSYSHEADER, CATEGORY) \ + SHOWINSYSHEADER, DEFERRABLE, CATEGORY) \ ENUM, #define ASTSTART #include "clang/Basic/DiagnosticASTKinds.inc" Index: clang/include/clang/Basic/DiagnosticAnalysis.h =================================================================== --- clang/include/clang/Basic/DiagnosticAnalysis.h +++ clang/include/clang/Basic/DiagnosticAnalysis.h @@ -15,7 +15,7 @@ namespace diag { enum { #define DIAG(ENUM, FLAGS, DEFAULT_MAPPING, DESC, GROUP, SFINAE, NOWERROR, \ - SHOWINSYSHEADER, CATEGORY) \ + SHOWINSYSHEADER, DEFERRABLE, CATEGORY) \ ENUM, #define ANALYSISSTART #include "clang/Basic/DiagnosticAnalysisKinds.inc" Index: clang/include/clang/Basic/DiagnosticComment.h =================================================================== --- clang/include/clang/Basic/DiagnosticComment.h +++ clang/include/clang/Basic/DiagnosticComment.h @@ -15,7 +15,7 @@ namespace diag { enum { #define DIAG(ENUM, FLAGS, DEFAULT_MAPPING, DESC, GROUP, SFINAE, NOWERROR, \ - SHOWINSYSHEADER, CATEGORY) \ + SHOWINSYSHEADER, DEFERRABLE, CATEGORY) \ ENUM, #define COMMENTSTART #include "clang/Basic/DiagnosticCommentKinds.inc" Index: clang/include/clang/Basic/DiagnosticCrossTU.h =================================================================== --- clang/include/clang/Basic/DiagnosticCrossTU.h +++ clang/include/clang/Basic/DiagnosticCrossTU.h @@ -15,7 +15,7 @@ namespace diag { enum { #define DIAG(ENUM, FLAGS, DEFAULT_MAPPING, DESC, GROUP, SFINAE, NOWERROR, \ - SHOWINSYSHEADER, CATEGORY) \ + SHOWINSYSHEADER, DEFERRABLE, CATEGORY) \ ENUM, #define CROSSTUSTART #include "clang/Basic/DiagnosticCrossTUKinds.inc" Index: clang/include/clang/Basic/DiagnosticDriver.h =================================================================== --- clang/include/clang/Basic/DiagnosticDriver.h +++ clang/include/clang/Basic/DiagnosticDriver.h @@ -15,7 +15,7 @@ namespace diag { enum { #define DIAG(ENUM, FLAGS, DEFAULT_MAPPING, DESC, GROUP, SFINAE, NOWERROR, \ - SHOWINSYSHEADER, CATEGORY) \ + SHOWINSYSHEADER, DEFERRABLE, CATEGORY) \ ENUM, #define DRIVERSTART #include "clang/Basic/DiagnosticDriverKinds.inc" Index: clang/include/clang/Basic/DiagnosticFrontend.h =================================================================== --- clang/include/clang/Basic/DiagnosticFrontend.h +++ clang/include/clang/Basic/DiagnosticFrontend.h @@ -15,7 +15,7 @@ namespace diag { enum { #define DIAG(ENUM, FLAGS, DEFAULT_MAPPING, DESC, GROUP, SFINAE, NOWERROR, \ - SHOWINSYSHEADER, CATEGORY) \ + SHOWINSYSHEADER, DEFERRABLE, CATEGORY) \ ENUM, #define FRONTENDSTART #include "clang/Basic/DiagnosticFrontendKinds.inc" Index: clang/include/clang/Basic/DiagnosticIDs.h =================================================================== --- clang/include/clang/Basic/DiagnosticIDs.h +++ clang/include/clang/Basic/DiagnosticIDs.h @@ -64,8 +64,9 @@ // Get typedefs for common diagnostics. enum { -#define DIAG(ENUM,FLAGS,DEFAULT_MAPPING,DESC,GROUP,\ - SFINAE,CATEGORY,NOWERROR,SHOWINSYSHEADER) ENUM, +#define DIAG(ENUM, FLAGS, DEFAULT_MAPPING, DESC, GROUP, SFINAE, CATEGORY, \ + NOWERROR, SHOWINSYSHEADER, DEFFERABLE) \ + ENUM, #define COMMONSTART #include "clang/Basic/DiagnosticCommonKinds.inc" NUM_BUILTIN_COMMON_DIAGNOSTICS @@ -280,6 +281,13 @@ /// 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 Index: clang/include/clang/Basic/DiagnosticLex.h =================================================================== --- clang/include/clang/Basic/DiagnosticLex.h +++ clang/include/clang/Basic/DiagnosticLex.h @@ -15,7 +15,7 @@ namespace diag { enum { #define DIAG(ENUM, FLAGS, DEFAULT_MAPPING, DESC, GROUP, SFINAE, NOWERROR, \ - SHOWINSYSHEADER, CATEGORY) \ + SHOWINSYSHEADER, DEFERRABLE, CATEGORY) \ ENUM, #define LEXSTART #include "clang/Basic/DiagnosticLexKinds.inc" Index: clang/include/clang/Basic/DiagnosticParse.h =================================================================== --- clang/include/clang/Basic/DiagnosticParse.h +++ clang/include/clang/Basic/DiagnosticParse.h @@ -15,7 +15,7 @@ namespace diag { enum { #define DIAG(ENUM, FLAGS, DEFAULT_MAPPING, DESC, GROUP, SFINAE, NOWERROR, \ - SHOWINSYSHEADER, CATEGORY) \ + SHOWINSYSHEADER, DEFERRABLE, CATEGORY) \ ENUM, #define PARSESTART #include "clang/Basic/DiagnosticParseKinds.inc" Index: clang/include/clang/Basic/DiagnosticRefactoring.h =================================================================== --- clang/include/clang/Basic/DiagnosticRefactoring.h +++ clang/include/clang/Basic/DiagnosticRefactoring.h @@ -15,7 +15,7 @@ namespace diag { enum { #define DIAG(ENUM, FLAGS, DEFAULT_MAPPING, DESC, GROUP, SFINAE, NOWERROR, \ - SHOWINSYSHEADER, CATEGORY) \ + SHOWINSYSHEADER, DEFERRABLE, CATEGORY) \ ENUM, #define REFACTORINGSTART #include "clang/Basic/DiagnosticRefactoringKinds.inc" Index: clang/include/clang/Basic/DiagnosticSema.h =================================================================== --- clang/include/clang/Basic/DiagnosticSema.h +++ clang/include/clang/Basic/DiagnosticSema.h @@ -15,7 +15,7 @@ namespace diag { enum { #define DIAG(ENUM, FLAGS, DEFAULT_MAPPING, DESC, GROUP, SFINAE, NOWERROR, \ - SHOWINSYSHEADER, CATEGORY) \ + SHOWINSYSHEADER, DEFERRABLE, CATEGORY) \ ENUM, #define SEMASTART #include "clang/Basic/DiagnosticSemaKinds.inc" Index: clang/include/clang/Basic/DiagnosticSemaKinds.td =================================================================== --- clang/include/clang/Basic/DiagnosticSemaKinds.td +++ clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -4050,6 +4050,8 @@ "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< @@ -4363,6 +4365,8 @@ 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" Index: clang/include/clang/Basic/DiagnosticSerialization.h =================================================================== --- clang/include/clang/Basic/DiagnosticSerialization.h +++ clang/include/clang/Basic/DiagnosticSerialization.h @@ -15,7 +15,7 @@ namespace diag { enum { #define DIAG(ENUM, FLAGS, DEFAULT_MAPPING, DESC, GROUP, SFINAE, NOWERROR, \ - SHOWINSYSHEADER, CATEGORY) \ + SHOWINSYSHEADER, DEFERRABLE, CATEGORY) \ ENUM, #define SERIALIZATIONSTART #include "clang/Basic/DiagnosticSerializationKinds.inc" Index: clang/include/clang/Basic/LangOptions.def =================================================================== --- clang/include/clang/Basic/LangOptions.def +++ clang/include/clang/Basic/LangOptions.def @@ -241,6 +241,7 @@ 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 all semantic diagnostic messages in host device functions for CUDA/HIP") LANGOPT(SYCL , 1, 0, "SYCL") LANGOPT(SYCLIsDevice , 1, 0, "Generate code for SYCL device") Index: clang/include/clang/Driver/Options.td =================================================================== --- clang/include/clang/Driver/Options.td +++ clang/include/clang/Driver/Options.td @@ -662,6 +662,9 @@ "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", " all semantic diagnostics in host device functions" + " 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">; Index: clang/include/clang/Sema/Sema.h =================================================================== --- clang/include/clang/Sema/Sema.h +++ clang/include/clang/Sema/Sema.h @@ -1462,28 +1462,31 @@ /// template instantiation stacks. /// /// This class provides a wrapper around the basic DiagnosticBuilder - /// class that emits diagnostics. SemaDiagnosticBuilder is + /// class that emits diagnostics. ImmediateDiagBuilder 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 SemaDiagnosticBuilder : public DiagnosticBuilder { + class ImmediateDiagBuilder : public DiagnosticBuilder { Sema &SemaRef; unsigned DiagID; public: - SemaDiagnosticBuilder(DiagnosticBuilder &DB, Sema &SemaRef, unsigned DiagID) - : DiagnosticBuilder(DB), SemaRef(SemaRef), DiagID(DiagID) { } + const static bool IsDiagBuilder = false; + 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) {} // 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 ~SemaDiagnosticBuilder is a safe no-op + // source anyway, so the user-defined ~ImmediateDiagBuilder is a safe no-op // in that case anwyay. - SemaDiagnosticBuilder(const SemaDiagnosticBuilder&) = default; + ImmediateDiagBuilder(const ImmediateDiagBuilder &) = default; - ~SemaDiagnosticBuilder() { + ~ImmediateDiagBuilder() { // If we aren't active, there is nothing to do. if (!isActive()) return; @@ -1504,28 +1507,138 @@ } /// Teach operator<< to produce an object of the correct type. - template - friend const SemaDiagnosticBuilder &operator<<( - const SemaDiagnosticBuilder &Diag, const T &Value) { + template + friend const ImmediateDiagBuilder & + operator<<(const ImmediateDiagBuilder &Diag, const T &Value) { const DiagnosticBuilder &BaseDiag = Diag; BaseDiag << Value; return Diag; } - const static bool IsDiagBuilder = false; }; + /// 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 + 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; + } + + 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 ImmediateDiag; + llvm::Optional 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) { - DiagnosticBuilder DB = Diags.Report(Loc, DiagID); - return SemaDiagnosticBuilder(DB, *this, DiagID); - } + SemaDiagnosticBuilder Diag(SourceLocation Loc, unsigned DiagID, + bool DeferHint = false); /// Emit a partial diagnostic. - SemaDiagnosticBuilder Diag(SourceLocation Loc, const PartialDiagnostic& PD); + SemaDiagnosticBuilder Diag(SourceLocation Loc, const PartialDiagnostic &PD, + bool DeferHint = false); /// 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. @@ -11659,84 +11772,11 @@ /* Caller = */ FunctionDeclAndLoc> DeviceKnownEmittedFns; - /// Diagnostic builder for CUDA/OpenMP devices errors which may or may not be - /// deferred. - /// - /// In CUDA, there exist constructs (e.g. variable-length arrays, try/catch) - /// which are not allowed to appear inside __device__ functions and are - /// allowed to appear in __host__ __device__ functions only if the host+device - /// function is never codegen'ed. - /// - /// To handle this, we use the notion of "deferred diagnostics", where we - /// attach a diagnostic to a FunctionDecl that's emitted iff it's codegen'ed. - /// - /// This class lets you emit either a regular diagnostic, a deferred - /// diagnostic, or no diagnostic at all, according to an argument you pass to - /// its constructor, thus simplifying the process of creating these "maybe - /// deferred" diagnostics. - class 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 - 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 ImmediateDiag; - llvm::Optional PartialDiagId; - }; - - /// Creates a DeviceDiagBuilder that emits the diagnostic if the current context - /// is "used as device code". + /// Creates a SemaDiagnosticBuilder 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 __host__ function, does not emit any diagnostics + /// unless \p EmitOnBothSides is true. /// - If CurContext is a __device__ or __global__ function, emits the /// diagnostics immediately. /// - If CurContext is a __host__ __device__ function and we are compiling for @@ -11749,15 +11789,16 @@ /// if (CUDADiagIfDeviceCode(Loc, diag::err_cuda_vla) << CurrentCUDATarget()) /// return ExprError(); /// // Otherwise, continue parsing as normal. - DeviceDiagBuilder CUDADiagIfDeviceCode(SourceLocation Loc, unsigned DiagID); + SemaDiagnosticBuilder CUDADiagIfDeviceCode(SourceLocation Loc, + unsigned DiagID); - /// Creates a DeviceDiagBuilder that emits the diagnostic if the current context - /// is "used as host code". + /// Creates a SemaDiagnosticBuilder that emits the diagnostic if the current + /// context is "used as host code". /// /// Same as CUDADiagIfDeviceCode, with "host" and "device" switched. - DeviceDiagBuilder CUDADiagIfHostCode(SourceLocation Loc, unsigned DiagID); + SemaDiagnosticBuilder CUDADiagIfHostCode(SourceLocation Loc, unsigned DiagID); - /// Creates a DeviceDiagBuilder that emits the diagnostic if the current + /// Creates a SemaDiagnosticBuilder 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 @@ -11772,9 +11813,10 @@ /// if (diagIfOpenMPDeviceCode(Loc, diag::err_vla_unsupported)) /// return ExprError(); /// // Otherwise, continue parsing as normal. - DeviceDiagBuilder diagIfOpenMPDeviceCode(SourceLocation Loc, unsigned DiagID); + SemaDiagnosticBuilder diagIfOpenMPDeviceCode(SourceLocation Loc, + unsigned DiagID); - /// Creates a DeviceDiagBuilder that emits the diagnostic if the current + /// Creates a SemaDiagnosticBuilder 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 @@ -11787,9 +11829,14 @@ /// if (diagIfOpenMPHostode(Loc, diag::err_vla_unsupported)) /// return ExprError(); /// // Otherwise, continue parsing as normal. - DeviceDiagBuilder diagIfOpenMPHostCode(SourceLocation Loc, unsigned DiagID); + SemaDiagnosticBuilder diagIfOpenMPHostCode(SourceLocation Loc, + unsigned DiagID); - DeviceDiagBuilder targetDiag(SourceLocation Loc, unsigned DiagID); + SemaDiagnosticBuilder targetDiag(SourceLocation Loc, unsigned DiagID); + SemaDiagnosticBuilder targetDiag(SourceLocation Loc, + const PartialDiagnostic &PD) { + return targetDiag(Loc, PD.getDiagID()) << PD; + } /// Check if the expression is allowed to be used in expressions for the /// offloading devices. @@ -12561,7 +12608,7 @@ ConstructorDestructor, BuiltinFunction }; - /// Creates a DeviceDiagBuilder that emits the diagnostic if the current + /// Creates a SemaDiagnosticBuilder 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 @@ -12579,7 +12626,8 @@ /// if (!S.Context.getTargetInfo().hasFloat128Type() && /// S.getLangOpts().SYCLIsDevice) /// SYCLDiagIfDeviceCode(Loc, diag::err_type_unsupported) << "__float128"; - DeviceDiagBuilder SYCLDiagIfDeviceCode(SourceLocation Loc, unsigned DiagID); + SemaDiagnosticBuilder SYCLDiagIfDeviceCode(SourceLocation Loc, + unsigned DiagID); /// Check whether we're allowed to call Callee from the current context. /// Index: clang/lib/Basic/DiagnosticIDs.cpp =================================================================== --- clang/lib/Basic/DiagnosticIDs.cpp +++ clang/lib/Basic/DiagnosticIDs.cpp @@ -42,6 +42,7 @@ unsigned SFINAE : 2; unsigned WarnNoWerror : 1; unsigned WarnShowInSystemHeader : 1; + unsigned Deferrable : 1; unsigned Category : 6; uint16_t OptionGroupIndex; @@ -96,12 +97,10 @@ static const StaticDiagInfoRec StaticDiagInfo[] = { #define DIAG(ENUM, CLASS, DEFAULT_SEVERITY, DESC, GROUP, SFINAE, NOWERROR, \ - SHOWINSYSHEADER, CATEGORY) \ - { \ - diag::ENUM, DEFAULT_SEVERITY, CLASS, DiagnosticIDs::SFINAE, NOWERROR, \ - SHOWINSYSHEADER, CATEGORY, GROUP, STR_SIZE(DESC, uint16_t), DESC \ - } \ - , + SHOWINSYSHEADER, DEFERRABLE, CATEGORY) \ + {diag::ENUM, DEFAULT_SEVERITY, CLASS, DiagnosticIDs::SFINAE, \ + NOWERROR, SHOWINSYSHEADER, DEFERRABLE, CATEGORY, \ + GROUP, STR_SIZE(DESC, uint16_t), DESC}, #include "clang/Basic/DiagnosticCommonKinds.inc" #include "clang/Basic/DiagnosticDriverKinds.inc" #include "clang/Basic/DiagnosticFrontendKinds.inc" @@ -253,6 +252,12 @@ 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) { Index: clang/lib/Driver/ToolChains/Cuda.cpp =================================================================== --- clang/lib/Driver/ToolChains/Cuda.cpp +++ clang/lib/Driver/ToolChains/Cuda.cpp @@ -634,6 +634,10 @@ 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)) Index: clang/lib/Driver/ToolChains/HIP.cpp =================================================================== --- clang/lib/Driver/ToolChains/HIP.cpp +++ clang/lib/Driver/ToolChains/HIP.cpp @@ -266,6 +266,10 @@ 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 Index: clang/lib/Frontend/CompilerInvocation.cpp =================================================================== --- clang/lib/Frontend/CompilerInvocation.cpp +++ clang/lib/Frontend/CompilerInvocation.cpp @@ -2628,6 +2628,9 @@ 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; Index: clang/lib/Sema/AnalysisBasedWarnings.cpp =================================================================== --- clang/lib/Sema/AnalysisBasedWarnings.cpp +++ clang/lib/Sema/AnalysisBasedWarnings.cpp @@ -2096,7 +2096,7 @@ if (cast(D)->isDependentContext()) return; - if (Diags.hasUncompilableErrorOccurred()) { + if (S.hasUncompilableErrorOccurred()) { // Flush out any possibly unreachable diagnostics. flushDiagnostics(S, fscope); return; Index: clang/lib/Sema/Sema.cpp =================================================================== --- clang/lib/Sema/Sema.cpp +++ clang/lib/Sema/Sema.cpp @@ -1435,11 +1435,24 @@ } Sema::SemaDiagnosticBuilder -Sema::Diag(SourceLocation Loc, const PartialDiagnostic& PD) { - SemaDiagnosticBuilder Builder(Diag(Loc, PD.getDiagID())); - PD.Emit(Builder); +Sema::Diag(SourceLocation Loc, const PartialDiagnostic &PD, bool DeferHint) { + return Diag(Loc, PD.getDiagID(), DeferHint) << PD; +} - return Builder; +bool Sema::hasUncompilableErrorOccurred() const { + if (getDiagnostics().hasUncompilableErrorOccurred()) + return true; + auto *FD = dyn_cast(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; } // Print notes showing how we can reach FD starting from an a priori @@ -1652,9 +1665,9 @@ // until we discover that the function is known-emitted, at which point we take // it out of this map and emit the diagnostic. -Sema::DeviceDiagBuilder::DeviceDiagBuilder(Kind K, SourceLocation Loc, - unsigned DiagID, FunctionDecl *Fn, - Sema &S) +Sema::SemaDiagnosticBuilder::SemaDiagnosticBuilder(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) { @@ -1662,7 +1675,8 @@ break; case K_Immediate: case K_ImmediateWithCallStack: - ImmediateDiag.emplace(S.Diag(Loc, DiagID)); + ImmediateDiag.emplace( + ImmediateDiagBuilder(S.Diags.Report(Loc, DiagID), S, DiagID)); break; case K_Deferred: assert(Fn && "Must have a function to attach the deferred diag to."); @@ -1673,7 +1687,7 @@ } } -Sema::DeviceDiagBuilder::DeviceDiagBuilder(DeviceDiagBuilder &&D) +Sema::SemaDiagnosticBuilder::SemaDiagnosticBuilder(SemaDiagnosticBuilder &&D) : S(D.S), Loc(D.Loc), DiagID(D.DiagID), Fn(D.Fn), ShowCallStack(D.ShowCallStack), ImmediateDiag(D.ImmediateDiag), PartialDiagId(D.PartialDiagId) { @@ -1683,7 +1697,7 @@ D.PartialDiagId.reset(); } -Sema::DeviceDiagBuilder::~DeviceDiagBuilder() { +Sema::SemaDiagnosticBuilder::~SemaDiagnosticBuilder() { if (ImmediateDiag) { // Emit our diagnostic and, if it was a warning or error, output a callstack // if Fn isn't a priori known-emitted. @@ -1698,7 +1712,8 @@ } } -Sema::DeviceDiagBuilder Sema::targetDiag(SourceLocation Loc, unsigned DiagID) { +Sema::SemaDiagnosticBuilder Sema::targetDiag(SourceLocation Loc, + unsigned DiagID) { if (LangOpts.OpenMP) return LangOpts.OpenMPIsDevice ? diagIfOpenMPDeviceCode(Loc, DiagID) : diagIfOpenMPHostCode(Loc, DiagID); @@ -1709,8 +1724,32 @@ if (getLangOpts().SYCLIsDevice) return SYCLDiagIfDeviceCode(Loc, DiagID); - return DeviceDiagBuilder(DeviceDiagBuilder::K_Immediate, Loc, DiagID, - getCurFunctionDecl(), *this); + 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; } void Sema::checkDeviceDecl(const ValueDecl *D, SourceLocation Loc) { Index: clang/lib/Sema/SemaAttr.cpp =================================================================== --- clang/lib/Sema/SemaAttr.cpp +++ clang/lib/Sema/SemaAttr.cpp @@ -380,8 +380,8 @@ // The user might have already reset the alignment, so suggest replacing // the reset with a pop. if (IsInnermost && PackStack.CurrentValue == PackStack.DefaultValue) { - DiagnosticBuilder DB = Diag(PackStack.CurrentPragmaLocation, - diag::note_pragma_pack_pop_instead_reset); + auto DB = Diag(PackStack.CurrentPragmaLocation, + diag::note_pragma_pack_pop_instead_reset); SourceLocation FixItLoc = Lexer::findLocationAfterToken( PackStack.CurrentPragmaLocation, tok::l_paren, SourceMgr, LangOpts, /*SkipTrailing=*/false); Index: clang/lib/Sema/SemaCUDA.cpp =================================================================== --- clang/lib/Sema/SemaCUDA.cpp +++ clang/lib/Sema/SemaCUDA.cpp @@ -639,58 +639,63 @@ } } -Sema::DeviceDiagBuilder Sema::CUDADiagIfDeviceCode(SourceLocation Loc, - unsigned DiagID) { +Sema::SemaDiagnosticBuilder Sema::CUDADiagIfDeviceCode(SourceLocation Loc, + unsigned DiagID) { assert(getLangOpts().CUDA && "Should only be called during CUDA compilation"); - DeviceDiagBuilder::Kind DiagKind = [this] { + SemaDiagnosticBuilder::Kind DiagKind = [&] { + if (!isa(CurContext)) + return SemaDiagnosticBuilder::K_Immediate; switch (CurrentCUDATarget()) { case CFT_Global: case CFT_Device: - return DeviceDiagBuilder::K_Immediate; + return SemaDiagnosticBuilder::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 (getEmissionStatus(cast(CurContext)) == - FunctionEmissionStatus::Emitted) - ? DeviceDiagBuilder::K_ImmediateWithCallStack - : DeviceDiagBuilder::K_Deferred; - } - return DeviceDiagBuilder::K_Nop; - + if (!getLangOpts().CUDAIsDevice) + return SemaDiagnosticBuilder::K_Nop; + if (IsLastErrorImmediate && Diags.getDiagnosticIDs()->isBuiltinNote(DiagID)) + return SemaDiagnosticBuilder::K_Immediate; + return (getEmissionStatus(cast(CurContext)) == + FunctionEmissionStatus::Emitted) + ? SemaDiagnosticBuilder::K_ImmediateWithCallStack + : SemaDiagnosticBuilder::K_Deferred; default: - return DeviceDiagBuilder::K_Nop; + return SemaDiagnosticBuilder::K_Nop; } }(); - return DeviceDiagBuilder(DiagKind, Loc, DiagID, - dyn_cast(CurContext), *this); + return SemaDiagnosticBuilder(DiagKind, Loc, DiagID, + dyn_cast(CurContext), *this); } -Sema::DeviceDiagBuilder Sema::CUDADiagIfHostCode(SourceLocation Loc, - unsigned DiagID) { +Sema::SemaDiagnosticBuilder Sema::CUDADiagIfHostCode(SourceLocation Loc, + unsigned DiagID) { assert(getLangOpts().CUDA && "Should only be called during CUDA compilation"); - DeviceDiagBuilder::Kind DiagKind = [this] { + SemaDiagnosticBuilder::Kind DiagKind = [&] { + if (!isa(CurContext)) + return SemaDiagnosticBuilder::K_Immediate; switch (CurrentCUDATarget()) { case CFT_Host: - return DeviceDiagBuilder::K_Immediate; + return SemaDiagnosticBuilder::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 DeviceDiagBuilder::K_Nop; - + return SemaDiagnosticBuilder::K_Nop; + if (IsLastErrorImmediate && Diags.getDiagnosticIDs()->isBuiltinNote(DiagID)) + return SemaDiagnosticBuilder::K_Immediate; return (getEmissionStatus(cast(CurContext)) == FunctionEmissionStatus::Emitted) - ? DeviceDiagBuilder::K_ImmediateWithCallStack - : DeviceDiagBuilder::K_Deferred; + ? SemaDiagnosticBuilder::K_ImmediateWithCallStack + : SemaDiagnosticBuilder::K_Deferred; default: - return DeviceDiagBuilder::K_Nop; + return SemaDiagnosticBuilder::K_Nop; } }(); - return DeviceDiagBuilder(DiagKind, Loc, DiagID, - dyn_cast(CurContext), *this); + return SemaDiagnosticBuilder(DiagKind, Loc, DiagID, + dyn_cast(CurContext), *this); } bool Sema::CheckCUDACall(SourceLocation Loc, FunctionDecl *Callee) { @@ -711,8 +716,8 @@ // Otherwise, mark the call in our call graph so we can traverse it later. bool CallerKnownEmitted = getEmissionStatus(Caller) == FunctionEmissionStatus::Emitted; - DeviceDiagBuilder::Kind DiagKind = [this, Caller, Callee, - CallerKnownEmitted] { + SemaDiagnosticBuilder::Kind DiagKind = [this, Caller, Callee, + CallerKnownEmitted] { switch (IdentifyCUDAPreference(Caller, Callee)) { case CFP_Never: case CFP_WrongSide: @@ -720,14 +725,15 @@ // 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 ? DeviceDiagBuilder::K_ImmediateWithCallStack - : DeviceDiagBuilder::K_Deferred; + return CallerKnownEmitted + ? SemaDiagnosticBuilder::K_ImmediateWithCallStack + : SemaDiagnosticBuilder::K_Deferred; default: - return DeviceDiagBuilder::K_Nop; + return SemaDiagnosticBuilder::K_Nop; } }(); - if (DiagKind == DeviceDiagBuilder::K_Nop) + if (DiagKind == SemaDiagnosticBuilder::K_Nop) return true; // Avoid emitting this error twice for the same location. Using a hashtable @@ -737,14 +743,14 @@ if (!LocsWithCUDACallDiags.insert({Caller, Loc}).second) return true; - DeviceDiagBuilder(DiagKind, Loc, diag::err_ref_bad_target, Caller, *this) + SemaDiagnosticBuilder(DiagKind, Loc, diag::err_ref_bad_target, Caller, *this) << IdentifyCUDATarget(Callee) << Callee << IdentifyCUDATarget(Caller); if (!Callee->getBuiltinID()) - DeviceDiagBuilder(DiagKind, Callee->getLocation(), diag::note_previous_decl, - Caller, *this) + SemaDiagnosticBuilder(DiagKind, Callee->getLocation(), + diag::note_previous_decl, Caller, *this) << Callee; - return DiagKind != DeviceDiagBuilder::K_Immediate && - DiagKind != DeviceDiagBuilder::K_ImmediateWithCallStack; + return DiagKind != SemaDiagnosticBuilder::K_Immediate && + DiagKind != SemaDiagnosticBuilder::K_ImmediateWithCallStack; } // Check the wrong-sided reference capture of lambda for CUDA/HIP. @@ -781,14 +787,14 @@ bool ShouldCheck = CalleeIsDevice && CallerIsHost; if (!ShouldCheck || !Capture.isReferenceCapture()) return; - auto DiagKind = DeviceDiagBuilder::K_Deferred; + auto DiagKind = SemaDiagnosticBuilder::K_Deferred; if (Capture.isVariableCapture()) { - DeviceDiagBuilder(DiagKind, Capture.getLocation(), - diag::err_capture_bad_target, Callee, *this) + SemaDiagnosticBuilder(DiagKind, Capture.getLocation(), + diag::err_capture_bad_target, Callee, *this) << Capture.getVariable(); } else if (Capture.isThisCapture()) { - DeviceDiagBuilder(DiagKind, Capture.getLocation(), - diag::err_capture_bad_target_this_ptr, Callee, *this); + SemaDiagnosticBuilder(DiagKind, Capture.getLocation(), + diag::err_capture_bad_target_this_ptr, Callee, *this); } return; } Index: clang/lib/Sema/SemaDecl.cpp =================================================================== --- clang/lib/Sema/SemaDecl.cpp +++ clang/lib/Sema/SemaDecl.cpp @@ -14542,11 +14542,11 @@ // 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 (getDiagnostics().hasUncompilableErrorOccurred() || + if (hasUncompilableErrorOccurred() || getDiagnostics().getSuppressAllDiagnostics()) { DiscardCleanupsInEvaluationContext(); } - if (!getDiagnostics().hasUncompilableErrorOccurred() && + if (!hasUncompilableErrorOccurred() && !isa(dcl)) { // Since the body is valid, issue any analysis-based warnings that are // enabled. @@ -14598,7 +14598,7 @@ // 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 (getDiagnostics().hasUncompilableErrorOccurred()) { + if (hasUncompilableErrorOccurred()) { DiscardCleanupsInEvaluationContext(); } Index: clang/lib/Sema/SemaExprObjC.cpp =================================================================== --- clang/lib/Sema/SemaExprObjC.cpp +++ clang/lib/Sema/SemaExprObjC.cpp @@ -2445,8 +2445,8 @@ SourceManager &SM = S.SourceMgr; edit::Commit ECommit(SM, S.LangOpts); if (refactor(Msg,*S.NSAPIObj, ECommit)) { - DiagnosticBuilder Builder = S.Diag(MsgLoc, DiagID) - << Msg->getSelector() << Msg->getSourceRange(); + auto 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,9 +3139,8 @@ if (ReceiverType->isObjCClassType() && !isImplicit && !(Receiver->isObjCSelfExpr() && getLangOpts().ObjCAutoRefCount)) { { - DiagnosticBuilder Builder = - Diag(Receiver->getExprLoc(), - diag::err_messaging_class_with_direct_method); + auto Builder = Diag(Receiver->getExprLoc(), + diag::err_messaging_class_with_direct_method); if (Receiver->isObjCSelfExpr()) { Builder.AddFixItHint(FixItHint::CreateReplacement( RecRange, Method->getClassInterface()->getName())); @@ -3153,7 +3152,7 @@ if (SuperLoc.isValid()) { { - DiagnosticBuilder Builder = + auto Builder = Diag(SuperLoc, diag::err_messaging_super_with_direct_method); if (ReceiverType->isObjCClassType()) { Builder.AddFixItHint(FixItHint::CreateReplacement( @@ -3736,15 +3735,11 @@ return LookupName(R, TUScope, false); } -static void addFixitForObjCARCConversion(Sema &S, - DiagnosticBuilder &DiagB, - Sema::CheckedConversionKind CCK, - SourceLocation afterLParen, - QualType castType, - Expr *castExpr, - Expr *realCast, - const char *bridgeKeyword, - const char *CFBridgeName) { +template +static void addFixitForObjCARCConversion( + Sema &S, DiagBuilderT &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: @@ -3921,9 +3916,9 @@ assert(CreateRule != ACC_bottom && "This cast should already be accepted."); if (CreateRule != ACC_plusOne) { - DiagnosticBuilder DiagB = - (CCK != Sema::CCK_OtherCast) ? S.Diag(noteLoc, diag::note_arc_bridge) - : S.Diag(noteLoc, diag::note_arc_cstyle_bridge); + auto 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 ", @@ -3931,12 +3926,12 @@ } if (CreateRule != ACC_plusZero) { - 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; + 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; addFixitForObjCARCConversion(S, DiagB, CCK, afterLParen, castType, castExpr, realCast, "__bridge_transfer ", @@ -3962,21 +3957,21 @@ assert(CreateRule != ACC_bottom && "This cast should already be accepted."); if (CreateRule != ACC_plusOne) { - DiagnosticBuilder DiagB = - (CCK != Sema::CCK_OtherCast) ? S.Diag(noteLoc, diag::note_arc_bridge) - : S.Diag(noteLoc, diag::note_arc_cstyle_bridge); + auto 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) { - 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; + 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; addFixitForObjCARCConversion(S, DiagB, CCK, afterLParen, castType, castExpr, realCast, "__bridge_retained ", Index: clang/lib/Sema/SemaOpenMP.cpp =================================================================== --- clang/lib/Sema/SemaOpenMP.cpp +++ clang/lib/Sema/SemaOpenMP.cpp @@ -1888,27 +1888,27 @@ }; } // anonymous namespace -Sema::DeviceDiagBuilder Sema::diagIfOpenMPDeviceCode(SourceLocation Loc, - unsigned DiagID) { +Sema::SemaDiagnosticBuilder Sema::diagIfOpenMPDeviceCode(SourceLocation Loc, + unsigned DiagID) { assert(LangOpts.OpenMP && LangOpts.OpenMPIsDevice && "Expected OpenMP device compilation."); FunctionDecl *FD = getCurFunctionDecl(); - DeviceDiagBuilder::Kind Kind = DeviceDiagBuilder::K_Nop; + SemaDiagnosticBuilder::Kind Kind = SemaDiagnosticBuilder::K_Nop; if (FD) { FunctionEmissionStatus FES = getEmissionStatus(FD); switch (FES) { case FunctionEmissionStatus::Emitted: - Kind = DeviceDiagBuilder::K_Immediate; + Kind = SemaDiagnosticBuilder::K_Immediate; break; case FunctionEmissionStatus::Unknown: Kind = isOpenMPDeviceDelayedContext(*this) - ? DeviceDiagBuilder::K_Deferred - : DeviceDiagBuilder::K_Immediate; + ? SemaDiagnosticBuilder::K_Deferred + : SemaDiagnosticBuilder::K_Immediate; break; case FunctionEmissionStatus::TemplateDiscarded: case FunctionEmissionStatus::OMPDiscarded: - Kind = DeviceDiagBuilder::K_Nop; + Kind = SemaDiagnosticBuilder::K_Nop; break; case FunctionEmissionStatus::CUDADiscarded: llvm_unreachable("CUDADiscarded unexpected in OpenMP device compilation"); @@ -1916,30 +1916,30 @@ } } - return DeviceDiagBuilder(Kind, Loc, DiagID, getCurFunctionDecl(), *this); + return SemaDiagnosticBuilder(Kind, Loc, DiagID, getCurFunctionDecl(), *this); } -Sema::DeviceDiagBuilder Sema::diagIfOpenMPHostCode(SourceLocation Loc, - unsigned DiagID) { +Sema::SemaDiagnosticBuilder Sema::diagIfOpenMPHostCode(SourceLocation Loc, + unsigned DiagID) { assert(LangOpts.OpenMP && !LangOpts.OpenMPIsDevice && "Expected OpenMP host compilation."); FunctionEmissionStatus FES = getEmissionStatus(getCurFunctionDecl()); - DeviceDiagBuilder::Kind Kind = DeviceDiagBuilder::K_Nop; + SemaDiagnosticBuilder::Kind Kind = SemaDiagnosticBuilder::K_Nop; switch (FES) { case FunctionEmissionStatus::Emitted: - Kind = DeviceDiagBuilder::K_Immediate; + Kind = SemaDiagnosticBuilder::K_Immediate; break; case FunctionEmissionStatus::Unknown: - Kind = DeviceDiagBuilder::K_Deferred; + Kind = SemaDiagnosticBuilder::K_Deferred; break; case FunctionEmissionStatus::TemplateDiscarded: case FunctionEmissionStatus::OMPDiscarded: case FunctionEmissionStatus::CUDADiscarded: - Kind = DeviceDiagBuilder::K_Nop; + Kind = SemaDiagnosticBuilder::K_Nop; break; } - return DeviceDiagBuilder(Kind, Loc, DiagID, getCurFunctionDecl(), *this); + return SemaDiagnosticBuilder(Kind, Loc, DiagID, getCurFunctionDecl(), *this); } static OpenMPDefaultmapClauseKind Index: clang/lib/Sema/SemaOverload.cpp =================================================================== --- clang/lib/Sema/SemaOverload.cpp +++ clang/lib/Sema/SemaOverload.cpp @@ -11528,9 +11528,18 @@ StringRef Opc, SourceLocation OpLoc, llvm::function_ref 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); + S.Diag(PD.first, PD.second, DeferHint); NoteCandidates(S, Args, Cands, Opc, OpLoc); Index: clang/lib/Sema/SemaSYCL.cpp =================================================================== --- clang/lib/Sema/SemaSYCL.cpp +++ clang/lib/Sema/SemaSYCL.cpp @@ -17,19 +17,19 @@ // SYCL device specific diagnostics implementation // ----------------------------------------------------------------------------- -Sema::DeviceDiagBuilder Sema::SYCLDiagIfDeviceCode(SourceLocation Loc, - unsigned DiagID) { +Sema::SemaDiagnosticBuilder Sema::SYCLDiagIfDeviceCode(SourceLocation Loc, + unsigned DiagID) { assert(getLangOpts().SYCLIsDevice && "Should only be called during SYCL compilation"); FunctionDecl *FD = dyn_cast(getCurLexicalContext()); - DeviceDiagBuilder::Kind DiagKind = [this, FD] { + SemaDiagnosticBuilder::Kind DiagKind = [this, FD] { if (!FD) - return DeviceDiagBuilder::K_Nop; + return SemaDiagnosticBuilder::K_Nop; if (getEmissionStatus(FD) == Sema::FunctionEmissionStatus::Emitted) - return DeviceDiagBuilder::K_ImmediateWithCallStack; - return DeviceDiagBuilder::K_Deferred; + return SemaDiagnosticBuilder::K_ImmediateWithCallStack; + return SemaDiagnosticBuilder::K_Deferred; }(); - return DeviceDiagBuilder(DiagKind, Loc, DiagID, FD, *this); + return SemaDiagnosticBuilder(DiagKind, Loc, DiagID, FD, *this); } bool Sema::checkSYCLDeviceFunction(SourceLocation Loc, FunctionDecl *Callee) { @@ -42,8 +42,8 @@ if (isUnevaluatedContext() || isConstantEvaluated()) return true; - DeviceDiagBuilder::Kind DiagKind = DeviceDiagBuilder::K_Nop; + SemaDiagnosticBuilder::Kind DiagKind = SemaDiagnosticBuilder::K_Nop; - return DiagKind != DeviceDiagBuilder::K_Immediate && - DiagKind != DeviceDiagBuilder::K_ImmediateWithCallStack; + return DiagKind != SemaDiagnosticBuilder::K_Immediate && + DiagKind != SemaDiagnosticBuilder::K_ImmediateWithCallStack; } Index: clang/lib/Sema/SemaStmt.cpp =================================================================== --- clang/lib/Sema/SemaStmt.cpp +++ clang/lib/Sema/SemaStmt.cpp @@ -1249,10 +1249,10 @@ // Produce a nice diagnostic if multiple values aren't handled. if (!UnhandledNames.empty()) { - DiagnosticBuilder DB = Diag(CondExpr->getExprLoc(), - TheDefaultStmt ? diag::warn_def_missing_case + auto 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) Index: clang/lib/Sema/SemaStmtAsm.cpp =================================================================== --- clang/lib/Sema/SemaStmtAsm.cpp +++ clang/lib/Sema/SemaStmtAsm.cpp @@ -448,9 +448,9 @@ unsigned Size = Context.getTypeSize(Ty); if (!Context.getTargetInfo().validateInputSize(FeatureMap, Literal->getString(), Size)) - return StmtResult( - targetDiag(InputExpr->getBeginLoc(), diag::err_asm_invalid_input_size) - << Info.getConstraintStr()); + return targetDiag(InputExpr->getBeginLoc(), + diag::err_asm_invalid_input_size) + << Info.getConstraintStr(); } // Check that the clobbers are valid. Index: clang/lib/Sema/SemaTemplateInstantiate.cpp =================================================================== --- clang/lib/Sema/SemaTemplateInstantiate.cpp +++ clang/lib/Sema/SemaTemplateInstantiate.cpp @@ -237,7 +237,7 @@ // 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.Diags.hasUncompilableErrorOccurred()) { + SemaRef.hasUncompilableErrorOccurred()) { Invalid = true; return; } Index: clang/lib/Sema/SemaTemplateInstantiateDecl.cpp =================================================================== --- clang/lib/Sema/SemaTemplateInstantiateDecl.cpp +++ clang/lib/Sema/SemaTemplateInstantiateDecl.cpp @@ -5979,7 +5979,7 @@ if (!Result) { if (isa(D)) { // UsingShadowDecls can instantiate to nothing because of using hiding. - } else if (Diags.hasUncompilableErrorOccurred()) { + } else if (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. Index: clang/lib/Sema/SemaTemplateVariadic.cpp =================================================================== --- clang/lib/Sema/SemaTemplateVariadic.cpp +++ clang/lib/Sema/SemaTemplateVariadic.cpp @@ -368,8 +368,8 @@ Locations.push_back(Unexpanded[I].second); } - DiagnosticBuilder DB = Diag(Loc, diag::err_unexpanded_parameter_pack) - << (int)UPPC << (int)Names.size(); + auto 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]; Index: clang/lib/Sema/SemaType.cpp =================================================================== --- clang/lib/Sema/SemaType.cpp +++ clang/lib/Sema/SemaType.cpp @@ -4131,7 +4131,8 @@ /// Creates a fix-it to insert a C-style nullability keyword at \p pointerLoc, /// taking into account whitespace before and after. -static void fixItNullability(Sema &S, DiagnosticBuilder &Diag, +template +static void fixItNullability(Sema &S, DiagBuilderT &Diag, SourceLocation PointerLoc, NullabilityKind Nullability) { assert(PointerLoc.isValid()); Index: clang/test/SemaCUDA/deferred-oeverload.cu =================================================================== --- /dev/null +++ clang/test/SemaCUDA/deferred-oeverload.cu @@ -0,0 +1,78 @@ +// 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 +__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 +__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 '}'}} Index: clang/test/TableGen/DiagnosticBase.inc =================================================================== --- clang/test/TableGen/DiagnosticBase.inc +++ clang/test/TableGen/DiagnosticBase.inc @@ -45,6 +45,7 @@ // diagnostics string Component = ""; string CategoryName = ""; + bit Deferrable = 0; } // Diagnostic Categories. These can be applied to groups or individual @@ -75,6 +76,7 @@ bit AccessControl = 0; bit WarningNoWerror = 0; bit ShowInSystemHeader = 0; + bit Deferrable = 0; Severity DefaultSeverity = defaultmapping; DiagGroup Group; string CategoryName = ""; @@ -98,6 +100,14 @@ 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 : Diagnostic, SFINAEFailure { bit ShowInSystemHeader = 1; Index: clang/test/TableGen/deferred-diag.td =================================================================== --- /dev/null +++ clang/test/TableGen/deferred-diag.td @@ -0,0 +1,27 @@ +// 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 Index: clang/tools/diagtool/DiagnosticNames.cpp =================================================================== --- clang/tools/diagtool/DiagnosticNames.cpp +++ clang/tools/diagtool/DiagnosticNames.cpp @@ -28,7 +28,7 @@ // out of sync easily? static const DiagnosticRecord BuiltinDiagnosticsByID[] = { #define DIAG(ENUM,CLASS,DEFAULT_MAPPING,DESC,GROUP, \ - SFINAE,NOWERROR,SHOWINSYSHEADER,CATEGORY) \ + SFINAE,NOWERROR,SHOWINSYSHEADER,DEFER,CATEGORY) \ { #ENUM, diag::ENUM, STR_SIZE(#ENUM, uint8_t) }, #include "clang/Basic/DiagnosticCommonKinds.inc" #include "clang/Basic/DiagnosticCrossTUKinds.inc" Index: clang/utils/TableGen/ClangDiagnosticsEmitter.cpp =================================================================== --- clang/utils/TableGen/ClangDiagnosticsEmitter.cpp +++ clang/utils/TableGen/ClangDiagnosticsEmitter.cpp @@ -1294,6 +1294,11 @@ else OS << ", false"; + if (R.getValueAsBit("Deferrable")) + OS << ", true"; + else + OS << ", false"; + // Category number. OS << ", " << CategoryIDs.getID(getDiagnosticCategory(&R, DGParentMap)); OS << ")\n";