[clang] 2c31aa2 - Speed up deferred diagnostic emitter
Yaxun Liu via cfe-commits
cfe-commits at lists.llvm.org
Mon Apr 6 10:08:31 PDT 2020
Author: Yaxun (Sam) Liu
Date: 2020-04-06T13:07:43-04:00
New Revision: 2c31aa2de13a23a00ced87123b92e905f2929c7b
URL: https://github.com/llvm/llvm-project/commit/2c31aa2de13a23a00ced87123b92e905f2929c7b
DIFF: https://github.com/llvm/llvm-project/commit/2c31aa2de13a23a00ced87123b92e905f2929c7b.diff
LOG: Speed up deferred diagnostic emitter
Move function emitDeferredDiags from Sema to DeferredDiagsEmitter since it
is only used by DeferredDiagsEmitter.
Also skip visited functions to avoid exponential compile time.
Differential Revision: https://reviews.llvm.org/D77028
Added:
clang/test/CodeGenCUDA/deferred-diag.cu
clang/test/SemaCUDA/deferred-diags-limit.cu
clang/test/SemaCUDA/deferred-diags.cu
Modified:
clang/include/clang/Sema/Sema.h
clang/lib/Sema/Sema.cpp
Removed:
################################################################################
diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h
index 10e2d69f3d9e..4645ef8d2cd1 100644
--- a/clang/include/clang/Sema/Sema.h
+++ b/clang/include/clang/Sema/Sema.h
@@ -1502,9 +1502,6 @@ class Sema final {
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.
diff --git a/clang/lib/Sema/Sema.cpp b/clang/lib/Sema/Sema.cpp
index 44b5115d9d2f..ff184aa7f48d 100644
--- a/clang/lib/Sema/Sema.cpp
+++ b/clang/lib/Sema/Sema.cpp
@@ -1440,59 +1440,72 @@ Sema::Diag(SourceLocation Loc, const PartialDiagnostic& PD) {
static void emitCallStackNotes(Sema &S, FunctionDecl *FD) {
auto FnIt = S.DeviceKnownEmittedFns.find(FD);
while (FnIt != S.DeviceKnownEmittedFns.end()) {
+ // Respect error limit.
+ if (S.Diags.hasFatalErrorOccurred())
+ return;
DiagnosticBuilder Builder(
S.Diags.Report(FnIt->second.Loc, diag::note_called_by));
Builder << FnIt->second.FD;
- Builder.setForceEmit();
-
FnIt = S.DeviceKnownEmittedFns.find(FnIt->second.FD);
}
}
-// Emit any deferred diagnostics for FD and erase them from the map in which
-// they're stored.
-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 |= getDiagnostics().getDiagnosticLevel(
- PD.getDiagID(), Loc) >= DiagnosticsEngine::Warning;
- {
- 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;
- }
- }
-
-}
-
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.
+///
+/// During parsing of AST, certain diagnostic messages are recorded as deferred
+/// diagnostics since it is unknown whether the functions containing such
+/// diagnostics will be emitted. A list of potentially emitted functions and
+/// variables that may potentially trigger emission of functions are also
+/// recorded. DeferredDiagnosticsEmitter recursively visits used functions
+/// by each function to emit deferred diagnostics.
+///
+/// During the visit, certain OpenMP directives or initializer of variables
+/// with certain OpenMP attributes will cause subsequent visiting of any
+/// functions enter a state which is called OpenMP device context in this
+/// implementation. The state is exited when the directive or initializer is
+/// exited. This state can change the emission states of subsequent uses
+/// of functions.
+///
+/// Conceptually the functions or variables to be visited form a use graph
+/// where the parent node uses the child node. At any point of the visit,
+/// the tree nodes traversed from the tree root to the current node form a use
+/// stack. The emission state of the current node depends on two factors:
+/// 1. the emission state of the root node
+/// 2. whether the current node is in OpenMP device context
+/// If the function is decided to be emitted, its contained deferred diagnostics
+/// are emitted, together with the information about the use stack.
+///
class DeferredDiagnosticsEmitter
: public UsedDeclVisitor<DeferredDiagnosticsEmitter> {
public:
typedef UsedDeclVisitor<DeferredDiagnosticsEmitter> Inherited;
- llvm::SmallSet<CanonicalDeclPtr<Decl>, 4> Visited;
- llvm::SmallVector<CanonicalDeclPtr<FunctionDecl>, 4> UseStack;
- bool ShouldEmit;
+
+ // Whether the function is already in the current use-path.
+ llvm::SmallSet<CanonicalDeclPtr<Decl>, 4> InUsePath;
+
+ // The current use-path.
+ llvm::SmallVector<CanonicalDeclPtr<FunctionDecl>, 4> UsePath;
+
+ // Whether the visiting of the function has been done. Done[0] is for the
+ // case not in OpenMP device context. Done[1] is for the case in OpenMP
+ // device context. We need two sets because diagnostics emission may be
+ //
diff erent depending on whether it is in OpenMP device context.
+ llvm::SmallSet<CanonicalDeclPtr<Decl>, 4> DoneMap[2];
+
+ // Emission state of the root node of the current use graph.
+ bool ShouldEmitRootNode;
+
+ // Current OpenMP device context level. It is initialized to 0 and each
+ // entering of device context increases it by 1 and each exit decreases
+ // it by 1. Non-zero value indicates it is currently in device context.
unsigned InOMPDeviceContext;
DeferredDiagnosticsEmitter(Sema &S)
- : Inherited(S), ShouldEmit(false), InOMPDeviceContext(0) {}
+ : Inherited(S), ShouldEmitRootNode(false), InOMPDeviceContext(0) {}
void VisitOMPTargetDirective(OMPTargetDirective *Node) {
++InOMPDeviceContext;
@@ -1525,36 +1538,72 @@ class DeferredDiagnosticsEmitter
}
void checkFunc(SourceLocation Loc, FunctionDecl *FD) {
- 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(FD))
+ auto &Done = DoneMap[InOMPDeviceContext];
+ FunctionDecl *Caller = UsePath.empty() ? nullptr : UsePath.back();
+ if ((!ShouldEmitRootNode && !S.getLangOpts().OpenMP && !Caller) ||
+ S.shouldIgnoreInHostDeviceCheck(FD) || InUsePath.count(FD))
return;
// Finalize analysis of OpenMP-specific constructs.
- if (Caller && S.LangOpts.OpenMP && UseStack.size() == 1)
+ if (Caller && S.LangOpts.OpenMP && UsePath.size() == 1)
S.finalizeOpenMPDelayedAnalysis(Caller, FD, Loc);
if (Caller)
S.DeviceKnownEmittedFns[FD] = {Caller, Loc};
- if (ShouldEmit || InOMPDeviceContext)
- S.emitDeferredDiags(FD, Caller);
- Visited.insert(FD);
- UseStack.push_back(FD);
+ // Always emit deferred diagnostics for the direct users. This does not
+ // lead to explosion of diagnostics since each user is visited at most
+ // twice.
+ if (ShouldEmitRootNode || InOMPDeviceContext)
+ emitDeferredDiags(FD, Caller);
+ // Do not revisit a function if the function body has been completely
+ // visited before.
+ if (Done.count(FD))
+ return;
+ InUsePath.insert(FD);
+ UsePath.push_back(FD);
if (auto *S = FD->getBody()) {
this->Visit(S);
}
- UseStack.pop_back();
- Visited.erase(FD);
+ UsePath.pop_back();
+ InUsePath.erase(FD);
+ Done.insert(FD);
}
void checkRecordedDecl(Decl *D) {
- if (auto *FD = dyn_cast<FunctionDecl>(D))
+ if (auto *FD = dyn_cast<FunctionDecl>(D)) {
+ ShouldEmitRootNode = S.getEmissionStatus(FD, /*Final=*/true) ==
+ Sema::FunctionEmissionStatus::Emitted;
checkFunc(SourceLocation(), FD);
- else
+ } else
checkVar(cast<VarDecl>(D));
}
+
+ // Emit any deferred diagnostics for FD
+ void emitDeferredDiags(FunctionDecl *FD, bool ShowCallStack) {
+ auto It = S.DeviceDeferredDiags.find(FD);
+ if (It == S.DeviceDeferredDiags.end())
+ return;
+ bool HasWarningOrError = false;
+ bool FirstDiag = true;
+ for (PartialDiagnosticAt &PDAt : It->second) {
+ // Respect error limit.
+ if (S.Diags.hasFatalErrorOccurred())
+ return;
+ const SourceLocation &Loc = PDAt.first;
+ const PartialDiagnostic &PD = PDAt.second;
+ HasWarningOrError |=
+ S.getDiagnostics().getDiagnosticLevel(PD.getDiagID(), Loc) >=
+ DiagnosticsEngine::Warning;
+ {
+ DiagnosticBuilder Builder(S.Diags.Report(Loc, PD.getDiagID()));
+ 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(S, FD);
+ FirstDiag = false;
+ }
+ }
+ }
};
} // namespace
diff --git a/clang/test/CodeGenCUDA/deferred-diag.cu b/clang/test/CodeGenCUDA/deferred-diag.cu
new file mode 100644
index 000000000000..e031c0c5aade
--- /dev/null
+++ b/clang/test/CodeGenCUDA/deferred-diag.cu
@@ -0,0 +1,25 @@
+// RUN: not %clang_cc1 -std=c++11 -triple x86_64-unknown-linux-gnu \
+// RUN: -emit-llvm -o - %s 2>&1 | FileCheck %s
+// RUN: not %clang_cc1 -std=c++11 -triple x86_64-unknown-linux-gnu \
+// RUN: -fcuda-is-device -emit-llvm -o - %s 2>&1 \
+// RUN: | FileCheck %s
+
+// Check no crash due to deferred diagnostics.
+
+#include "Inputs/cuda.h"
+
+// CHECK: error: invalid output constraint '=h' in asm
+// CHECK-NOT: core dump
+inline __host__ __device__ int foo() {
+ short h;
+ __asm__("dont care" : "=h"(h) : "f"(0.0), "d"(0.0), "h"(0), "r"(0), "l"(0));
+ return 0;
+}
+
+void host_fun() {
+ foo();
+}
+
+__global__ void kernel() {
+ foo();
+}
diff --git a/clang/test/SemaCUDA/deferred-diags-limit.cu b/clang/test/SemaCUDA/deferred-diags-limit.cu
new file mode 100644
index 000000000000..59328134da90
--- /dev/null
+++ b/clang/test/SemaCUDA/deferred-diags-limit.cu
@@ -0,0 +1,20 @@
+// RUN: not %clang_cc1 -fcxx-exceptions -fcuda-is-device -fsyntax-only \
+// RUN: -ferror-limit 2 2>&1 %s | FileCheck %s
+
+#include "Inputs/cuda.h"
+
+// CHECK: cannot use 'throw' in __host__ __device__ function
+// CHECK: cannot use 'throw' in __host__ __device__ function
+// CHECK-NOT: cannot use 'throw' in __host__ __device__ function
+// CHECK: too many errors emitted, stopping now
+
+inline __host__ __device__ void hasInvalid() {
+ throw NULL;
+}
+
+__global__ void use0() {
+ hasInvalid();
+ hasInvalid();
+ hasInvalid();
+ hasInvalid();
+}
diff --git a/clang/test/SemaCUDA/deferred-diags.cu b/clang/test/SemaCUDA/deferred-diags.cu
new file mode 100644
index 000000000000..856a5e06a58f
--- /dev/null
+++ b/clang/test/SemaCUDA/deferred-diags.cu
@@ -0,0 +1,36 @@
+// RUN: %clang_cc1 -fcxx-exceptions -fcuda-is-device -fsyntax-only -verify %s
+
+#include "Inputs/cuda.h"
+
+// Error, instantiated on device.
+inline __host__ __device__ void hasInvalid() {
+ throw NULL;
+ // expected-error at -1 2{{cannot use 'throw' in __host__ __device__ function}}
+}
+
+static __device__ void use0() {
+ hasInvalid(); // expected-note {{called by 'use0'}}
+ hasInvalid(); // expected-note {{called by 'use0'}}
+}
+
+// To avoid excessive diagnostic messages, deferred diagnostics are only
+// emitted the first time a function is called.
+static __device__ void use1() {
+ use0(); // expected-note 2{{called by 'use1'}}
+ use0();
+}
+
+static __device__ void use2() {
+ use1(); // expected-note 2{{called by 'use2'}}
+ use1();
+}
+
+static __device__ void use3() {
+ use2(); // expected-note 2{{called by 'use3'}}
+ use2();
+}
+
+__global__ void use4() {
+ use3(); // expected-note 2{{called by 'use4'}}
+ use3();
+}
More information about the cfe-commits
mailing list