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,136 @@ } /// 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); /// Emit a partial diagnostic. - SemaDiagnosticBuilder Diag(SourceLocation Loc, const PartialDiagnostic& PD); + 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. @@ -11616,84 +11727,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 @@ -11706,15 +11744,18 @@ /// 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, + bool EmitOnBothSides = false); - /// 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, + bool EmitOnBothSides = false); - /// 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 @@ -11729,9 +11770,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 @@ -11744,9 +11786,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. @@ -12518,7 +12565,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 @@ -12536,7 +12583,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/Sema/AnalysisBasedWarnings.cpp =================================================================== --- clang/lib/Sema/AnalysisBasedWarnings.cpp +++ clang/lib/Sema/AnalysisBasedWarnings.cpp @@ -2089,7 +2089,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 @@ -1432,12 +1432,25 @@ PrintContextStack(); } -Sema::SemaDiagnosticBuilder -Sema::Diag(SourceLocation Loc, const PartialDiagnostic& PD) { - SemaDiagnosticBuilder Builder(Diag(Loc, PD.getDiagID())); - PD.Emit(Builder); +Sema::SemaDiagnosticBuilder Sema::Diag(SourceLocation Loc, + const PartialDiagnostic &PD) { + return Diag(Loc, PD.getDiagID()) << 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 @@ -1649,9 +1662,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) { @@ -1659,7 +1672,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."); @@ -1670,7 +1684,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) { @@ -1680,7 +1694,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. @@ -1695,7 +1709,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); @@ -1706,8 +1721,23 @@ 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) { + if (!getLangOpts().CUDA) + return SemaDiagnosticBuilder(SemaDiagnosticBuilder::K_Immediate, Loc, + DiagID, getCurFunctionDecl(), *this); + + SemaDiagnosticBuilder DB = + getLangOpts().CUDAIsDevice + ? CUDADiagIfDeviceCode(Loc, DiagID, /*EmitOnBothSides=*/true) + : CUDADiagIfHostCode(Loc, DiagID, /*EmitOnBothSides=*/true); + + if (Diags.getDiagnosticIDs()->isDefaultMappingAsError(DiagID)) + IsLastErrorImmediate = 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 @@ -382,8 +382,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,67 @@ } } -Sema::DeviceDiagBuilder Sema::CUDADiagIfDeviceCode(SourceLocation Loc, - unsigned DiagID) { +Sema::SemaDiagnosticBuilder Sema::CUDADiagIfDeviceCode(SourceLocation Loc, + unsigned DiagID, + bool EmitOnBothSides) { 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 EmitOnBothSides ? SemaDiagnosticBuilder::K_Immediate + : 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, + bool EmitOnBothSides) { 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 EmitOnBothSides ? SemaDiagnosticBuilder::K_Immediate + : 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 +720,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 +729,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 +747,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 +791,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 @@ -14451,11 +14451,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. @@ -14507,7 +14507,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 @@ -1859,27 +1859,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"); @@ -1887,30 +1887,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/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 @@ -1244,10 +1244,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 @@ -5999,7 +5999,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 @@ -4087,7 +4087,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/asm-constraints-mixed.cu =================================================================== --- clang/test/SemaCUDA/asm-constraints-mixed.cu +++ clang/test/SemaCUDA/asm-constraints-mixed.cu @@ -1,7 +1,7 @@ // REQUIRES: x86-registered-target // REQUIRES: nvptx-registered-target -// RUN: %clang_cc1 -triple nvptx-unknown-cuda -fsyntax-only -fcuda-is-device -verify %s -// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fsyntax-only -verify %s +// RUN: %clang_cc1 -triple nvptx-unknown-cuda -fsyntax-only -fcuda-is-device -verify=dev %s +// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fsyntax-only -verify=host %s __attribute__((device)) register long global_dev_reg asm("r0"); __attribute__((device)) register long @@ -30,10 +30,10 @@ // We should only see errors relevant to current compilation mode. #if defined(__CUDA_ARCH__) // Device-side compilation: -// expected-error@8 {{unknown register name 'rsp' in asm}} -// expected-error@15 {{unknown register name 'rsp' in asm}} +// dev-error@8 {{unknown register name 'rsp' in asm}} +// dev-error@15 {{unknown register name 'rsp' in asm}} #else // Host-side compilation: -// expected-error@11 {{unknown register name 'r0' in asm}} -// expected-error@22 {{unknown register name 'r0' in asm}} +// host-error@11 {{unknown register name 'r0' in asm}} +// host-error@22 {{unknown register name 'r0' in asm}} #endif Index: clang/test/SemaCUDA/constexpr-variables.cu =================================================================== --- clang/test/SemaCUDA/constexpr-variables.cu +++ clang/test/SemaCUDA/constexpr-variables.cu @@ -1,23 +1,23 @@ // RUN: %clang_cc1 -std=c++14 %s -emit-llvm -o - -triple nvptx64-nvidia-cuda \ -// RUN: -fcuda-is-device -verify -fsyntax-only +// RUN: -fcuda-is-device -verify=expected,dev -fsyntax-only // RUN: %clang_cc1 -std=c++17 %s -emit-llvm -o - -triple nvptx64-nvidia-cuda \ -// RUN: -fcuda-is-device -verify -fsyntax-only +// RUN: -fcuda-is-device -verify=expected,dev -fsyntax-only // RUN: %clang_cc1 -std=c++14 %s -emit-llvm -o - \ -// RUN: -triple x86_64-unknown-linux-gnu -verify -fsyntax-only +// RUN: -triple x86_64-unknown-linux-gnu -verify=expected,host -fsyntax-only // RUN: %clang_cc1 -std=c++17 %s -emit-llvm -o - \ -// RUN: -triple x86_64-unknown-linux-gnu -verify -fsyntax-only +// RUN: -triple x86_64-unknown-linux-gnu -verify=expected,host -fsyntax-only #include "Inputs/cuda.h" template __host__ __device__ void foo(const T **a) { - // expected-note@-1 {{declared here}} + // expected-note@-1 2{{declared here}} static const T b = sizeof(a); static constexpr T c = sizeof(a); const T d = sizeof(a); constexpr T e = sizeof(a); constexpr T f = **a; - // expected-error@-1 {{constexpr variable 'f' must be initialized by a constant expression}} - // expected-note@-2 {{read of non-constexpr variable 'a' is not allowed in a constant expression}} + // expected-error@-1 2{{constexpr variable 'f' must be initialized by a constant expression}} + // expected-note@-2 2{{read of non-constexpr variable 'a' is not allowed in a constant expression}} a[0] = &b; a[1] = &c; a[2] = &d; @@ -34,7 +34,7 @@ a[0] = &b; a[1] = &c; foo(a); - // expected-note@-1 {{in instantiation of function template specialization 'foo' requested here}} + // dev-note@-1 {{called by 'device_fun'}} } void host_fun(const int **a) { @@ -47,6 +47,7 @@ a[0] = &b; a[1] = &c; foo(a); + // host-note@-1 {{called by 'host_fun'}} } __host__ __device__ void host_device_fun(const int **a) { @@ -59,6 +60,7 @@ a[0] = &b; a[1] = &c; foo(a); + // expected-note@-1 {{called by 'host_device_fun'}} } template Index: clang/test/SemaCUDA/cxx11-kernel-call.cu =================================================================== --- clang/test/SemaCUDA/cxx11-kernel-call.cu +++ clang/test/SemaCUDA/cxx11-kernel-call.cu @@ -7,4 +7,10 @@ template void k1Wrapper() { void (*f)() = [] { k1<<>>(); }; // expected-error {{initializer contains unexpanded parameter pack 'Dimensions'}} void (*g[])() = { [] { k1<<>>(); } ... }; // ok + f(); + g(); +} + +void foo() { + k1Wrapper<1,1>(); } Index: clang/test/SemaCUDA/exceptions.cu =================================================================== --- clang/test/SemaCUDA/exceptions.cu +++ clang/test/SemaCUDA/exceptions.cu @@ -1,5 +1,5 @@ -// RUN: %clang_cc1 -fcxx-exceptions -fcuda-is-device -fsyntax-only -verify %s -// RUN: %clang_cc1 -fcxx-exceptions -fsyntax-only -verify %s +// RUN: %clang_cc1 -fcxx-exceptions -fcuda-is-device -fsyntax-only -verify=expected,dev %s +// RUN: %clang_cc1 -fcxx-exceptions -fsyntax-only -verify=expected,host %s #include "Inputs/cuda.h" Index: clang/test/SemaCUDA/implicit-member-target-inherited.cu =================================================================== --- clang/test/SemaCUDA/implicit-member-target-inherited.cu +++ clang/test/SemaCUDA/implicit-member-target-inherited.cu @@ -1,4 +1,5 @@ -// RUN: %clang_cc1 -std=c++11 -fsyntax-only -verify %s -Wno-defaulted-function-deleted +// RUN: %clang_cc1 -std=c++11 -fsyntax-only -verify=expected,host %s -Wno-defaulted-function-deleted +// RUN: %clang_cc1 -std=c++11 -fsyntax-only -fcuda-is-device -verify=expected,dev %s -Wno-defaulted-function-deleted #include "Inputs/cuda.h" Index: clang/test/SemaCUDA/lambda.cu =================================================================== --- clang/test/SemaCUDA/lambda.cu +++ clang/test/SemaCUDA/lambda.cu @@ -1,4 +1,4 @@ -// RUN: %clang_cc1 -std=c++17 -fsyntax-only -verify=com %s +// RUN: %clang_cc1 -std=c++17 -fsyntax-only -verify=com,host %s // RUN: %clang_cc1 -std=c++17 -fsyntax-only -fcuda-is-device -verify=com,dev %s #include "Inputs/cuda.h" Index: clang/test/SemaCUDA/no-call-stack-for-immediate-errs.cu =================================================================== --- clang/test/SemaCUDA/no-call-stack-for-immediate-errs.cu +++ clang/test/SemaCUDA/no-call-stack-for-immediate-errs.cu @@ -2,11 +2,6 @@ #include "Inputs/cuda.h" -// Here we should dump an error about the VLA in device_fn, but we should not -// print a callstack indicating how device_fn becomes known-emitted, because -// it's an error to use a VLA in any __device__ function, even one that doesn't -// get emitted. - inline __device__ void device_fn(int n); inline __device__ void device_fn2() { device_fn(42); } Index: clang/test/SemaCUDA/union-init.cu =================================================================== --- clang/test/SemaCUDA/union-init.cu +++ clang/test/SemaCUDA/union-init.cu @@ -1,4 +1,7 @@ -// RUN: %clang_cc1 %s --std=c++11 -triple x86_64-linux-unknown -fsyntax-only -o - -verify +// RUN: %clang_cc1 %s --std=c++11 -triple x86_64-linux-unknown \ +// RUN: -fsyntax-only -o - -verify=com,host +// RUN: %clang_cc1 %s --std=c++11 -triple nvptx -fcuda-is-device \ +// RUN: -fsyntax-only -o - -verify=com,dev #include "Inputs/cuda.h" @@ -31,14 +34,14 @@ __device__ B b; __device__ C c; -// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}} +// com-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}} __device__ D d; -// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}} +// com-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}} __device__ void foo() { __shared__ B b; __shared__ C c; - // expected-error@-1 {{initialization is not supported for __shared__ variables.}} + // com-error@-1 {{initialization is not supported for __shared__ variables.}} __shared__ D d; - // expected-error@-1 {{initialization is not supported for __shared__ variables.}} + // com-error@-1 {{initialization is not supported for __shared__ variables.}} } Index: clang/test/SemaCUDA/usual-deallocators.cu =================================================================== --- clang/test/SemaCUDA/usual-deallocators.cu +++ clang/test/SemaCUDA/usual-deallocators.cu @@ -81,9 +81,9 @@ __host__ __device__ void tests_hd(void *t) { test_hd(t); test_hd(t); - // host-note@-1 {{in instantiation of function template specialization 'test_hd' requested here}} + // host-note@-1 {{called by 'tests_hd'}} test_hd(t); - // device-note@-1 {{in instantiation of function template specialization 'test_hd' requested here}} + // device-note@-1 {{called by 'tests_hd'}} test_hd(t); test_hd(t); test_hd(t);