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
This commit is contained in:
parent
e1730cfeb3
commit
2c31aa2de1
|
@ -1502,9 +1502,6 @@ public:
|
|||
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.
|
||||
|
|
|
@ -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
|
||||
// different 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 @@ public:
|
|||
}
|
||||
|
||||
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
|
||||
|
||||
|
|
25
clang/test/CodeGenCUDA/deferred-diag.cu
Normal file
25
clang/test/CodeGenCUDA/deferred-diag.cu
Normal file
|
@ -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();
|
||||
}
|
20
clang/test/SemaCUDA/deferred-diags-limit.cu
Normal file
20
clang/test/SemaCUDA/deferred-diags-limit.cu
Normal file
|
@ -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();
|
||||
}
|
36
clang/test/SemaCUDA/deferred-diags.cu
Normal file
36
clang/test/SemaCUDA/deferred-diags.cu
Normal file
|
@ -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@-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();
|
||||
}
|
Loading…
Reference in a new issue