Speed up deferred diagnostic emitter

[yaxunl]

Summary

Move function emitDeferredDiags from Sema to DeferredDiagsEmitter since it
is only used by DeferredDiagsEmitter.

Also skip visited functions to avoid exponential compile time.

[rjmccall]

I made a suggestion in the other patch about restructuring things out of visitUsedDecl. I'll review this when that's done.

[yaxunl]

rebase

[rjmccall]

clang/lib/Sema/Sema.cpp

1541

It makes me a little uncomfortable to be holding an iterator this long while calling a fair amount of other stuff in the meantime.

Your use of DiagsCount is subtle enough that it really needs to be explained in some comments. You're doing stuff conditionally based on both whether the entry exists but also whether it's non-zero.

[yaxunl]

added comments

[yaxunl]

clang/lib/Sema/Sema.cpp

1541

added comments

[rjmccall]

clang/lib/Sema/Sema.cpp

1540–1541

This becomes global state for the visitor; that doesn't seem like it can be right.

1541

Okay, thanks for that. Two points, then. First, it looks like the count is really just a boolean for whether the function recursively triggers any diagnostics. And second, can't it just be as simple as whether we've visited that function at all in a context that's forcing diagnostics to be emitted? The logic seems to be to try to emit the diagnostics for each use-path, but why would we want that?

[yaxunl]

clang/lib/Sema/Sema.cpp

1540–1541

This state is for the root node of each traversal initiated from the items in DeclsToCheckForDeferredDiags. It only needs to be set once. Will rename it as ShouldEmitRootNode and move it to checkRecordedDecl.

1541

For the second comment, we need to visit a function again for each use-path because we need to report each use-path that triggers a diagnostic, otherwise users will see a new error after they fix one error, instead of seeing all the errors at once.

For the first comment, I will change the count to two flags: one for the case where the function is not in device context, the other is for the case where the function is in device context. This will allow us to avoid redundant visits whether or not we are in a device context.

[yaxunl]

revised by John's comments

[rjmccall]

clang/lib/Sema/Sema.cpp

1541

For the second comment, we need to visit a function again for each use-path because we need to report each use-path that triggers a diagnostic, otherwise users will see a new error after they fix one error, instead of seeing all the errors at once.

This is not what we do in analogous cases where errors are triggered by a use, like in template instantiation. The bug might be that the device program is using a function that it shouldn't be using, or the bug might be that a function that's supposed to be usable from the device is written incorrectly. In the former case, yes, not reporting the errors for each use-path may force the programmer to build multiple times to find all the problematic uses. However, in the latter case you can easily end up emitting a massive number of errors that completely drowns out everything else. It's also non-linear: the number of different use-paths of a particular function can be combinatoric.

[yaxunl]

clang/lib/Sema/Sema.cpp

1541

The deferred diagnostics fall into the first case mostly.

Not all diagnostic messages happen in device host functions are deferred. Most of diagnostic messages are emitted immediately, disregarding whether the function is emitted or not.

Only a few special types of diagnostic messages e.g. inline assembly errors, exceptions, varags are deferred. This is because clang has to use pragmas to make all functions device and host in some system headers to be able to use them in device compilation, whereas some of these functions will cause error only if they are emitted in device compilation. Since we cannot change these headers, we have to defer such diagnostics to the point where they are actually triggered. Therefore we are mostly interested in "who is calling these functions" instead of the diagnostics themselves.

For normal programs containing no diagnostics, the current approach should sufficiently avoid all redundant visits and all functions will be visited once.

The functions may be visited multiple times when there are deferred diagnostics, however this should be limited by the maximum number of errors.

[rjmccall]

clang/lib/Sema/Sema.cpp

1541

Okay. I understand the point now about being mostly concerned about using functions in headers.

My concern about visiting things a combinatorial number of times is less about generating an excessive number of errors and more about taking an excessive amount of time to finish compilation. The compiler really should not be using any exponential algorithms; a simple visited set will ensure that, but I think what you're doing does not. Can you make a case for why that's not true?

[yaxunl]

clang/lib/Sema/Sema.cpp

1541

The current approach is linear for the number of visits to functions.

Let's assume the number of functions is N.

The current approach already restricted revisiting a node that's already in the current use-path to prevent infinite recursion, therefore a use-path has maximum length N.

The functions can be classified as two categories:

Those who trigger deferred diags directly or indirectly, which we call trigger nodes.

Those who do not trigger deferred diags directly or indirectly, which we call non-trigger nodes.

Non-trigger nodes can be identified at most with two visits. All other visits are skipped. Therefore the total number of visits of non-trigger nodes is less than 2*N.

