Magic-Mirror/llvm-capstone
1
0
Fork 0
mirror of https://github.com/capstone-engine/llvm-capstone.git synced 2025-02-03 07:38:57 +00:00
Code Issues Actions Packages Projects Releases Wiki Activity

recommit 1b978ddba05c [CUDA][HIP][OpenMP] Emit deferred diagnostics by a post-parsing AST travese

Browse Source 
Differential Revision: https://reviews.llvm.org/D70172
This commit is contained in:
Yaxun (Sam) Liu 2020-02-26 10:57:39 -05:00
parent 278c00c4ff
commit b670ab7b6b
21 changed files with 258 additions and 297 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

9
clang/include/clang/Sema/ExternalSemaSource.h
View File

@ -193,6 +193,15 @@ public:
llvm::MapVector<const FunctionDecl *, std::unique_ptr<LateParsedTemplate>>
&LPTMap) {}
/// Read the set of decls to be checked for deferred diags.
///
/// The external source should append its own potentially emitted function
/// and variable decls which may cause deferred diags. Note that this routine
/// may be invoked multiple times; the external source should take care not to
/// introduce the same declarations repeatedly.
virtual void ReadDeclsToCheckForDeferredDiags(
llvm::SmallVector<Decl *, 4> &Decls) {}
/// \copydoc Sema::CorrectTypo
/// \note LookupKind must correspond to a valid Sema::LookupNameKind
///

9
clang/include/clang/Sema/MultiplexExternalSemaSource.h
View File

@ -332,6 +332,15 @@ public:
llvm::MapVector<const FunctionDecl *, std::unique_ptr<LateParsedTemplate>>
&LPTMap) override;
/// Read the set of decls to be checked for deferred diags.
///
/// The external source should append its own potentially emitted function
/// and variable decls which may cause deferred diags. Note that this routine
/// may be invoked multiple times; the external source should take care not to
/// introduce the same declarations repeatedly.
void ReadDeclsToCheckForDeferredDiags(
llvm::SmallVector<Decl *, 4> &Decls) override;
/// \copydoc ExternalSemaSource::CorrectTypo
/// \note Returns the first nonempty correction.
TypoCorrection CorrectTypo(const DeclarationNameInfo &Typo,

52
clang/include/clang/Sema/Sema.h
View File

@ -1492,6 +1492,18 @@ public:
void emitAndClearUnusedLocalTypedefWarnings();
private:
/// Function or variable declarations to be checked for whether the deferred
/// diagnostics should be emitted.
SmallVector<Decl *, 4> DeclsToCheckForDeferredDiags;
public:
// Emit all deferred diagnostics.
void emitDeferredDiags();
// Emit any deferred diagnostics for FD and erase them from the map in which
// they're stored.
void emitDeferredDiags(FunctionDecl *FD, bool ShowCallStack);
enum TUFragmentKind {
/// The global module fragment, between 'module;' and a module-declaration.
Global,
@ -3767,7 +3779,8 @@ public:
TemplateDiscarded, // Discarded due to uninstantiated templates
Unknown,
};
FunctionEmissionStatus getEmissionStatus(FunctionDecl *Decl);
FunctionEmissionStatus getEmissionStatus(FunctionDecl *Decl,
bool Final = false);
// Whether the callee should be ignored in CUDA/HIP/OpenMP host/device check.
bool shouldIgnoreInHostDeviceCheck(FunctionDecl *Callee);
@ -9767,22 +9780,10 @@ private:
/// Pop OpenMP function region for non-capturing function.
void popOpenMPFunctionRegion(const sema::FunctionScopeInfo *OldFSI);
/// Check whether we're allowed to call Callee from the current function.
void checkOpenMPDeviceFunction(SourceLocation Loc, FunctionDecl *Callee,
bool CheckForDelayedContext = true);
/// Check whether we're allowed to call Callee from the current function.
void checkOpenMPHostFunction(SourceLocation Loc, FunctionDecl *Callee,
bool CheckCaller = true);
/// Check if the expression is allowed to be used in expressions for the
/// OpenMP devices.
void checkOpenMPDeviceExpr(const Expr *E);
/// Finishes analysis of the deferred functions calls that may be declared as
/// host/nohost during device/host compilation.
void finalizeOpenMPDelayedAnalysis();
/// Checks if a type or a declaration is disabled due to the owning extension
/// being disabled, and emits diagnostic messages if it is disabled.
/// \param D type or declaration to be checked.
@ -9973,6 +9974,11 @@ public:
void
checkDeclIsAllowedInOpenMPTarget(Expr *E, Decl *D,
SourceLocation IdLoc = SourceLocation());
/// Finishes analysis of the deferred functions calls that may be declared as
/// host/nohost during device/host compilation.
void finalizeOpenMPDelayedAnalysis(const FunctionDecl *Caller,
const FunctionDecl *Callee,
SourceLocation Loc);
/// Return true inside OpenMP declare target region.
bool isInOpenMPDeclareTargetContext() const {
return DeclareTargetNestingLevel > 0;
@ -11359,18 +11365,6 @@ public:
/* Caller = */ FunctionDeclAndLoc>
DeviceKnownEmittedFns;
/// A partial call graph maintained during CUDA/OpenMP device code compilation
/// to support deferred diagnostics.
///
/// Functions are only added here if, at the time they're considered, they are
/// not known-emitted. As soon as we discover that a function is
/// known-emitted, we remove it and everything it transitively calls from this
/// set and add those functions to DeviceKnownEmittedFns.
llvm::DenseMap</* Caller = */ CanonicalDeclPtr<FunctionDecl>,
/* Callees = */ llvm::MapVector<CanonicalDeclPtr<FunctionDecl>,
SourceLocation>>
DeviceCallGraph;
/// Diagnostic builder for CUDA/OpenMP devices errors which may or may not be
/// deferred.
///
@ -11445,14 +11439,6 @@ public:
llvm::Optional<unsigned> PartialDiagId;
};
/// 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 markKnownEmitted(
Sema &S, FunctionDecl *OrigCaller, FunctionDecl *OrigCallee,
SourceLocation OrigLoc,
const llvm::function_ref<bool(Sema &, FunctionDecl *)> IsKnownEmitted);
/// Creates a DeviceDiagBuilder that emits the diagnostic if the current context
/// is "used as device code".
///

