const char *getDiagnosticCode(unsigned ID) {
switch (ID) {
#define DIAG(ENUM, CLASS, DEFAULT_MAPPING, DESC, GROPU, SFINAE, NOWERROR, \
- SHOWINSYSHEADER, CATEGORY) \
+ SHOWINSYSHEADER, DEFERRABLE, CATEGORY) \
case clang::diag::ENUM: \
return #ENUM;
#include "clang/Basic/DiagnosticASTKinds.inc"
// diagnostics
string Component = "";
string CategoryName = "";
+ bit Deferrable = 0;
}
// Diagnostic Categories. These can be applied to groups or individual
bit AccessControl = 0;
bit WarningNoWerror = 0;
bit ShowInSystemHeader = 0;
+ bit Deferrable = 0;
Severity DefaultSeverity = defaultmapping;
DiagGroup Group;
string CategoryName = "";
bit ShowInSystemHeader = 0;
}
+class Deferrable {
+ bit Deferrable = 1;
+}
+
+class NonDeferrable {
+ bit Deferrable = 0;
+}
+
// FIXME: ExtWarn and Extension should also be SFINAEFailure by default.
class Error<string str> : Diagnostic<str, CLASS_ERROR, SEV_Error>, SFINAEFailure {
bit ShowInSystemHeader = 1;
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"
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"
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"
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"
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"
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"
// 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
/// 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
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"
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"
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"
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"
"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<
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"
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"
LANGOPT(GPURelocatableDeviceCode, 1, 0, "generate relocatable device code")
LANGOPT(GPUAllowDeviceInit, 1, 0, "allowing device side global init functions for HIP")
LANGOPT(GPUMaxThreadsPerBlock, 32, 256, "default max threads per block for kernel launch bounds for HIP")
+LANGOPT(GPUDeferDiag, 1, 0, "defer host/device related diagnostic messages for CUDA/HIP")
LANGOPT(SYCL , 1, 0, "SYCL")
LANGOPT(SYCLIsDevice , 1, 0, "Generate code for SYCL device")
"Use", "Don't use", " new kernel launching API for HIP">;
defm gpu_allow_device_init : OptInFFlag<"gpu-allow-device-init",
"Allow", "Don't allow", " device side init function in HIP">;
+defm gpu_defer_diag : OptInFFlag<"gpu-defer-diag",
+ "Defer", "Don't defer", " host/device related diagnostic messages"
+ " for CUDA/HIP">;
def gpu_max_threads_per_block_EQ : Joined<["--"], "gpu-max-threads-per-block=">,
Flags<[CC1Option]>,
HelpText<"Default max threads per block for kernel launch bounds for HIP">;
/// 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) { }
+ 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;
}
/// Teach operator<< to produce an object of the correct type.
- template<typename T>
- friend const SemaDiagnosticBuilder &operator<<(
- const SemaDiagnosticBuilder &Diag, const T &Value) {
+ template <typename T>
+ friend const ImmediateDiagBuilder &
+ operator<<(const ImmediateDiagBuilder &Diag, const T &Value) {
const DiagnosticBuilder &BaseDiag = Diag;
BaseDiag << Value;
return Diag;
// bitfield is not allowed.
template <typename T, typename = typename std::enable_if<
!std::is_lvalue_reference<T>::value>::type>
+ const ImmediateDiagBuilder &operator<<(T &&V) const {
+ const DiagnosticBuilder &BaseDiag = *this;
+ BaseDiag << std::move(V);
+ return *this;
+ }
+ };
+
+ /// A generic diagnostic builder for errors which may or may not be deferred.
+ ///
+ /// In CUDA, there exist constructs (e.g. variable-length arrays, try/catch)
+ /// which are not allowed to appear inside __device__ functions and are
+ /// allowed to appear in __host__ __device__ functions only if the host+device
+ /// function is never codegen'ed.
+ ///
+ /// To handle this, we use the notion of "deferred diagnostics", where we
+ /// attach a diagnostic to a FunctionDecl that's emitted iff it's codegen'ed.
+ ///
+ /// This class lets you emit either a regular diagnostic, a deferred
+ /// diagnostic, or no diagnostic at all, according to an argument you pass to
+ /// its constructor, thus simplifying the process of creating these "maybe
+ /// deferred" diagnostics.
+ class SemaDiagnosticBuilder {
+ public:
+ enum Kind {
+ /// Emit no diagnostics.
+ K_Nop,
+ /// Emit the diagnostic immediately (i.e., behave like Sema::Diag()).
+ K_Immediate,
+ /// Emit the diagnostic immediately, and, if it's a warning or error, also
+ /// emit a call stack showing how this function can be reached by an a
+ /// priori known-emitted function.
+ K_ImmediateWithCallStack,
+ /// Create a deferred diagnostic, which is emitted only if the function
+ /// it's attached to is codegen'ed. Also emit a call stack as with
+ /// K_ImmediateWithCallStack.
+ K_Deferred
+ };
+
+ SemaDiagnosticBuilder(Kind K, SourceLocation Loc, unsigned DiagID,
+ FunctionDecl *Fn, Sema &S);
+ SemaDiagnosticBuilder(SemaDiagnosticBuilder &&D);
+ SemaDiagnosticBuilder(const SemaDiagnosticBuilder &) = default;
+ ~SemaDiagnosticBuilder();
+
+ bool isImmediate() const { return ImmediateDiag.hasValue(); }
+
+ /// Convertible to bool: True if we immediately emitted an error, false if
+ /// we didn't emit an error or we created a deferred error.
+ ///
+ /// Example usage:
+ ///
+ /// if (SemaDiagnosticBuilder(...) << foo << bar)
+ /// return ExprError();
+ ///
+ /// But see CUDADiagIfDeviceCode() and CUDADiagIfHostCode() -- you probably
+ /// want to use these instead of creating a SemaDiagnosticBuilder yourself.
+ operator bool() const { return isImmediate(); }
+
+ template <typename T>
+ friend const SemaDiagnosticBuilder &
+ operator<<(const SemaDiagnosticBuilder &Diag, const T &Value) {
+ if (Diag.ImmediateDiag.hasValue())
+ *Diag.ImmediateDiag << Value;
+ else if (Diag.PartialDiagId.hasValue())
+ Diag.S.DeviceDeferredDiags[Diag.Fn][*Diag.PartialDiagId].second
+ << Value;
+ return Diag;
+ }
+
+ // It is necessary to limit this to rvalue reference to avoid calling this
+ // function with a bitfield lvalue argument since non-const reference to
+ // bitfield is not allowed.
+ template <typename T, typename = typename std::enable_if<
+ !std::is_lvalue_reference<T>::value>::type>
const SemaDiagnosticBuilder &operator<<(T &&V) const {
- const StreamableDiagnosticBase &DB = *this;
- DB << std::move(V);
+ if (ImmediateDiag.hasValue())
+ *ImmediateDiag << std::move(V);
+ else if (PartialDiagId.hasValue())
+ S.DeviceDeferredDiags[Fn][*PartialDiagId].second << std::move(V);
return *this;
}
+
+ friend const SemaDiagnosticBuilder &
+ operator<<(const SemaDiagnosticBuilder &Diag, const PartialDiagnostic &PD) {
+ if (Diag.ImmediateDiag.hasValue())
+ PD.Emit(*Diag.ImmediateDiag);
+ else if (Diag.PartialDiagId.hasValue())
+ Diag.S.DeviceDeferredDiags[Diag.Fn][*Diag.PartialDiagId].second = PD;
+ return Diag;
+ }
+
+ void AddFixItHint(const FixItHint &Hint) const {
+ if (ImmediateDiag.hasValue())
+ ImmediateDiag->AddFixItHint(Hint);
+ else if (PartialDiagId.hasValue())
+ S.DeviceDeferredDiags[Fn][*PartialDiagId].second.AddFixItHint(Hint);
+ }
+
+ friend ExprResult ExprError(const SemaDiagnosticBuilder &) {
+ return ExprError();
+ }
+ friend StmtResult StmtError(const SemaDiagnosticBuilder &) {
+ return StmtError();
+ }
+ operator ExprResult() const { return ExprError(); }
+ operator StmtResult() const { return StmtError(); }
+ operator TypeResult() const { return TypeError(); }
+ operator DeclResult() const { return DeclResult(true); }
+ operator MemInitResult() const { return MemInitResult(true); }
+
+ private:
+ Sema &S;
+ SourceLocation Loc;
+ unsigned DiagID;
+ FunctionDecl *Fn;
+ bool ShowCallStack;
+
+ // Invariant: At most one of these Optionals has a value.
+ // FIXME: Switch these to a Variant once that exists.
+ llvm::Optional<ImmediateDiagBuilder> ImmediateDiag;
+ llvm::Optional<unsigned> PartialDiagId;
};
+ using DiagBuilderT = SemaDiagnosticBuilder;
+
+ /// Is the last error level diagnostic immediate. This is used to determined
+ /// whether the next info diagnostic should be immediate.
+ bool IsLastErrorImmediate = true;
/// Emit a diagnostic.
- SemaDiagnosticBuilder Diag(SourceLocation Loc, unsigned DiagID) {
- 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.
/* 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 <typename T>
- friend const DeviceDiagBuilder &operator<<(const DeviceDiagBuilder &Diag,
- const T &Value) {
- if (Diag.ImmediateDiag.hasValue())
- *Diag.ImmediateDiag << Value;
- else if (Diag.PartialDiagId.hasValue())
- Diag.S.DeviceDeferredDiags[Diag.Fn][*Diag.PartialDiagId].second
- << Value;
- return Diag;
- }
-
- private:
- Sema &S;
- SourceLocation Loc;
- unsigned DiagID;
- FunctionDecl *Fn;
- bool ShowCallStack;
-
- // Invariant: At most one of these Optionals has a value.
- // FIXME: Switch these to a Variant once that exists.
- llvm::Optional<SemaDiagnosticBuilder> ImmediateDiag;
- llvm::Optional<unsigned> PartialDiagId;
- };
-
- /// Creates a DeviceDiagBuilder that emits the diagnostic if the current context
- /// is "used as device code".
+ /// 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
/// 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
/// 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
/// 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.
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
/// 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.
///
unsigned SFINAE : 2;
unsigned WarnNoWerror : 1;
unsigned WarnShowInSystemHeader : 1;
+ unsigned Deferrable : 1;
unsigned Category : 6;
uint16_t OptionGroupIndex;
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"
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) {
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))
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
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;
if (cast<DeclContext>(D)->isDependentContext())
return;
- if (Diags.hasUncompilableErrorOccurred()) {
+ if (S.hasUncompilableErrorOccurred()) {
// Flush out any possibly unreachable diagnostics.
flushDiagnostics(S, fscope);
return;
}
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<FunctionDecl>(CurContext);
+ if (!FD)
+ return false;
+ auto Loc = DeviceDeferredDiags.find(FD);
+ if (Loc == DeviceDeferredDiags.end())
+ return false;
+ for (auto PDAt : Loc->second) {
+ if (DiagnosticIDs::isDefaultMappingAsError(PDAt.second.getDiagID()))
+ return true;
+ }
+ return false;
}
// Print notes showing how we can reach FD starting from an a priori
// 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) {
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.");
}
}
-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) {
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.
}
}
-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);
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) {
// 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);
}
}
-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<FunctionDecl>(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<FunctionDecl>(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<FunctionDecl>(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<FunctionDecl>(CurContext), *this);
+ return SemaDiagnosticBuilder(DiagKind, Loc, DiagID,
+ dyn_cast<FunctionDecl>(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<FunctionDecl>(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<FunctionDecl>(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<FunctionDecl>(CurContext), *this);
+ return SemaDiagnosticBuilder(DiagKind, Loc, DiagID,
+ dyn_cast<FunctionDecl>(CurContext), *this);
}
bool Sema::CheckCUDACall(SourceLocation Loc, FunctionDecl *Callee) {
// Otherwise, mark the call in our call graph so we can traverse it later.
bool CallerKnownEmitted =
getEmissionStatus(Caller) == FunctionEmissionStatus::Emitted;
- DeviceDiagBuilder::Kind DiagKind = [this, Caller, Callee,
- CallerKnownEmitted] {
+ SemaDiagnosticBuilder::Kind DiagKind = [this, Caller, Callee,
+ CallerKnownEmitted] {
switch (IdentifyCUDAPreference(Caller, Callee)) {
case CFP_Never:
case CFP_WrongSide:
// 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
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.
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;
}
// 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<FunctionTemplateDecl>(dcl)) {
// Since the body is valid, issue any analysis-based warnings that are
// enabled.
// 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();
}
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;
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()));
if (SuperLoc.isValid()) {
{
- DiagnosticBuilder Builder =
+ auto Builder =
Diag(SuperLoc, diag::err_messaging_super_with_direct_method);
if (ReceiverType->isObjCClassType()) {
Builder.AddFixItHint(FixItHint::CreateReplacement(
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 <typename DiagBuilderT>
+static void addFixitForObjCARCConversion(
+ Sema &S, DiagBuilderT &DiagB, Sema::CheckedConversionKind CCK,
+ SourceLocation afterLParen, QualType castType, Expr *castExpr,
+ Expr *realCast, const char *bridgeKeyword, const char *CFBridgeName) {
// We handle C-style and implicit casts here.
switch (CCK) {
case Sema::CCK_ImplicitConversion:
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 ",
}
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 ",
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 ",
};
} // 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");
}
}
- 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
StringRef Opc, SourceLocation OpLoc,
llvm::function_ref<bool(OverloadCandidate &)> Filter) {
+ bool DeferHint = false;
+ if (S.getLangOpts().CUDA && S.getLangOpts().GPUDeferDiag) {
+ // Defer diagnostic for CUDA/HIP if there are wrong-sided candidates.
+ auto WrongSidedCands =
+ CompleteCandidates(S, OCD_AllCandidates, Args, OpLoc, [](auto &Cand) {
+ return Cand.Viable == false &&
+ Cand.FailureKind == ovl_fail_bad_target;
+ });
+ DeferHint = WrongSidedCands.size();
+ }
auto Cands = CompleteCandidates(S, OCD, Args, OpLoc, Filter);
- S.Diag(PD.first, PD.second);
+ S.Diag(PD.first, PD.second, DeferHint);
NoteCandidates(S, Args, Cands, Opc, OpLoc);
// 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<FunctionDecl>(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) {
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;
}
// 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)
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.
// 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;
}
if (!Result) {
if (isa<UsingShadowDecl>(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.
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];
/// 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 <typename DiagBuilderT>
+static void fixItNullability(Sema &S, DiagBuilderT &Diag,
SourceLocation PointerLoc,
NullabilityKind Nullability) {
assert(PointerLoc.isValid());
--- /dev/null
+// RUN: %clang_cc1 -fcuda-is-device -fsyntax-only -verify=dev,com %s \
+// RUN: -std=c++11 -fgpu-defer-diag
+// RUN: %clang_cc1 -fsyntax-only -verify=host,com %s \
+// RUN: -std=c++11 -fgpu-defer-diag
+
+#include "Inputs/cuda.h"
+
+// When callee is called by a host function with integer arguments, there is an error for ambiguity.
+// It should be deferred since it involves wrong-sided candidates.
+__device__ void callee(int);
+__host__ void callee(float); // host-note {{candidate function}}
+__host__ void callee(double); // host-note {{candidate function}}
+
+// When callee2 is called by a device function without arguments, there is an error for 'no matching function'.
+// It should be deferred since it involves wrong-sided candidates.
+__host__ void callee2(); // dev-note{{candidate function not viable: call to __host__ function from __device__ function}}
+
+// When callee3 is called by a device function without arguments, there is an error for 'no matching function'.
+// It should be deferred since it involves wrong-sided candidates.
+__host__ void callee3(); // dev-note{{candidate function not viable: call to __host__ function from __device__ function}}
+__device__ void callee3(int); // dev-note{{candidate function not viable: requires 1 argument, but 0 were provided}}
+
+// When callee4 is called by a host or device function without arguments, there is an error for 'no matching function'.
+// It should be immediate since it involves no wrong-sided candidates (it is not a viable candiate due to signature).
+__host__ void callee4(int); // com-note 2{{candidate function not viable: requires 1 argument, but 0 were provided}}
+
+// When callee5 is called by a host function with integer arguments, there is an error for ambiguity.
+// It should be immediate since it involves no wrong-sided candidates.
+__host__ void callee5(float); // com-note {{candidate function}}
+__host__ void callee5(double); // com-note {{candidate function}}
+
+__host__ void hf() {
+ callee(1); // host-error {{call to 'callee' is ambiguous}}
+ callee2();
+ callee3();
+ callee4(); // com-error {{no matching function for call to 'callee4'}}
+ callee5(1); // com-error {{call to 'callee5' is ambiguous}}
+ undeclared_func(); // com-error {{use of undeclared identifier 'undeclared_func'}}
+}
+
+__device__ void df() {
+ callee(1);
+ callee2(); // dev-error {{no matching function for call to 'callee2'}}
+ callee3(); // dev-error {{no matching function for call to 'callee3'}}
+ callee4(); // com-error {{no matching function for call to 'callee4'}}
+}
+
+struct A { int x; typedef int isA; };
+struct B { int x; };
+
+// This function is invalid for A and B by SFINAE.
+// This fails to substitue for A but no diagnostic
+// should be emitted.
+template<typename T, typename T::foo* = nullptr>
+__host__ __device__ void sfinae(T t) { // com-note {{candidate template ignored: substitution failure [with T = B]}}
+ t.x = 1;
+}
+
+// This function is defined for A only by SFINAE.
+// Calling it with A should succeed, with B should fail.
+// The error should not be deferred since it happens in
+// file scope.
+
+template<typename T, typename T::isA* = nullptr>
+__host__ __device__ void sfinae(T t) { // com-note {{candidate template ignored: substitution failure [with T = B]}}
+ t.x = 1;
+}
+
+void test_sfinae() {
+ sfinae(A());
+ sfinae(B()); // com-error{{no matching function for call to 'sfinae'}}
+}
+
+// If a syntax error causes a function not declared, it cannot
+// be deferred.
+
+inline __host__ __device__ void bad_func() { // com-note {{to match this '{'}}
+// com-error {{expected '}'}}
// diagnostics
string Component = "";
string CategoryName = "";
+ bit Deferrable = 0;
}
// Diagnostic Categories. These can be applied to groups or individual
bit AccessControl = 0;
bit WarningNoWerror = 0;
bit ShowInSystemHeader = 0;
+ bit Deferrable = 0;
Severity DefaultSeverity = defaultmapping;
DiagGroup Group;
string CategoryName = "";
bit ShowInSystemHeader = 0;
}
+class Deferrable {
+ bit Deferrable = 1;
+}
+
+class NonDeferrable {
+ bit Deferrable = 0;
+}
+
// FIXME: ExtWarn and Extension should also be SFINAEFailure by default.
class Error<string str> : Diagnostic<str, CLASS_ERROR, SEV_Error>, SFINAEFailure {
bit ShowInSystemHeader = 1;
--- /dev/null
+// 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
// 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"
else
OS << ", false";
+ if (R.getValueAsBit("Deferrable"))
+ OS << ", true";
+ else
+ OS << ", false";
+
// Category number.
OS << ", " << CategoryIDs.getID(getDiagnosticCategory(&R, DGParentMap));
OS << ")\n";