Skip to content

Commit 1b978dd

Browse files
yxsamliuyxsamliu
committed
[CUDA][HIP][OpenMP] Emit deferred diagnostics by a post-parsing AST travese
This patch removes the explicit call graph for CUDA/HIP/OpenMP deferred diagnostics generated during parsing since it is error prone due to incomplete information about function declarations during parsing. In stead, this patch does a post-parsing AST traverse and emits deferred diagnostics based on the use graph implicitly generated during the traverse. Differential Revision: https://reviews.llvm.org/D70172
1 parent e8e078c commit 1b978dd

13 files changed

+201
-349
lines changed

‎clang/include/clang/Sema/Sema.h

Lines changed: 13 additions & 33 deletions
Original file line numberDiff line numberDiff line change
@@ -1464,6 +1464,12 @@ class Sema final {
14641464

14651465
void emitAndClearUnusedLocalTypedefWarnings();
14661466

1467+
// Emit all deferred diagnostics.
1468+
void emitDeferredDiags();
1469+
// Emit any deferred diagnostics for FD and erase them from the map in which
1470+
// they're stored.
1471+
void emitDeferredDiags(FunctionDecl *FD, bool ShowCallStack);
1472+
14671473
enum TUFragmentKind {
14681474
/// The global module fragment, between 'module;' and a module-declaration.
14691475
Global,
@@ -3683,7 +3689,8 @@ class Sema final {
36833689
TemplateDiscarded, // Discarded due to uninstantiated templates
36843690
Unknown,
36853691
};
3686-
FunctionEmissionStatus getEmissionStatus(FunctionDecl *Decl);
3692+
FunctionEmissionStatus getEmissionStatus(FunctionDecl *Decl,
3693+
bool Final = false);
36873694

36883695
// Whether the callee should be ignored in CUDA/HIP/OpenMP host/device check.
36893696
bool shouldIgnoreInHostDeviceCheck(FunctionDecl *Callee);
@@ -9677,22 +9684,10 @@ class Sema final {
96779684
/// Pop OpenMP function region for non-capturing function.
96789685
void popOpenMPFunctionRegion(const sema::FunctionScopeInfo *OldFSI);
96799686

9680-
/// Check whether we're allowed to call Callee from the current function.
9681-
void checkOpenMPDeviceFunction(SourceLocation Loc, FunctionDecl *Callee,
9682-
bool CheckForDelayedContext = true);
9683-
9684-
/// Check whether we're allowed to call Callee from the current function.
9685-
void checkOpenMPHostFunction(SourceLocation Loc, FunctionDecl *Callee,
9686-
bool CheckCaller = true);
9687-
96889687
/// Check if the expression is allowed to be used in expressions for the
96899688
/// OpenMP devices.
96909689
void checkOpenMPDeviceExpr(const Expr *E);
96919690

9692-
/// Finishes analysis of the deferred functions calls that may be declared as
9693-
/// host/nohost during device/host compilation.
9694-
void finalizeOpenMPDelayedAnalysis();
9695-
96969691
/// Checks if a type or a declaration is disabled due to the owning extension
96979692
/// being disabled, and emits diagnostic messages if it is disabled.
96989693
/// \param D type or declaration to be checked.
@@ -9875,6 +9870,11 @@ class Sema final {
98759870
void
98769871
checkDeclIsAllowedInOpenMPTarget(Expr *E, Decl *D,
98779872
SourceLocation IdLoc = SourceLocation());
9873+
/// Finishes analysis of the deferred functions calls that may be declared as
9874+
/// host/nohost during device/host compilation.
9875+
void finalizeOpenMPDelayedAnalysis(const FunctionDecl *Caller,
9876+
const FunctionDecl *Callee,
9877+
SourceLocation Loc);
98789878
/// Return true inside OpenMP declare target region.
98799879
bool isInOpenMPDeclareTargetContext() const {
98809880
return DeclareTargetNestingLevel > 0;
@@ -11223,18 +11223,6 @@ class Sema final {
1122311223
/* Caller = */ FunctionDeclAndLoc>
1122411224
DeviceKnownEmittedFns;
1122511225

11226-
/// A partial call graph maintained during CUDA/OpenMP device code compilation
11227-
/// to support deferred diagnostics.
11228-
///
11229-
/// Functions are only added here if, at the time they're considered, they are
11230-
/// not known-emitted. As soon as we discover that a function is
11231-
/// known-emitted, we remove it and everything it transitively calls from this
11232-
/// set and add those functions to DeviceKnownEmittedFns.
11233-
llvm::DenseMap</* Caller = */ CanonicalDeclPtr<FunctionDecl>,
11234-
/* Callees = */ llvm::MapVector<CanonicalDeclPtr<FunctionDecl>,
11235-
SourceLocation>>
11236-
DeviceCallGraph;
11237-
1123811226
/// Diagnostic builder for CUDA/OpenMP devices errors which may or may not be
1123911227
/// deferred.
1124011228
///
@@ -11309,14 +11297,6 @@ class Sema final {
1130911297
llvm::Optional<unsigned> PartialDiagId;
1131011298
};
1131111299

11312-
/// Indicate that this function (and thus everything it transtively calls)
11313-
/// will be codegen'ed, and emit any deferred diagnostics on this function and
11314-
/// its (transitive) callees.
11315-
void markKnownEmitted(
11316-
Sema &S, FunctionDecl *OrigCaller, FunctionDecl *OrigCallee,
11317-
SourceLocation OrigLoc,
11318-
const llvm::function_ref<bool(Sema &, FunctionDecl *)> IsKnownEmitted);
11319-
1132011300
/// Creates a DeviceDiagBuilder that emits the diagnostic if the current context
1132111301
/// is "used as device code".
1132211302
///

‎clang/lib/Sema/Sema.cpp

Lines changed: 110 additions & 75 deletions
Original file line numberDiff line numberDiff line change
@@ -11,6 +11,7 @@
1111
//
1212
//===----------------------------------------------------------------------===//
1313

14+
#include "UsedDeclVisitor.h"
1415
#include "clang/AST/ASTContext.h"
1516
#include "clang/AST/ASTDiagnostic.h"
1617
#include "clang/AST/DeclCXX.h"
@@ -954,9 +955,7 @@ void Sema::ActOnEndOfTranslationUnitFragment(TUFragmentKind Kind) {
954955
PerformPendingInstantiations();
955956
}
956957

957-
// Finalize analysis of OpenMP-specific constructs.
958-
if (LangOpts.OpenMP)
959-
finalizeOpenMPDelayedAnalysis();
958+
emitDeferredDiags();
960959

961960
assert(LateParsedInstantiations.empty() &&
962961
"end of TU template instantiation should not create more "
@@ -1451,27 +1450,128 @@ static void emitCallStackNotes(Sema &S, FunctionDecl *FD) {
14511450

14521451
// Emit any deferred diagnostics for FD and erase them from the map in which
14531452
// they're stored.
1454-
static void emitDeferredDiags(Sema &S, FunctionDecl *FD, bool ShowCallStack) {
1455-
auto It = S.DeviceDeferredDiags.find(FD);
1456-
if (It == S.DeviceDeferredDiags.end())
1453+
void Sema::emitDeferredDiags(FunctionDecl *FD, bool ShowCallStack) {
1454+
auto It = DeviceDeferredDiags.find(FD);
1455+
if (It == DeviceDeferredDiags.end())
14571456
return;
14581457
bool HasWarningOrError = false;
14591458
for (PartialDiagnosticAt &PDAt : It->second) {
14601459
const SourceLocation &Loc = PDAt.first;
14611460
const PartialDiagnostic &PD = PDAt.second;
1462-
HasWarningOrError |= S.getDiagnostics().getDiagnosticLevel(
1461+
HasWarningOrError |= getDiagnostics().getDiagnosticLevel(
14631462
PD.getDiagID(), Loc) >= DiagnosticsEngine::Warning;
1464-
DiagnosticBuilder Builder(S.Diags.Report(Loc, PD.getDiagID()));
1463+
DiagnosticBuilder Builder(Diags.Report(Loc, PD.getDiagID()));
14651464
Builder.setForceEmit();
14661465
PD.Emit(Builder);
14671466
}
1468-
S.DeviceDeferredDiags.erase(It);
14691467

14701468
// FIXME: Should this be called after every warning/error emitted in the loop
14711469
// above, instead of just once per function? That would be consistent with
14721470
// how we handle immediate errors, but it also seems like a bit much.
14731471
if (HasWarningOrError && ShowCallStack)
1474-
emitCallStackNotes(S, FD);
1472+
emitCallStackNotes(*this, FD);
1473+
}
1474+
1475+
namespace {
1476+
/// Helper class that emits deferred diagnostic messages if an entity directly
1477+
/// or indirectly using the function that causes the deferred diagnostic
1478+
/// messages is known to be emitted.
1479+
class DeferredDiagnosticsEmitter
1480+
: public UsedDeclVisitor<DeferredDiagnosticsEmitter> {
1481+
public:
1482+
typedef UsedDeclVisitor<DeferredDiagnosticsEmitter> Inherited;
1483+
llvm::SmallSet<CanonicalDeclPtr<Decl>, 4> Visited;
1484+
llvm::SmallVector<CanonicalDeclPtr<FunctionDecl>, 4> UseStack;
1485+
bool ShouldEmit;
1486+
unsigned InOMPDeviceContext;
1487+
1488+
DeferredDiagnosticsEmitter(Sema &S)
1489+
: Inherited(S), ShouldEmit(false), InOMPDeviceContext(0) {}
1490+
1491+
void VisitDeclRefExpr(DeclRefExpr *E) {
1492+
if (FunctionDecl *FD = dyn_cast<FunctionDecl>(E->getDecl())) {
1493+
visitUsedDecl(E->getLocation(), FD);
1494+
}
1495+
}
1496+
1497+
void VisitMemberExpr(MemberExpr *E) {
1498+
if (FunctionDecl *FD = dyn_cast<FunctionDecl>(E->getMemberDecl()))
1499+
visitUsedDecl(E->getMemberLoc(), FD);
1500+
}
1501+
1502+
void VisitOMPTargetDirective(OMPTargetDirective *Node) {
1503+
++InOMPDeviceContext;
1504+
Inherited::VisitOMPTargetDirective(Node);
1505+
--InOMPDeviceContext;
1506+
}
1507+
1508+
void VisitCapturedStmt(CapturedStmt *Node) {
1509+
visitUsedDecl(Node->getBeginLoc(), Node->getCapturedDecl());
1510+
Inherited::VisitCapturedStmt(Node);
1511+
}
1512+
1513+
void visitUsedDecl(SourceLocation Loc, Decl *D) {
1514+
if (auto *TD = dyn_cast<TranslationUnitDecl>(D)) {
1515+
for (auto *DD : TD->decls()) {
1516+
visitUsedDecl(Loc, DD);
1517+
}
1518+
} else if (auto *FTD = dyn_cast<FunctionTemplateDecl>(D)) {
1519+
for (auto *DD : FTD->specializations()) {
1520+
visitUsedDecl(Loc, DD);
1521+
}
1522+
} else if (auto *FD = dyn_cast<FunctionDecl>(D)) {
1523+
FunctionDecl *Caller = UseStack.empty() ? nullptr : UseStack.back();
1524+
auto IsKnownEmitted = S.getEmissionStatus(FD, /*Final=*/true) ==
1525+
Sema::FunctionEmissionStatus::Emitted;
1526+
if (!Caller)
1527+
ShouldEmit = IsKnownEmitted;
1528+
if ((!ShouldEmit && !S.getLangOpts().OpenMP && !Caller) ||
1529+
S.shouldIgnoreInHostDeviceCheck(FD) || Visited.count(D))
1530+
return;
1531+
// Finalize analysis of OpenMP-specific constructs.
1532+
if (Caller && S.LangOpts.OpenMP && UseStack.size() == 1)
1533+
S.finalizeOpenMPDelayedAnalysis(Caller, FD, Loc);
1534+
if (Caller)
1535+
S.DeviceKnownEmittedFns[FD] = {Caller, Loc};
1536+
if (ShouldEmit || InOMPDeviceContext)
1537+
S.emitDeferredDiags(FD, Caller);
1538+
Visited.insert(D);
1539+
UseStack.push_back(FD);
1540+
if (auto *S = FD->getBody()) {
1541+
this->Visit(S);
1542+
}
1543+
UseStack.pop_back();
1544+
Visited.erase(D);
1545+
} else if (auto *RD = dyn_cast<RecordDecl>(D)) {
1546+
for (auto *DD : RD->decls()) {
1547+
visitUsedDecl(Loc, DD);
1548+
}
1549+
} else if (auto *CD = dyn_cast<CapturedDecl>(D)) {
1550+
if (auto *S = CD->getBody()) {
1551+
this->Visit(S);
1552+
}
1553+
} else if (auto *VD = dyn_cast<VarDecl>(D)) {
1554+
if (auto *Init = VD->getInit()) {
1555+
auto DevTy = OMPDeclareTargetDeclAttr::getDeviceType(VD);
1556+
bool IsDev = DevTy && (*DevTy == OMPDeclareTargetDeclAttr::DT_NoHost ||
1557+
*DevTy == OMPDeclareTargetDeclAttr::DT_Any);
1558+
if (IsDev)
1559+
++InOMPDeviceContext;
1560+
this->Visit(Init);
1561+
if (IsDev)
1562+
--InOMPDeviceContext;
1563+
}
1564+
}
1565+
}
1566+
};
1567+
} // namespace
1568+
1569+
void Sema::emitDeferredDiags() {
1570+
if (DeviceDeferredDiags.empty() && !LangOpts.OpenMP)
1571+
return;
1572+
1573+
DeferredDiagnosticsEmitter(*this).visitUsedDecl(
1574+
SourceLocation(), Context.getTranslationUnitDecl());
14751575
}
14761576

14771577
// In CUDA, there are some constructs which may appear in semantically-valid
@@ -1544,71 +1644,6 @@ Sema::DeviceDiagBuilder::~DeviceDiagBuilder() {
15441644
}
15451645
}
15461646

1547-
// Indicate that this function (and thus everything it transtively calls) will
1548-
// be codegen'ed, and emit any deferred diagnostics on this function and its
1549-
// (transitive) callees.
1550-
void Sema::markKnownEmitted(
1551-
Sema &S, FunctionDecl *OrigCaller, FunctionDecl *OrigCallee,
1552-
SourceLocation OrigLoc,
1553-
const llvm::function_ref<bool(Sema &, FunctionDecl *)> IsKnownEmitted) {
1554-
// Nothing to do if we already know that FD is emitted.
1555-
if (IsKnownEmitted(S, OrigCallee)) {
1556-
assert(!S.DeviceCallGraph.count(OrigCallee));
1557-
return;
1558-
}
1559-
1560-
// We've just discovered that OrigCallee is known-emitted. Walk our call
1561-
// graph to see what else we can now discover also must be emitted.
1562-
1563-
struct CallInfo {
1564-
FunctionDecl *Caller;
1565-
FunctionDecl *Callee;
1566-
SourceLocation Loc;
1567-
};
1568-
llvm::SmallVector<CallInfo, 4> Worklist = {{OrigCaller, OrigCallee, OrigLoc}};
1569-
llvm::SmallSet<CanonicalDeclPtr<FunctionDecl>, 4> Seen;
1570-
Seen.insert(OrigCallee);
1571-
while (!Worklist.empty()) {
1572-
CallInfo C = Worklist.pop_back_val();
1573-
assert(!IsKnownEmitted(S, C.Callee) &&
1574-
"Worklist should not contain known-emitted functions.");
1575-
S.DeviceKnownEmittedFns[C.Callee] = {C.Caller, C.Loc};
1576-
emitDeferredDiags(S, C.Callee, C.Caller);
1577-
1578-
// If this is a template instantiation, explore its callgraph as well:
1579-
// Non-dependent calls are part of the template's callgraph, while dependent
1580-
// calls are part of to the instantiation's call graph.
1581-
if (auto *Templ = C.Callee->getPrimaryTemplate()) {
1582-
FunctionDecl *TemplFD = Templ->getAsFunction();
1583-
if (!Seen.count(TemplFD) && !S.DeviceKnownEmittedFns.count(TemplFD)) {
1584-
Seen.insert(TemplFD);
1585-
Worklist.push_back(
1586-
{/* Caller = */ C.Caller, /* Callee = */ TemplFD, C.Loc});
1587-
}
1588-
}
1589-
1590-
// Add all functions called by Callee to our worklist.
1591-
auto CGIt = S.DeviceCallGraph.find(C.Callee);
1592-
if (CGIt == S.DeviceCallGraph.end())
1593-
continue;
1594-
1595-
for (std::pair<CanonicalDeclPtr<FunctionDecl>, SourceLocation> FDLoc :
1596-
CGIt->second) {
1597-
FunctionDecl *NewCallee = FDLoc.first;
1598-
SourceLocation CallLoc = FDLoc.second;
1599-
if (Seen.count(NewCallee) || IsKnownEmitted(S, NewCallee))
1600-
continue;
1601-
Seen.insert(NewCallee);
1602-
Worklist.push_back(
1603-
{/* Caller = */ C.Callee, /* Callee = */ NewCallee, CallLoc});
1604-
}
1605-
1606-
// C.Callee is now known-emitted, so we no longer need to maintain its list
1607-
// of callees in DeviceCallGraph.
1608-
S.DeviceCallGraph.erase(CGIt);
1609-
}
1610-
}
1611-
16121647
Sema::DeviceDiagBuilder Sema::targetDiag(SourceLocation Loc, unsigned DiagID) {
16131648
if (LangOpts.OpenMP)
16141649
return LangOpts.OpenMPIsDevice ? diagIfOpenMPDeviceCode(Loc, DiagID)

‎clang/lib/Sema/SemaCUDA.cpp

Lines changed: 0 additions & 19 deletions
Original file line numberDiff line numberDiff line change
@@ -674,25 +674,6 @@ bool Sema::CheckCUDACall(SourceLocation Loc, FunctionDecl *Callee) {
674674
// Otherwise, mark the call in our call graph so we can traverse it later.
675675
bool CallerKnownEmitted =
676676
getEmissionStatus(Caller) == FunctionEmissionStatus::Emitted;
677-
if (CallerKnownEmitted) {
678-
// Host-side references to a __global__ function refer to the stub, so the
679-
// function itself is never emitted and therefore should not be marked.
680-
if (!shouldIgnoreInHostDeviceCheck(Callee))
681-
markKnownEmitted(
682-
*this, Caller, Callee, Loc, [](Sema &S, FunctionDecl *FD) {
683-
return S.getEmissionStatus(FD) == FunctionEmissionStatus::Emitted;
684-
});
685-
} else {
686-
// If we have
687-
// host fn calls kernel fn calls host+device,
688-
// the HD function does not get instantiated on the host. We model this by
689-
// omitting at the call to the kernel from the callgraph. This ensures
690-
// that, when compiling for host, only HD functions actually called from the
691-
// host get marked as known-emitted.
692-
if (!shouldIgnoreInHostDeviceCheck(Callee))
693-
DeviceCallGraph[Caller].insert({Callee, Loc});
694-
}
695-
696677
DeviceDiagBuilder::Kind DiagKind = [this, Caller, Callee,
697678
CallerKnownEmitted] {
698679
switch (IdentifyCUDAPreference(Caller, Callee)) {

‎clang/lib/Sema/SemaDecl.cpp

Lines changed: 10 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -17929,7 +17929,8 @@ Decl *Sema::getObjCDeclContext() const {
1792917929
return (dyn_cast_or_null<ObjCContainerDecl>(CurContext));
1793017930
}
1793117931

17932-
Sema::FunctionEmissionStatus Sema::getEmissionStatus(FunctionDecl *FD) {
17932+
Sema::FunctionEmissionStatus Sema::getEmissionStatus(FunctionDecl *FD,
17933+
bool Final) {
1793317934
// Templates are emitted when they're instantiated.
1793417935
if (FD->isDependentContext())
1793517936
return FunctionEmissionStatus::TemplateDiscarded;
@@ -17941,8 +17942,10 @@ Sema::FunctionEmissionStatus Sema::getEmissionStatus(FunctionDecl *FD) {
1794117942
if (DevTy.hasValue()) {
1794217943
if (*DevTy == OMPDeclareTargetDeclAttr::DT_Host)
1794317944
OMPES = FunctionEmissionStatus::OMPDiscarded;
17944-
else if (DeviceKnownEmittedFns.count(FD) > 0)
17945+
else if (*DevTy == OMPDeclareTargetDeclAttr::DT_NoHost ||
17946+
*DevTy == OMPDeclareTargetDeclAttr::DT_Any) {
1794517947
OMPES = FunctionEmissionStatus::Emitted;
17948+
}
1794617949
}
1794717950
} else if (LangOpts.OpenMP) {
1794817951
// In OpenMP 4.5 all the functions are host functions.
@@ -17958,10 +17961,11 @@ Sema::FunctionEmissionStatus Sema::getEmissionStatus(FunctionDecl *FD) {
1795817961
if (DevTy.hasValue()) {
1795917962
if (*DevTy == OMPDeclareTargetDeclAttr::DT_NoHost) {
1796017963
OMPES = FunctionEmissionStatus::OMPDiscarded;
17961-
} else if (DeviceKnownEmittedFns.count(FD) > 0) {
17964+
} else if (*DevTy == OMPDeclareTargetDeclAttr::DT_Host ||
17965+
*DevTy == OMPDeclareTargetDeclAttr::DT_Any)
1796217966
OMPES = FunctionEmissionStatus::Emitted;
17963-
}
17964-
}
17967+
} else if (Final)
17968+
OMPES = FunctionEmissionStatus::Emitted;
1796517969
}
1796617970
}
1796717971
if (OMPES == FunctionEmissionStatus::OMPDiscarded ||
@@ -17996,9 +18000,7 @@ Sema::FunctionEmissionStatus Sema::getEmissionStatus(FunctionDecl *FD) {
1799618000

1799718001
// Otherwise, the function is known-emitted if it's in our set of
1799818002
// known-emitted functions.
17999-
return (DeviceKnownEmittedFns.count(FD) > 0)
18000-
? FunctionEmissionStatus::Emitted
18001-
: FunctionEmissionStatus::Unknown;
18003+
return FunctionEmissionStatus::Unknown;
1800218004
}
1800318005

1800418006
bool Sema::shouldIgnoreInHostDeviceCheck(FunctionDecl *Callee) {

0 commit comments

Comments
 (0)