5
clang/include/clang/Serialization/ASTBitCodes.h
View File

@ -650,7 +650,10 @@ namespace serialization {
PP_CONDITIONAL_STACK = 62,
/// A table of skipped ranges within the preprocessing record.
PPD_SKIPPED_RANGES = 63
PPD_SKIPPED_RANGES = 63,
/// Record code for the Decls to be checked for deferred diags.
DECLS_TO_CHECK_FOR_DEFERRED_DIAGS = 64,
};
/// Record types used within a source manager block.

9
clang/include/clang/Serialization/ASTReader.h
View File

@ -890,6 +890,12 @@ private:
// A list of late parsed template function data.
SmallVector<uint64_t, 1> LateParsedTemplates;
/// The IDs of all decls to be checked for deferred diags.
///
/// Sema tracks these to emit deferred diags.
SmallVector<uint64_t, 4> DeclsToCheckForDeferredDiags;
public:
struct ImportedSubmodule {
serialization::SubmoduleID ID;
@ -1983,6 +1989,9 @@ public:
void ReadUnusedLocalTypedefNameCandidates(
llvm::SmallSetVector<const TypedefNameDecl *, 4> &Decls) override;
void ReadDeclsToCheckForDeferredDiags(
llvm::SmallVector<Decl *, 4> &Decls) override;
void ReadReferencedSelectors(
SmallVectorImpl<std::pair<Selector, SourceLocation>> &Sels) override;

6
clang/lib/Sema/MultiplexExternalSemaSource.cpp
View File

@ -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)