We can also classify the use-paths as trigger paths and non-trigger paths. Each trigger path causes at least one diagnostic emitted. Each non-trigger path does not cause any diagnostics emitted.

Due to limit of maximum diagnostics allowed, the algorithm will stop after a finite number of use-paths visited. Let's say it is M. (In reality, M could be much smaller than the maximum diagnostic allowed.) Then the number of trigger paths is at most M.

Let's list all the use-paths visited and count all the nodes visited in these use-paths.

First we consider all the use-paths that containing only non-trigger nodes, since each non-trigger nodes can only be visited twice. The total number of visits of nodes is less than 2*N.

Then we consider all the trigger paths. Since the maximum path length is N, the total number of nodes visited is less than N*M.

Now we only need to consider the non-trigger paths that starts with trigger nodes and end with a non-trigger node. For example, ABCD, where A, B, C, D denote nodes (functions) visited in the path. We can prove that A, B, C must be trigger nodes, since if any of A, B, C is a non-trigger node, the path will ends earlier, not reaching D. We can also prove that the sub-path ABC must be shared with a trigger path e.g. ABCEF, since otherwise C will be a non-trigger node. Then we do not need to count visit of A, B, and C, since they have already been counted when we count the visited nodes by trigger paths.

Therefore, the total number of visits for functions is less than (2+M)*N, where M is actual number of diagnostics emitted.

[rjmccall]

clang/lib/Sema/Sema.cpp

1541

Due to limit of maximum diagnostics allowed, the algorithm will stop after a finite number of use-paths visited.

Why do you think this is true? I don't see anything in your actual code which asks whether the error limit has been reached; it just does all the work and trusts the diagnostics engine to stop emitting diagnostics eventually.

Try turning the error limit off (which will tell you how much work you're doing, since you're not short-circuiting when it's reached) and seeing how many diagnostics you get if you have a chain of functions like this:

// All of these functions should be emitted on demand.
void hasInvalid() { /* this function should trigger a deferred diagnostic of some sort */ }
void use0() { hasInvalid(); hasInvalid(); }
void use1() { use0(); use0(); }
void use2() { use1(); use1(); }
void use3() { use2(); use2(); }
void use4() { use3(); use3(); }
//  The final function should be a kernel or something like that.

Also, in general, since the error limit can be disabled, I don't think it's a good idea for it to be the only thing stopping you from doing an exponential amount of work.

[yaxunl]

clang/lib/Sema/Sema.cpp

1541

It seems not practical to emit diagnostic for each triggering use-path since this will cause exponential time and also exponential diagnostics.

I will make change so that each function is visited at most twice.

[yaxunl]

Revised by John's comments. Skip visited functions and check error limit.

[rjmccall]

Thank you. Minor request, but feel free to commit with that change.

clang/lib/Sema/Sema.cpp

1558

insert returns whether it changed the set, so you can combine this check with the insertion unless the exact ordering is important.

[yaxunl]

clang/lib/Sema/Sema.cpp

1558

The order is important here. We cannot inert until the body of the function has been traversed.

[rjmccall]

clang/lib/Sema/Sema.cpp

1558

Why? It only matters if we encounter FD again while recursing into its body, during which time InUsePath.count(FD) will be true and so we'll have bailed out above.

[yaxunl]

clang/lib/Sema/Sema.cpp

1558

You are right. I forgot InUsePath. Will fix this. Thanks.

clang/include/clang/Sema/Sema.h

public:
   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,
     /// A normal translation unit fragment. For a non-module unit, this is the
     /// entire translation unit. Otherwise, it runs from the module-declaration
     /// to the private-module-fragment (if any) or the end of the TU (if not).
     Normal,

clang/lib/Sema/Sema.cpp

Sema::Diag(SourceLocation Loc, const PartialDiagnostic& PD) {
   return Builder;
 }
 
 // Print notes showing how we can reach FD starting from an a priori
 // known-callable function.
 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;
     Inherited::VisitOMPTargetDirective(Node);
     --InOMPDeviceContext;
   }
 
   void visitUsedDecl(SourceLocation Loc, Decl *D) {
if (auto *Init = VD->getInit()) {
       if (IsDev)
         ++InOMPDeviceContext;
       this->Visit(Init);
       if (IsDev)
         --InOMPDeviceContext;
     }
   }
 
   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
 
 void Sema::emitDeferredDiags() {
   if (ExternalSource)
     ExternalSource->ReadDeclsToCheckForDeferredDiags(
         DeclsToCheckForDeferredDiags);
 

clang/test/CodeGenCUDA/deferred-diag.cu

+// 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();
+}

clang/test/SemaCUDA/deferred-diags-limit.cu

+// 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();
+}

clang/test/SemaCUDA/deferred-diags.cu

+// 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();
+}