aboutsummaryrefslogtreecommitdiff
path: root/clang/lib
diff options
context:
space:
mode:
Diffstat (limited to 'clang/lib')
-rw-r--r--clang/lib/Sema/MultiplexExternalSemaSource.cpp6
-rw-r--r--clang/lib/Sema/Sema.cpp177
-rw-r--r--clang/lib/Sema/SemaCUDA.cpp19
-rw-r--r--clang/lib/Sema/SemaDecl.cpp27
-rw-r--r--clang/lib/Sema/SemaExpr.cpp7
-rw-r--r--clang/lib/Sema/SemaOpenMP.cpp158
-rw-r--r--clang/lib/Sema/UsedDeclVisitor.h12
-rw-r--r--clang/lib/Serialization/ASTReader.cpp18
-rw-r--r--clang/lib/Serialization/ASTWriter.cpp11
9 files changed, 190 insertions, 245 deletions
diff --git a/clang/lib/Sema/MultiplexExternalSemaSource.cpp b/clang/lib/Sema/MultiplexExternalSemaSource.cpp
index 2b0cd6b..80333e6 100644
--- a/clang/lib/Sema/MultiplexExternalSemaSource.cpp
+++ b/clang/lib/Sema/MultiplexExternalSemaSource.cpp
@@ -275,6 +275,12 @@ void MultiplexExternalSemaSource::ReadExtVectorDecls(
Sources[i]->ReadExtVectorDecls(Decls);
}
+void MultiplexExternalSemaSource::ReadDeclsToCheckForDeferredDiags(
+ llvm::SmallVector<Decl *, 4> &Decls) {
+ for(size_t i = 0; i < Sources.size(); ++i)
+ Sources[i]->ReadDeclsToCheckForDeferredDiags(Decls);
+}
+
void MultiplexExternalSemaSource::ReadUnusedLocalTypedefNameCandidates(
llvm::SmallSetVector<const TypedefNameDecl *, 4> &Decls) {
for(size_t i = 0; i < Sources.size(); ++i)
diff --git a/clang/lib/Sema/Sema.cpp b/clang/lib/Sema/Sema.cpp
index 186c0c7..0347524 100644
--- a/clang/lib/Sema/Sema.cpp
+++ b/clang/lib/Sema/Sema.cpp
@@ -11,6 +11,7 @@
//
//===----------------------------------------------------------------------===//
+#include "UsedDeclVisitor.h"
#include "clang/AST/ASTContext.h"
#include "clang/AST/ASTDiagnostic.h"
#include "clang/AST/DeclCXX.h"
@@ -955,9 +956,7 @@ void Sema::ActOnEndOfTranslationUnitFragment(TUFragmentKind Kind) {
PerformPendingInstantiations();
}
- // Finalize analysis of OpenMP-specific constructs.
- if (LangOpts.OpenMP)
- finalizeOpenMPDelayedAnalysis();
+ emitDeferredDiags();
assert(LateParsedInstantiations.empty() &&
"end of TU template instantiation should not create more "
@@ -1452,27 +1451,108 @@ static void emitCallStackNotes(Sema &S, FunctionDecl *FD) {
// Emit any deferred diagnostics for FD and erase them from the map in which
// they're stored.
-static void emitDeferredDiags(Sema &S, FunctionDecl *FD, bool ShowCallStack) {
- auto It = S.DeviceDeferredDiags.find(FD);
- if (It == S.DeviceDeferredDiags.end())
+void Sema::emitDeferredDiags(FunctionDecl *FD, bool ShowCallStack) {
+ auto It = DeviceDeferredDiags.find(FD);
+ if (It == DeviceDeferredDiags.end())
return;
bool HasWarningOrError = false;
+ bool FirstDiag = true;
for (PartialDiagnosticAt &PDAt : It->second) {
const SourceLocation &Loc = PDAt.first;
const PartialDiagnostic &PD = PDAt.second;
- HasWarningOrError |= S.getDiagnostics().getDiagnosticLevel(
+ HasWarningOrError |= getDiagnostics().getDiagnosticLevel(
PD.getDiagID(), Loc) >= DiagnosticsEngine::Warning;
- DiagnosticBuilder Builder(S.Diags.Report(Loc, PD.getDiagID()));
- Builder.setForceEmit();
- PD.Emit(Builder);
+ {
+ DiagnosticBuilder Builder(Diags.Report(Loc, PD.getDiagID()));
+ Builder.setForceEmit();
+ PD.Emit(Builder);
+ }
+
+ // Emit the note on the first diagnostic in case too many diagnostics cause
+ // the note not emitted.
+ if (FirstDiag && HasWarningOrError && ShowCallStack) {
+ emitCallStackNotes(*this, FD);
+ FirstDiag = false;
+ }
}
- S.DeviceDeferredDiags.erase(It);
- // FIXME: Should this be called after every warning/error emitted in the loop
- // above, instead of just once per function? That would be consistent with
- // how we handle immediate errors, but it also seems like a bit much.
- if (HasWarningOrError && ShowCallStack)
- emitCallStackNotes(S, FD);
+}
+
+namespace {
+/// Helper class that emits deferred diagnostic messages if an entity directly
+/// or indirectly using the function that causes the deferred diagnostic
+/// messages is known to be emitted.
+class DeferredDiagnosticsEmitter
+ : public UsedDeclVisitor<DeferredDiagnosticsEmitter> {
+public:
+ typedef UsedDeclVisitor<DeferredDiagnosticsEmitter> Inherited;
+ llvm::SmallSet<CanonicalDeclPtr<Decl>, 4> Visited;
+ llvm::SmallVector<CanonicalDeclPtr<FunctionDecl>, 4> UseStack;
+ bool ShouldEmit;
+ unsigned InOMPDeviceContext;
+
+ DeferredDiagnosticsEmitter(Sema &S)
+ : Inherited(S), ShouldEmit(false), InOMPDeviceContext(0) {}
+
+ void VisitOMPTargetDirective(OMPTargetDirective *Node) {
+ ++InOMPDeviceContext;
+ Inherited::VisitOMPTargetDirective(Node);
+ --InOMPDeviceContext;
+ }
+
+ void visitUsedDecl(SourceLocation Loc, Decl *D) {
+ if (auto *FD = dyn_cast<FunctionDecl>(D)) {
+ FunctionDecl *Caller = UseStack.empty() ? nullptr : UseStack.back();
+ auto IsKnownEmitted = S.getEmissionStatus(FD, /*Final=*/true) ==
+ Sema::FunctionEmissionStatus::Emitted;
+ if (!Caller)
+ ShouldEmit = IsKnownEmitted;
+ if ((!ShouldEmit && !S.getLangOpts().OpenMP && !Caller) ||
+ S.shouldIgnoreInHostDeviceCheck(FD) || Visited.count(D))
+ return;
+ // Finalize analysis of OpenMP-specific constructs.
+ if (Caller && S.LangOpts.OpenMP && UseStack.size() == 1)
+ S.finalizeOpenMPDelayedAnalysis(Caller, FD, Loc);
+ if (Caller)
+ S.DeviceKnownEmittedFns[FD] = {Caller, Loc};
+ if (ShouldEmit || InOMPDeviceContext)
+ S.emitDeferredDiags(FD, Caller);
+ Visited.insert(D);
+ UseStack.push_back(FD);
+ if (auto *S = FD->getBody()) {
+ this->Visit(S);
+ }
+ UseStack.pop_back();
+ Visited.erase(D);
+ } else if (auto *VD = dyn_cast<VarDecl>(D)) {
+ if (auto *Init = VD->getInit()) {
+ auto DevTy = OMPDeclareTargetDeclAttr::getDeviceType(VD);
+ bool IsDev = DevTy && (*DevTy == OMPDeclareTargetDeclAttr::DT_NoHost ||
+ *DevTy == OMPDeclareTargetDeclAttr::DT_Any);
+ if (IsDev)
+ ++InOMPDeviceContext;
+ this->Visit(Init);
+ if (IsDev)
+ --InOMPDeviceContext;
+ }
+ } else
+ Inherited::visitUsedDecl(Loc, D);
+ }
+};
+} // namespace
+
+void Sema::emitDeferredDiags() {
+ if (ExternalSource)
+ ExternalSource->ReadDeclsToCheckForDeferredDiags(
+ DeclsToCheckForDeferredDiags);
+
+ if ((DeviceDeferredDiags.empty() && !LangOpts.OpenMP) ||
+ DeclsToCheckForDeferredDiags.empty())
+ return;
+
+ DeferredDiagnosticsEmitter DDE(*this);
+ for (auto D : DeclsToCheckForDeferredDiags)
+ DDE.visitUsedDecl(SourceLocation(), D);
}
// In CUDA, there are some constructs which may appear in semantically-valid
@@ -1545,71 +1625,6 @@ Sema::DeviceDiagBuilder::~DeviceDiagBuilder() {
}
}
-// Indicate that this function (and thus everything it transtively calls) will
-// be codegen'ed, and emit any deferred diagnostics on this function and its
-// (transitive) callees.
-void Sema::markKnownEmitted(
- Sema &S, FunctionDecl *OrigCaller, FunctionDecl *OrigCallee,
- SourceLocation OrigLoc,
- const llvm::function_ref<bool(Sema &, FunctionDecl *)> IsKnownEmitted) {
- // Nothing to do if we already know that FD is emitted.
- if (IsKnownEmitted(S, OrigCallee)) {
- assert(!S.DeviceCallGraph.count(OrigCallee));
- return;
- }
-
- // We've just discovered that OrigCallee is known-emitted. Walk our call
- // graph to see what else we can now discover also must be emitted.
-
- struct CallInfo {
- FunctionDecl *Caller;
- FunctionDecl *Callee;
- SourceLocation Loc;
- };
- llvm::SmallVector<CallInfo, 4> Worklist = {{OrigCaller, OrigCallee, OrigLoc}};
- llvm::SmallSet<CanonicalDeclPtr<FunctionDecl>, 4> Seen;
- Seen.insert(OrigCallee);
- while (!Worklist.empty()) {
- CallInfo C = Worklist.pop_back_val();
- assert(!IsKnownEmitted(S, C.Callee) &&
- "Worklist should not contain known-emitted functions.");
- S.DeviceKnownEmittedFns[C.Callee] = {C.Caller, C.Loc};
- emitDeferredDiags(S, C.Callee, C.Caller);
-
- // If this is a template instantiation, explore its callgraph as well:
- // Non-dependent calls are part of the template's callgraph, while dependent
- // calls are part of to the instantiation's call graph.
- if (auto *Templ = C.Callee->getPrimaryTemplate()) {
- FunctionDecl *TemplFD = Templ->getAsFunction();
- if (!Seen.count(TemplFD) && !S.DeviceKnownEmittedFns.count(TemplFD)) {
- Seen.insert(TemplFD);
- Worklist.push_back(
- {/* Caller = */ C.Caller, /* Callee = */ TemplFD, C.Loc});
- }
- }
-
- // Add all functions called by Callee to our worklist.
- auto CGIt = S.DeviceCallGraph.find(C.Callee);
- if (CGIt == S.DeviceCallGraph.end())
- continue;
-
- for (std::pair<CanonicalDeclPtr<FunctionDecl>, SourceLocation> FDLoc :
- CGIt->second) {
- FunctionDecl *NewCallee = FDLoc.first;
- SourceLocation CallLoc = FDLoc.second;
- if (Seen.count(NewCallee) || IsKnownEmitted(S, NewCallee))
- continue;
- Seen.insert(NewCallee);
- Worklist.push_back(
- {/* Caller = */ C.Callee, /* Callee = */ NewCallee, CallLoc});
- }
-
- // C.Callee is now known-emitted, so we no longer need to maintain its list
- // of callees in DeviceCallGraph.
- S.DeviceCallGraph.erase(CGIt);
- }
-}
-
Sema::DeviceDiagBuilder Sema::targetDiag(SourceLocation Loc, unsigned DiagID) {
if (LangOpts.OpenMP)
return LangOpts.OpenMPIsDevice ? diagIfOpenMPDeviceCode(Loc, DiagID)
diff --git a/clang/lib/Sema/SemaCUDA.cpp b/clang/lib/Sema/SemaCUDA.cpp
index 5bea9c8..faab250e 100644
--- a/clang/lib/Sema/SemaCUDA.cpp
+++ b/clang/lib/Sema/SemaCUDA.cpp
@@ -675,25 +675,6 @@ bool Sema::CheckCUDACall(SourceLocation Loc, FunctionDecl *Callee) {
// Otherwise, mark the call in our call graph so we can traverse it later.
bool CallerKnownEmitted =
getEmissionStatus(Caller) == FunctionEmissionStatus::Emitted;
- if (CallerKnownEmitted) {
- // Host-side references to a __global__ function refer to the stub, so the
- // function itself is never emitted and therefore should not be marked.
- if (!shouldIgnoreInHostDeviceCheck(Callee))
- markKnownEmitted(
- *this, Caller, Callee, Loc, [](Sema &S, FunctionDecl *FD) {
- return S.getEmissionStatus(FD) == FunctionEmissionStatus::Emitted;
- });
- } else {
- // If we have
- // host fn calls kernel fn calls host+device,
- // the HD function does not get instantiated on the host. We model this by
- // omitting at the call to the kernel from the callgraph. This ensures
- // that, when compiling for host, only HD functions actually called from the
- // host get marked as known-emitted.
- if (!shouldIgnoreInHostDeviceCheck(Callee))
- DeviceCallGraph[Caller].insert({Callee, Loc});
- }
-
DeviceDiagBuilder::Kind DiagKind = [this, Caller, Callee,
CallerKnownEmitted] {
switch (IdentifyCUDAPreference(Caller, Callee)) {
diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp
index b33596e..225f720 100644
--- a/clang/lib/Sema/SemaDecl.cpp
+++ b/clang/lib/Sema/SemaDecl.cpp
@@ -12247,6 +12247,8 @@ void Sema::AddInitializerToDecl(Decl *RealDecl, Expr *Init, bool DirectInit) {
VDecl->setInitStyle(VarDecl::ListInit);
}
+ if (LangOpts.OpenMP && VDecl->hasGlobalStorage())
+ DeclsToCheckForDeferredDiags.push_back(VDecl);
CheckCompleteVariableDeclaration(VDecl);
}
@@ -14363,6 +14365,13 @@ Decl *Sema::ActOnFinishFunctionBody(Decl *dcl, Stmt *Body,
DiscardCleanupsInEvaluationContext();
}
+ if (LangOpts.OpenMP || LangOpts.CUDA) {
+ auto ES = getEmissionStatus(FD);
+ if (ES == Sema::FunctionEmissionStatus::Emitted ||
+ ES == Sema::FunctionEmissionStatus::Unknown)
+ DeclsToCheckForDeferredDiags.push_back(FD);
+ }
+
return dcl;
}
@@ -18026,7 +18035,8 @@ Decl *Sema::getObjCDeclContext() const {
return (dyn_cast_or_null<ObjCContainerDecl>(CurContext));
}
-Sema::FunctionEmissionStatus Sema::getEmissionStatus(FunctionDecl *FD) {
+Sema::FunctionEmissionStatus Sema::getEmissionStatus(FunctionDecl *FD,
+ bool Final) {
// Templates are emitted when they're instantiated.
if (FD->isDependentContext())
return FunctionEmissionStatus::TemplateDiscarded;
@@ -18038,8 +18048,10 @@ Sema::FunctionEmissionStatus Sema::getEmissionStatus(FunctionDecl *FD) {
if (DevTy.hasValue()) {
if (*DevTy == OMPDeclareTargetDeclAttr::DT_Host)
OMPES = FunctionEmissionStatus::OMPDiscarded;
- else if (DeviceKnownEmittedFns.count(FD) > 0)
+ else if (*DevTy == OMPDeclareTargetDeclAttr::DT_NoHost ||
+ *DevTy == OMPDeclareTargetDeclAttr::DT_Any) {
OMPES = FunctionEmissionStatus::Emitted;
+ }
}
} else if (LangOpts.OpenMP) {
// In OpenMP 4.5 all the functions are host functions.
@@ -18055,10 +18067,11 @@ Sema::FunctionEmissionStatus Sema::getEmissionStatus(FunctionDecl *FD) {
if (DevTy.hasValue()) {
if (*DevTy == OMPDeclareTargetDeclAttr::DT_NoHost) {
OMPES = FunctionEmissionStatus::OMPDiscarded;
- } else if (DeviceKnownEmittedFns.count(FD) > 0) {
+ } else if (*DevTy == OMPDeclareTargetDeclAttr::DT_Host ||
+ *DevTy == OMPDeclareTargetDeclAttr::DT_Any)
OMPES = FunctionEmissionStatus::Emitted;
- }
- }
+ } else if (Final)
+ OMPES = FunctionEmissionStatus::Emitted;
}
}
if (OMPES == FunctionEmissionStatus::OMPDiscarded ||
@@ -18093,9 +18106,7 @@ Sema::FunctionEmissionStatus Sema::getEmissionStatus(FunctionDecl *FD) {
// Otherwise, the function is known-emitted if it's in our set of
// known-emitted functions.
- return (DeviceKnownEmittedFns.count(FD) > 0)
- ? FunctionEmissionStatus::Emitted
- : FunctionEmissionStatus::Unknown;
+ return FunctionEmissionStatus::Unknown;
}
bool Sema::shouldIgnoreInHostDeviceCheck(FunctionDecl *Callee) {
diff --git a/clang/lib/Sema/SemaExpr.cpp b/clang/lib/Sema/SemaExpr.cpp
index a176125..afffd96 100644
--- a/clang/lib/Sema/SemaExpr.cpp
+++ b/clang/lib/Sema/SemaExpr.cpp
@@ -16017,13 +16017,8 @@ void Sema::MarkFunctionReferenced(SourceLocation Loc, FunctionDecl *Func,
Func->markUsed(Context);
}
- if (LangOpts.OpenMP) {
+ if (LangOpts.OpenMP)
markOpenMPDeclareVariantFuncsReferenced(Loc, Func, MightBeOdrUse);
- if (LangOpts.OpenMPIsDevice)
- checkOpenMPDeviceFunction(Loc, Func);
- else
- checkOpenMPHostFunction(Loc, Func);
- }
}
/// Directly mark a variable odr-used. Given a choice, prefer to use
diff --git a/clang/lib/Sema/SemaOpenMP.cpp b/clang/lib/Sema/SemaOpenMP.cpp
index 309ab4a..8c488da 100644
--- a/clang/lib/Sema/SemaOpenMP.cpp
+++ b/clang/lib/Sema/SemaOpenMP.cpp
@@ -1785,92 +1785,6 @@ Sema::DeviceDiagBuilder Sema::diagIfOpenMPHostCode(SourceLocation Loc,
return DeviceDiagBuilder(Kind, Loc, DiagID, getCurFunctionDecl(), *this);
}
-void Sema::checkOpenMPDeviceFunction(SourceLocation Loc, FunctionDecl *Callee,
- bool CheckForDelayedContext) {
- assert(LangOpts.OpenMP && LangOpts.OpenMPIsDevice &&
- "Expected OpenMP device compilation.");
- assert(Callee && "Callee may not be null.");
- Callee = Callee->getMostRecentDecl();
- FunctionDecl *Caller = getCurFunctionDecl();
-
- // host only function are not available on the device.
- if (Caller) {
- FunctionEmissionStatus CallerS = getEmissionStatus(Caller);
- FunctionEmissionStatus CalleeS = getEmissionStatus(Callee);
- assert(CallerS != FunctionEmissionStatus::CUDADiscarded &&
- CalleeS != FunctionEmissionStatus::CUDADiscarded &&
- "CUDADiscarded unexpected in OpenMP device function check");
- if ((CallerS == FunctionEmissionStatus::Emitted ||
- (!isOpenMPDeviceDelayedContext(*this) &&
- CallerS == FunctionEmissionStatus::Unknown)) &&
- CalleeS == FunctionEmissionStatus::OMPDiscarded) {
- StringRef HostDevTy = getOpenMPSimpleClauseTypeName(
- OMPC_device_type, OMPC_DEVICE_TYPE_host);
- Diag(Loc, diag::err_omp_wrong_device_function_call) << HostDevTy << 0;
- Diag(Callee->getAttr<OMPDeclareTargetDeclAttr>()->getLocation(),
- diag::note_omp_marked_device_type_here)
- << HostDevTy;
- return;
- }
- }
- // If the caller is known-emitted, mark the callee as known-emitted.
- // Otherwise, mark the call in our call graph so we can traverse it later.
- if ((CheckForDelayedContext && !isOpenMPDeviceDelayedContext(*this)) ||
- (!Caller && !CheckForDelayedContext) ||
- (Caller && getEmissionStatus(Caller) == FunctionEmissionStatus::Emitted))
- markKnownEmitted(*this, Caller, Callee, Loc,
- [CheckForDelayedContext](Sema &S, FunctionDecl *FD) {
- return CheckForDelayedContext &&
- S.getEmissionStatus(FD) ==
- FunctionEmissionStatus::Emitted;
- });
- else if (Caller)
- DeviceCallGraph[Caller].insert({Callee, Loc});
-}
-
-void Sema::checkOpenMPHostFunction(SourceLocation Loc, FunctionDecl *Callee,
- bool CheckCaller) {
- assert(LangOpts.OpenMP && !LangOpts.OpenMPIsDevice &&
- "Expected OpenMP host compilation.");
- assert(Callee && "Callee may not be null.");
- Callee = Callee->getMostRecentDecl();
- FunctionDecl *Caller = getCurFunctionDecl();
-
- // device only function are not available on the host.
- if (Caller) {
- FunctionEmissionStatus CallerS = getEmissionStatus(Caller);
- FunctionEmissionStatus CalleeS = getEmissionStatus(Callee);
- assert(
- (LangOpts.CUDA || (CallerS != FunctionEmissionStatus::CUDADiscarded &&
- CalleeS != FunctionEmissionStatus::CUDADiscarded)) &&
- "CUDADiscarded unexpected in OpenMP host function check");
- if (CallerS == FunctionEmissionStatus::Emitted &&
- CalleeS == FunctionEmissionStatus::OMPDiscarded) {
- StringRef NoHostDevTy = getOpenMPSimpleClauseTypeName(
- OMPC_device_type, OMPC_DEVICE_TYPE_nohost);
- Diag(Loc, diag::err_omp_wrong_device_function_call) << NoHostDevTy << 1;
- Diag(Callee->getAttr<OMPDeclareTargetDeclAttr>()->getLocation(),
- diag::note_omp_marked_device_type_here)
- << NoHostDevTy;
- return;
- }
- }
- // If the caller is known-emitted, mark the callee as known-emitted.
- // Otherwise, mark the call in our call graph so we can traverse it later.
- if (!shouldIgnoreInHostDeviceCheck(Callee)) {
- if ((!CheckCaller && !Caller) ||
- (Caller &&
- getEmissionStatus(Caller) == FunctionEmissionStatus::Emitted))
- markKnownEmitted(
- *this, Caller, Callee, Loc, [CheckCaller](Sema &S, FunctionDecl *FD) {
- return CheckCaller &&
- S.getEmissionStatus(FD) == FunctionEmissionStatus::Emitted;
- });
- else if (Caller)
- DeviceCallGraph[Caller].insert({Callee, Loc});
- }
-}
-
void Sema::checkOpenMPDeviceExpr(const Expr *E) {
assert(getLangOpts().OpenMP && getLangOpts().OpenMPIsDevice &&
"OpenMP device compilation mode is expected.");
@@ -2330,52 +2244,43 @@ bool Sema::isOpenMPGlobalCapturedDecl(ValueDecl *D, unsigned Level,
void Sema::DestroyDataSharingAttributesStack() { delete DSAStack; }
-void Sema::finalizeOpenMPDelayedAnalysis() {
+void Sema::finalizeOpenMPDelayedAnalysis(const FunctionDecl *Caller,
+ const FunctionDecl *Callee,
+ SourceLocation Loc) {
assert(LangOpts.OpenMP && "Expected OpenMP compilation mode.");
- // Diagnose implicit declare target functions and their callees.
- for (const auto &CallerCallees : DeviceCallGraph) {
- Optional<OMPDeclareTargetDeclAttr::DevTypeTy> DevTy =
- OMPDeclareTargetDeclAttr::getDeviceType(
- CallerCallees.getFirst()->getMostRecentDecl());
- // Ignore host functions during device analyzis.
- if (LangOpts.OpenMPIsDevice && DevTy &&
- *DevTy == OMPDeclareTargetDeclAttr::DT_Host)
- continue;
- // Ignore nohost functions during host analyzis.
- if (!LangOpts.OpenMPIsDevice && DevTy &&
- *DevTy == OMPDeclareTargetDeclAttr::DT_NoHost)
- continue;
- for (const std::pair<CanonicalDeclPtr<FunctionDecl>, SourceLocation>
- &Callee : CallerCallees.getSecond()) {
- const FunctionDecl *FD = Callee.first->getMostRecentDecl();
- Optional<OMPDeclareTargetDeclAttr::DevTypeTy> DevTy =
- OMPDeclareTargetDeclAttr::getDeviceType(FD);
- if (LangOpts.OpenMPIsDevice && DevTy &&
- *DevTy == OMPDeclareTargetDeclAttr::DT_Host) {
- // Diagnose host function called during device codegen.
- StringRef HostDevTy = getOpenMPSimpleClauseTypeName(
- OMPC_device_type, OMPC_DEVICE_TYPE_host);
- Diag(Callee.second, diag::err_omp_wrong_device_function_call)
- << HostDevTy << 0;
- Diag(FD->getAttr<OMPDeclareTargetDeclAttr>()->getLocation(),
- diag::note_omp_marked_device_type_here)
- << HostDevTy;
- continue;
- }
+ Optional<OMPDeclareTargetDeclAttr::DevTypeTy> DevTy =
+ OMPDeclareTargetDeclAttr::getDeviceType(Caller->getMostRecentDecl());
+ // Ignore host functions during device analyzis.
+ if (LangOpts.OpenMPIsDevice && DevTy &&
+ *DevTy == OMPDeclareTargetDeclAttr::DT_Host)
+ return;
+ // Ignore nohost functions during host analyzis.
+ if (!LangOpts.OpenMPIsDevice && DevTy &&
+ *DevTy == OMPDeclareTargetDeclAttr::DT_NoHost)
+ return;
+ const FunctionDecl *FD = Callee->getMostRecentDecl();
+ DevTy = OMPDeclareTargetDeclAttr::getDeviceType(FD);
+ if (LangOpts.OpenMPIsDevice && DevTy &&
+ *DevTy == OMPDeclareTargetDeclAttr::DT_Host) {
+ // Diagnose host function called during device codegen.
+ StringRef HostDevTy =
+ getOpenMPSimpleClauseTypeName(OMPC_device_type, OMPC_DEVICE_TYPE_host);
+ Diag(Loc, diag::err_omp_wrong_device_function_call) << HostDevTy << 0;
+ Diag(FD->getAttr<OMPDeclareTargetDeclAttr>()->getLocation(),
+ diag::note_omp_marked_device_type_here)
+ << HostDevTy;
+ return;
+ }
if (!LangOpts.OpenMPIsDevice && DevTy &&
*DevTy == OMPDeclareTargetDeclAttr::DT_NoHost) {
// Diagnose nohost function called during host codegen.
StringRef NoHostDevTy = getOpenMPSimpleClauseTypeName(
OMPC_device_type, OMPC_DEVICE_TYPE_nohost);
- Diag(Callee.second, diag::err_omp_wrong_device_function_call)
- << NoHostDevTy << 1;
+ Diag(Loc, diag::err_omp_wrong_device_function_call) << NoHostDevTy << 1;
Diag(FD->getAttr<OMPDeclareTargetDeclAttr>()->getLocation(),
diag::note_omp_marked_device_type_here)
<< NoHostDevTy;
- continue;
}
- }
- }
}
void Sema::StartOpenMPDSABlock(OpenMPDirectiveKind DKind,
@@ -17747,15 +17652,6 @@ void Sema::checkDeclIsAllowedInOpenMPTarget(Expr *E, Decl *D,
Diag(FD->getLocation(), diag::note_defined_here) << FD;
return;
}
- // Mark the function as must be emitted for the device.
- Optional<OMPDeclareTargetDeclAttr::DevTypeTy> DevTy =
- OMPDeclareTargetDeclAttr::getDeviceType(FD);
- if (LangOpts.OpenMPIsDevice && Res.hasValue() && IdLoc.isValid() &&
- *DevTy != OMPDeclareTargetDeclAttr::DT_Host)
- checkOpenMPDeviceFunction(IdLoc, FD, /*CheckForDelayedContext=*/false);
- if (!LangOpts.OpenMPIsDevice && Res.hasValue() && IdLoc.isValid() &&
- *DevTy != OMPDeclareTargetDeclAttr::DT_NoHost)
- checkOpenMPHostFunction(IdLoc, FD, /*CheckCaller=*/false);
}
if (auto *VD = dyn_cast<ValueDecl>(D)) {
// Problem if any with var declared with incomplete type will be reported
diff --git a/clang/lib/Sema/UsedDeclVisitor.h b/clang/lib/Sema/UsedDeclVisitor.h
index be46f0d..d207e07 100644
--- a/clang/lib/Sema/UsedDeclVisitor.h
+++ b/clang/lib/Sema/UsedDeclVisitor.h
@@ -84,6 +84,18 @@ public:
void VisitCXXDefaultArgExpr(CXXDefaultArgExpr *E) {
asImpl().Visit(E->getExpr());
}
+
+ void visitUsedDecl(SourceLocation Loc, Decl *D) {
+ if (auto *CD = dyn_cast<CapturedDecl>(D)) {
+ if (auto *S = CD->getBody()) {
+ asImpl().Visit(S);
+ }
+ } else if (auto *CD = dyn_cast<BlockDecl>(D)) {
+ if (auto *S = CD->getBody()) {
+ asImpl().Visit(S);
+ }
+ }
+ }
};
} // end namespace clang
diff --git a/clang/lib/Serialization/ASTReader.cpp b/clang/lib/Serialization/ASTReader.cpp
index 12a8061..ac11177 100644
--- a/clang/lib/Serialization/ASTReader.cpp
+++ b/clang/lib/Serialization/ASTReader.cpp
@@ -3773,6 +3773,11 @@ ASTReader::ReadASTBlock(ModuleFile &F, unsigned ClientLoadCapabilities) {
}
break;
}
+
+ case DECLS_TO_CHECK_FOR_DEFERRED_DIAGS:
+ for (unsigned I = 0, N = Record.size(); I != N; ++I)
+ DeclsToCheckForDeferredDiags.push_back(getGlobalDeclID(F, Record[I]));
+ break;
}
}
}
@@ -8180,6 +8185,19 @@ void ASTReader::ReadUnusedLocalTypedefNameCandidates(
UnusedLocalTypedefNameCandidates.clear();
}
+void ASTReader::ReadDeclsToCheckForDeferredDiags(
+ llvm::SmallVector<Decl *, 4> &Decls) {
+ for (unsigned I = 0, N = DeclsToCheckForDeferredDiags.size(); I != N;
+ ++I) {
+ auto *D = dyn_cast_or_null<Decl>(
+ GetDecl(DeclsToCheckForDeferredDiags[I]));
+ if (D)
+ Decls.push_back(D);
+ }
+ DeclsToCheckForDeferredDiags.clear();
+}
+
+
void ASTReader::ReadReferencedSelectors(
SmallVectorImpl<std::pair<Selector, SourceLocation>> &Sels) {
if (ReferencedSelectorsData.empty())
diff --git a/clang/lib/Serialization/ASTWriter.cpp b/clang/lib/Serialization/ASTWriter.cpp
index c95cabc..a5d5b0f 100644
--- a/clang/lib/Serialization/ASTWriter.cpp
+++ b/clang/lib/Serialization/ASTWriter.cpp
@@ -756,6 +756,7 @@ void ASTWriter::WriteBlockInfoBlock() {
RECORD(DELETE_EXPRS_TO_ANALYZE);
RECORD(CUDA_PRAGMA_FORCE_HOST_DEVICE_DEPTH);
RECORD(PP_CONDITIONAL_STACK);
+ RECORD(DECLS_TO_CHECK_FOR_DEFERRED_DIAGS);
// SourceManager Block.
BLOCK(SOURCE_MANAGER_BLOCK);
@@ -4671,6 +4672,11 @@ ASTFileSignature ASTWriter::WriteASTCore(Sema &SemaRef, StringRef isysroot,
Buffer.data(), Buffer.size());
}
+ // Build a record containing all of the DeclsToCheckForDeferredDiags.
+ RecordData DeclsToCheckForDeferredDiags;
+ for (auto *D : SemaRef.DeclsToCheckForDeferredDiags)
+ AddDeclRef(D, DeclsToCheckForDeferredDiags);
+
RecordData DeclUpdatesOffsetsRecord;
// Keep writing types, declarations, and declaration update records
@@ -4762,6 +4768,11 @@ ASTFileSignature ASTWriter::WriteASTCore(Sema &SemaRef, StringRef isysroot,
if (!SemaDeclRefs.empty())
Stream.EmitRecord(SEMA_DECL_REFS, SemaDeclRefs);
+ // Write the record containing decls to be checked for deferred diags.
+ if (!DeclsToCheckForDeferredDiags.empty())
+ Stream.EmitRecord(DECLS_TO_CHECK_FOR_DEFERRED_DIAGS,
+ DeclsToCheckForDeferredDiags);
+
// Write the record containing CUDA-specific declaration references.
if (!CUDASpecialDeclRefs.empty())
Stream.EmitRecord(CUDA_SPECIAL_DECL_REFS, CUDASpecialDeclRefs);
generated by cgit v1.1 at 2025-08-03 03:45:57 +0000