yaxunl updated this revision to Diff 279950.
yaxunl added a comment.
update for depending patch change
CHANGES SINCE LAST ACTION
https://reviews.llvm.org/D84364/new/
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,28 +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;
}
- const static bool IsDiagBuilder = false;
};
+ /// A generic diagnostic builder for errors which may or may not be deferred.
+ ///
+ /// In CUDA, there exist constructs (e.g. variable-length arrays, try/catch)
+ /// which are not allowed to appear inside __device__ functions and are
+ /// allowed to appear in __host__ __device__ functions only if the host+device
+ /// function is never codegen'ed.
+ ///
+ /// To handle this, we use the notion of "deferred diagnostics", where we
+ /// attach a diagnostic to a FunctionDecl that's emitted iff it's codegen'ed.
+ ///
+ /// This class lets you emit either a regular diagnostic, a deferred
+ /// diagnostic, or no diagnostic at all, according to an argument you pass to
+ /// its constructor, thus simplifying the process of creating these "maybe
+ /// deferred" diagnostics.
+ class SemaDiagnosticBuilder {
+ public:
+ enum Kind {
+ /// Emit no diagnostics.
+ K_Nop,
+ /// Emit the diagnostic immediately (i.e., behave like Sema::Diag()).
+ K_Immediate,
+ /// Emit the diagnostic immediately, and, if it's a warning or error, also
+ /// emit a call stack showing how this function can be reached by an a
+ /// priori known-emitted function.
+ K_ImmediateWithCallStack,
+ /// Create a deferred diagnostic, which is emitted only if the function
+ /// it's attached to is codegen'ed. Also emit a call stack as with
+ /// K_ImmediateWithCallStack.
+ K_Deferred
+ };
+
+ SemaDiagnosticBuilder(Kind K, SourceLocation Loc, unsigned DiagID,
+ FunctionDecl *Fn, Sema &S);
+ SemaDiagnosticBuilder(SemaDiagnosticBuilder &&D);
+ SemaDiagnosticBuilder(const SemaDiagnosticBuilder &) = default;
+ ~SemaDiagnosticBuilder();
+
+ bool isImmediate() const { return ImmediateDiag.hasValue(); }
+
+ /// Convertible to bool: True if we immediately emitted an error, false if
+ /// we didn't emit an error or we created a deferred error.
+ ///
+ /// Example usage:
+ ///
+ /// if (SemaDiagnosticBuilder(...) << foo << bar)
+ /// return ExprError();
+ ///
+ /// But see CUDADiagIfDeviceCode() and CUDADiagIfHostCode() -- you probably
+ /// want to use these instead of creating a SemaDiagnosticBuilder yourself.
+ operator bool() const { return isImmediate(); }
+
+ template <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.
@@ -11616,84 +11727,11 @@
/* Caller = */ FunctionDeclAndLoc>
DeviceKnownEmittedFns;
- /// Diagnostic builder for CUDA/OpenMP devices errors which may or may not be
- /// deferred.
- ///
- /// In CUDA, there exist constructs (e.g. variable-length arrays, try/catch)
- /// which are not allowed to appear inside __device__ functions and are
- /// allowed to appear in __host__ __device__ functions only if the host+device
- /// function is never codegen'ed.
- ///
- /// To handle this, we use the notion of "deferred diagnostics", where we
- /// attach a diagnostic to a FunctionDecl that's emitted iff it's codegen'ed.
- ///
- /// This class lets you emit either a regular diagnostic, a deferred
- /// diagnostic, or no diagnostic at all, according to an argument you pass to
- /// its constructor, thus simplifying the process of creating these "maybe
- /// deferred" diagnostics.
- class DeviceDiagBuilder {
- public:
- enum Kind {
- /// Emit no diagnostics.
- K_Nop,
- /// Emit the diagnostic immediately (i.e., behave like Sema::Diag()).
- K_Immediate,
- /// Emit the diagnostic immediately, and, if it's a warning or error, also
- /// emit a call stack showing how this function can be reached by an a
- /// priori known-emitted function.
- K_ImmediateWithCallStack,
- /// Create a deferred diagnostic, which is emitted only if the function
- /// it's attached to is codegen'ed. Also emit a call stack as with
- /// K_ImmediateWithCallStack.
- K_Deferred
- };
-
- DeviceDiagBuilder(Kind K, SourceLocation Loc, unsigned DiagID,
- FunctionDecl *Fn, Sema &S);
- DeviceDiagBuilder(DeviceDiagBuilder &&D);
- DeviceDiagBuilder(const DeviceDiagBuilder &) = default;
- ~DeviceDiagBuilder();
-
- /// Convertible to bool: True if we immediately emitted an error, false if
- /// we didn't emit an error or we created a deferred error.
- ///
- /// Example usage:
- ///
- /// if (DeviceDiagBuilder(...) << foo << bar)
- /// return ExprError();
- ///
- /// But see CUDADiagIfDeviceCode() and CUDADiagIfHostCode() -- you probably
- /// want to use these instead of creating a DeviceDiagBuilder yourself.
- operator bool() const { return ImmediateDiag.hasValue(); }
-
- template <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
@@ -11706,15 +11744,18 @@
/// if (CUDADiagIfDeviceCode(Loc, diag::err_cuda_vla) << CurrentCUDATarget())
/// return ExprError();
/// // Otherwise, continue parsing as normal.
- DeviceDiagBuilder CUDADiagIfDeviceCode(SourceLocation Loc, unsigned DiagID);
+ SemaDiagnosticBuilder CUDADiagIfDeviceCode(SourceLocation Loc,
+ unsigned DiagID,
+ bool EmitOnBothSides = false);
- /// Creates a DeviceDiagBuilder that emits the diagnostic if the current context
- /// is "used as host code".
+ /// Creates a SemaDiagnosticBuilder that emits the diagnostic if the current
+ /// context is "used as host code".
///
/// Same as CUDADiagIfDeviceCode, with "host" and "device" switched.
- DeviceDiagBuilder CUDADiagIfHostCode(SourceLocation Loc, unsigned DiagID);
+ SemaDiagnosticBuilder CUDADiagIfHostCode(SourceLocation Loc, unsigned DiagID,
+ bool EmitOnBothSides = false);
- /// Creates a DeviceDiagBuilder that emits the diagnostic if the current
+ /// Creates a SemaDiagnosticBuilder that emits the diagnostic if the current
/// context is "used as device code".
///
/// - If CurContext is a `declare target` function or it is known that the
@@ -11729,9 +11770,10 @@
/// if (diagIfOpenMPDeviceCode(Loc, diag::err_vla_unsupported))
/// return ExprError();
/// // Otherwise, continue parsing as normal.
- DeviceDiagBuilder diagIfOpenMPDeviceCode(SourceLocation Loc, unsigned DiagID);
+ SemaDiagnosticBuilder diagIfOpenMPDeviceCode(SourceLocation Loc,
+ unsigned DiagID);
- /// Creates a DeviceDiagBuilder that emits the diagnostic if the current
+ /// Creates a SemaDiagnosticBuilder that emits the diagnostic if the current
/// context is "used as host code".
///
/// - If CurContext is a `declare target` function or it is known that the
@@ -11744,9 +11786,14 @@
/// if (diagIfOpenMPHostode(Loc, diag::err_vla_unsupported))
/// return ExprError();
/// // Otherwise, continue parsing as normal.
- DeviceDiagBuilder diagIfOpenMPHostCode(SourceLocation Loc, unsigned DiagID);
+ SemaDiagnosticBuilder diagIfOpenMPHostCode(SourceLocation Loc,
+ unsigned DiagID);
- DeviceDiagBuilder targetDiag(SourceLocation Loc, unsigned DiagID);
+ SemaDiagnosticBuilder targetDiag(SourceLocation Loc, unsigned DiagID);
+ SemaDiagnosticBuilder targetDiag(SourceLocation Loc,
+ const PartialDiagnostic &PD) {
+ return targetDiag(Loc, PD.getDiagID()) << PD;
+ }
/// Check if the expression is allowed to be used in expressions for the
/// offloading devices.
@@ -12518,7 +12565,7 @@
ConstructorDestructor,
BuiltinFunction
};
- /// Creates a DeviceDiagBuilder that emits the diagnostic if the current
+ /// Creates a SemaDiagnosticBuilder that emits the diagnostic if the current
/// context is "used as device code".
///
/// - If CurLexicalContext is a kernel function or it is known that the
@@ -12536,7 +12583,8 @@
/// if (!S.Context.getTargetInfo().hasFloat128Type() &&
/// S.getLangOpts().SYCLIsDevice)
/// SYCLDiagIfDeviceCode(Loc, diag::err_type_unsupported) << "__float128";
- DeviceDiagBuilder SYCLDiagIfDeviceCode(SourceLocation Loc, unsigned DiagID);
+ SemaDiagnosticBuilder SYCLDiagIfDeviceCode(SourceLocation Loc,
+ unsigned DiagID);
/// Check whether we're allowed to call Callee from the current context.
///
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
