yaxunl created this revision.
yaxunl added a reviewer: tra.
Herald added subscribers: sstefan1, kristof.beyls.
Herald added a reviewer: jdoerfert.
In CUDA/HIP a function may become implicit host device function by
pragma or constexpr. A host device function is checked in both
host and device compilation. However it may be emitted only
on host or device side, therefore the diagnostics should be
deferred until it is known to be emitted.
Currently clang is only able to defer certain diagnostics. This causes
false alarms and limits the usefulness of host device functions.
This patch lets clang defer all diagnostics for host device functions.
It is NFC for other languages.
https://reviews.llvm.org/D84364
Files:
clang/include/clang/Sema/Sema.h
clang/lib/Sema/AnalysisBasedWarnings.cpp
clang/lib/Sema/Sema.cpp
clang/lib/Sema/SemaAttr.cpp
clang/lib/Sema/SemaCUDA.cpp
clang/lib/Sema/SemaDecl.cpp
clang/lib/Sema/SemaExprObjC.cpp
clang/lib/Sema/SemaOpenMP.cpp
clang/lib/Sema/SemaSYCL.cpp
clang/lib/Sema/SemaStmt.cpp
clang/lib/Sema/SemaStmtAsm.cpp
clang/lib/Sema/SemaTemplateInstantiate.cpp
clang/lib/Sema/SemaTemplateInstantiateDecl.cpp
clang/lib/Sema/SemaTemplateVariadic.cpp
clang/lib/Sema/SemaType.cpp
clang/test/SemaCUDA/asm-constraints-mixed.cu
clang/test/SemaCUDA/constexpr-variables.cu
clang/test/SemaCUDA/cxx11-kernel-call.cu
clang/test/SemaCUDA/exceptions.cu
clang/test/SemaCUDA/implicit-member-target-inherited.cu
clang/test/SemaCUDA/lambda.cu
clang/test/SemaCUDA/no-call-stack-for-immediate-errs.cu
clang/test/SemaCUDA/union-init.cu
clang/test/SemaCUDA/usual-deallocators.cu
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<H1D1>(t);
test_hd<h1D1>(t);
- // host-note@-1 {{in instantiation of function template specialization 'test_hd<h1D1>' requested here}}
+ // host-note@-1 {{called by 'tests_hd'}}
test_hd<H1d1>(t);
- // device-note@-1 {{in instantiation of function template specialization 'test_hd<H1d1>' requested here}}
+ // device-note@-1 {{called by 'tests_hd'}}
test_hd<H1D2>(t);
test_hd<H2D1>(t);
test_hd<H2D2>(t);
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/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/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/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/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/cxx11-kernel-call.cu
===================================================================
--- clang/test/SemaCUDA/cxx11-kernel-call.cu
+++ clang/test/SemaCUDA/cxx11-kernel-call.cu
@@ -7,4 +7,10 @@
template<int ...Dimensions> void k1Wrapper() {
void (*f)() = [] { k1<<<Dimensions, Dimensions>>>(); }; // expected-error {{initializer contains unexpanded parameter pack 'Dimensions'}}
void (*g[])() = { [] { k1<<<Dimensions, Dimensions>>>(); } ... }; // ok
+ f();
+ g();
+}
+
+void foo() {
+ k1Wrapper<1,1>();
}
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<typename T>
__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<int>' 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 <class T>
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/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 <typename DiagBuilderT>
+static void fixItNullability(Sema &S, DiagBuilderT &Diag,
SourceLocation PointerLoc,
NullabilityKind Nullability) {
assert(PointerLoc.isValid());
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/SemaTemplateInstantiateDecl.cpp
===================================================================
--- clang/lib/Sema/SemaTemplateInstantiateDecl.cpp
+++ clang/lib/Sema/SemaTemplateInstantiateDecl.cpp
@@ -5999,7 +5999,7 @@
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.
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/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/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/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<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) {
@@ -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/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/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 <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:
@@ -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/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<FunctionTemplateDecl>(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/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<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 EmitOnBothSides ? SemaDiagnosticBuilder::K_Immediate
+ : 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,
+ bool EmitOnBothSides) {
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 EmitOnBothSides ? SemaDiagnosticBuilder::K_Immediate
+ : 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) {
@@ -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/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/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<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
@@ -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/AnalysisBasedWarnings.cpp
===================================================================
--- clang/lib/Sema/AnalysisBasedWarnings.cpp
+++ clang/lib/Sema/AnalysisBasedWarnings.cpp
@@ -2089,7 +2089,7 @@
if (cast<DeclContext>(D)->isDependentContext())
return;
- if (Diags.hasUncompilableErrorOccurred()) {
+ if (S.hasUncompilableErrorOccurred()) {
// Flush out any possibly unreachable diagnostics.
flushDiagnostics(S, fscope);
return;
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,27 +1507,136 @@
}
/// 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;
}
};
+ /// 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;
+ }
+
+ 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);
/// 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.
@@ -11615,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 <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
@@ -11705,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
@@ -11728,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
@@ -11743,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.
@@ -12517,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
@@ -12535,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.
///
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
