diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h --- a/clang/include/clang/Sema/Sema.h +++ b/clang/include/clang/Sema/Sema.h @@ -1502,9 +1502,6 @@ 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 --- a/clang/lib/Sema/Sema.cpp +++ b/clang/lib/Sema/Sema.cpp @@ -1440,59 +1440,72 @@ 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 { public: typedef UsedDeclVisitor Inherited; - llvm::SmallSet, 4> Visited; - llvm::SmallVector, 4> UseStack; - bool ShouldEmit; + + // Whether the function is already in the current use-path. + llvm::SmallSet, 4> InUsePath; + + // The current use-path. + llvm::SmallVector, 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, 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 @@ } 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(D)) + if (auto *FD = dyn_cast(D)) { + ShouldEmitRootNode = S.getEmissionStatus(FD, /*Final=*/true) == + Sema::FunctionEmissionStatus::Emitted; checkFunc(SourceLocation(), FD); - else + } else checkVar(cast(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 --- /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 --- /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 --- /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@-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(); +}