179
clang/lib/Sema/Sema.cpp
View File

@ -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);
}
S.DeviceDeferredDiags.erase(It);
{
DiagnosticBuilder Builder(Diags.Report(Loc, PD.getDiagID()));
Builder.setForceEmit();
PD.Emit(Builder);
}
// 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);
// 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;
}
}
}
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)

19
clang/lib/Sema/SemaCUDA.cpp
View File

@ -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)) {

27
clang/lib/Sema/SemaDecl.cpp
View File

@ -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) {

7
clang/lib/Sema/SemaExpr.cpp
View File

@ -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

158
clang/lib/Sema/SemaOpenMP.cpp
View File

@ -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

12
clang/lib/Sema/UsedDeclVisitor.h
View File

@ -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

18
clang/lib/Serialization/ASTReader.cpp
View File

@ -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())

11
clang/lib/Serialization/ASTWriter.cpp
View File

@ -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);

12
clang/test/OpenMP/declare_target_messages.cpp
View File

@ -162,17 +162,17 @@ namespace {
#pragma omp declare target link(x) // expected-error {{'x' must not appear in both clauses 'to' and 'link'}}
void bazz() {}
#pragma omp declare target to(bazz) device_type(nohost) // omp45-error {{unexpected 'device_type' clause, only 'to' or 'link' clauses expected}} host5-note {{marked as 'device_type(nohost)' here}}
#pragma omp declare target to(bazz) device_type(nohost) // omp45-error {{unexpected 'device_type' clause, only 'to' or 'link' clauses expected}} host5-note 3{{marked as 'device_type(nohost)' here}}
void bazzz() {bazz();}
#pragma omp declare target to(bazzz) device_type(nohost) // omp45-error {{unexpected 'device_type' clause, only 'to' or 'link' clauses expected}}
void any() {bazz();} // host5-error {{function with 'device_type(nohost)' is not available on host}}
void host1() {bazz();}
#pragma omp declare target to(host1) device_type(host) // omp45-error {{unexpected 'device_type' clause, only 'to' or 'link' clauses expected}} dev5-note 2 {{marked as 'device_type(host)' here}}
void host2() {bazz();}
void host1() {bazz();} // host5-error {{function with 'device_type(nohost)' is not available on host}}
#pragma omp declare target to(host1) device_type(host) // omp45-error {{unexpected 'device_type' clause, only 'to' or 'link' clauses expected}} dev5-note 4 {{marked as 'device_type(host)' here}}
void host2() {bazz();} //host5-error {{function with 'device_type(nohost)' is not available on host}}
#pragma omp declare target to(host2)
void device() {host1();}
void device() {host1();} // dev5-error {{function with 'device_type(host)' is not available on device}}
#pragma omp declare target to(device) device_type(nohost) // omp45-error {{unexpected 'device_type' clause, only 'to' or 'link' clauses expected}} host5-note 2 {{marked as 'device_type(nohost)' here}}
void host3() {host1();}
void host3() {host1();} // dev5-error {{function with 'device_type(host)' is not available on device}}
#pragma omp declare target to(host3)
#pragma omp declare target

4
clang/test/OpenMP/nvptx_target_exceptions_messages.cpp
View File

@ -38,7 +38,7 @@ int d;
#pragma omp end declare target
int c;
int bar() { return 1 + foo() + bar() + baz1() + baz2(); }
int bar() { return 1 + foo() + bar() + baz1() + baz2(); } // expected-note {{called by 'bar'}}
int maini1() {
int a;
@ -49,7 +49,7 @@ int maini1() {
{
S s(a);
static long aaa = 23;
a = foo() + bar() + b + c + d + aa + aaa + FA<int>();
a = foo() + bar() + b + c + d + aa + aaa + FA<int>(); // expected-note{{called by 'maini1'}}
if (!a)
throw "Error"; // expected-error {{cannot use 'throw' with exceptions disabled}}
}

4
clang/test/SemaCUDA/bad-calls-on-same-line.cu
View File

@ -33,8 +33,8 @@ inline __host__ __device__ void hd() {
void host_fn() {
hd<int>();
hd<double>(); // expected-note {{function template specialization 'hd<double>'}}
hd<double>();
// expected-note@-1 {{called by 'host_fn'}}
hd<float>(); // expected-note {{function template specialization 'hd<float>'}}
hd<float>();
// expected-note@-1 {{called by 'host_fn'}}
}

4
clang/test/SemaCUDA/call-device-fn-from-host.cu
View File

@ -1,7 +1,7 @@
// RUN: %clang_cc1 %s --std=c++11 -triple x86_64-unknown-linux -emit-llvm -o - \
// RUN: -verify -verify-ignore-unexpected=note
// RUN: %clang_cc1 %s --std=c++11 -triple x86_64-unknown-linux -emit-llvm -o - \
// RUN: -verify -verify-ignore-unexpected=note -fopenmp
// RUN: -verify=expected,omp -verify-ignore-unexpected=note -fopenmp
// Note: This test won't work with -fsyntax-only, because some of these errors
// are emitted during codegen.
@ -39,7 +39,7 @@ __host__ __device__ void T::hd3() {
}
template <typename T> __host__ __device__ void hd2() { device_fn(); }
// expected-error@-1 2 {{reference to __device__ function 'device_fn' in __host__ __device__ function}}
// expected-error@-1 {{reference to __device__ function 'device_fn' in __host__ __device__ function}}
void host_fn() { hd2<int>(); }
__host__ __device__ void hd() { device_fn(); }

4
clang/test/SemaCUDA/call-host-fn-from-device.cu
View File

@ -56,14 +56,14 @@ __host__ __device__ void T::hd3() {
}
template <typename T> __host__ __device__ void hd2() { host_fn(); }
// expected-error@-1 2 {{reference to __host__ function 'host_fn' in __host__ __device__ function}}
// expected-error@-1 {{reference to __host__ function 'host_fn' in __host__ __device__ function}}
__global__ void kernel() { hd2<int>(); }
__host__ __device__ void hd() { host_fn(); }
// expected-error@-1 {{reference to __host__ function 'host_fn' in __host__ __device__ function}}
template <typename T> __host__ __device__ void hd3() { host_fn(); }
// expected-error@-1 2 {{reference to __host__ function 'host_fn' in __host__ __device__ function}}
// expected-error@-1 {{reference to __host__ function 'host_fn' in __host__ __device__ function}}
__device__ void device_fn() { hd3<int>(); }
// No error because this is never instantiated.

4
clang/test/SemaCUDA/openmp-target.cu
View File

@ -16,9 +16,9 @@ void bazz() {}
void bazzz() {bazz();}
#pragma omp declare target to(bazzz) device_type(nohost)
void any() {bazz();} // expected-error {{function with 'device_type(nohost)' is not available on host}}
void host1() {bazz();}
void host1() {bazz();} // expected-error {{function with 'device_type(nohost)' is not available on host}}
#pragma omp declare target to(host1) device_type(host)
void host2() {bazz();}
void host2() {bazz();} // expected-error {{function with 'device_type(nohost)' is not available on host}}
#pragma omp declare target to(host2)
void device() {host1();}
#pragma omp declare target to(device) device_type(nohost)

2
clang/test/SemaCUDA/trace-through-global.cu
View File

@ -38,7 +38,7 @@ void launch_kernel() {
// Notice that these two diagnostics are different: Because the call to hd1
// is not dependent on T, the call to hd1 comes from 'launch_kernel', while
// the call to hd3, being dependent, comes from 'launch_kernel<int>'.
hd1(); // expected-note {{called by 'launch_kernel'}}
hd1(); // expected-note {{called by 'launch_kernel<int>'}}
hd3(T()); // expected-note {{called by 'launch_kernel<int>'}}
}