Fix infinite recursion in deferred diagnostic emitter

[yaxunl]

Summary

Currently deferred diagnostic emitter checks variable decl in DeclRefExpr, which
causes infinite recursion for cases like long a = (long)&a;.

Deferred diagnostic emitter does not need check variable decls in DeclRefExpr
since reference of a variable does not cause emission of functions directly or
indirectly. Therefore there is no need to check variable decls in DeclRefExpr.

[rjmccall]

Can you explain what exactly the emission/semantic model is for variables? Normal code-generation absolutely triggers the emission of many variables lazily (e.g. internal-linkage globals, C++ inline variables); and any variable that's *not* being emitted lazily actually needs to be treated as a potential root into the delayed-diagnostic graph.

[yaxunl]

In D76937#1950077, @rjmccall wrote:

Can you explain what exactly the emission/semantic model is for variables? Normal code-generation absolutely triggers the emission of many variables lazily (e.g. internal-linkage globals, C++ inline variables); and any variable that's *not* being emitted lazily actually needs to be treated as a potential root into the delayed-diagnostic graph.

Currently only in the following case a global variable can change the emission state of a function:

int foobar3() { return 1; }
#pragma omp declare target
int (*C)() = &foobar3;
#pragma omp end declare target

The global variable needs to be in a omp declare target directive. Its initializer references a host function.

This will cause the function emitted in device compilation.

This is not transitive for variable declaration/references, e.g. the following case will not cause foobar3 to be emitted in device compilation, unless variable C itself is enclosed in omp declare target directive.

int foobar3() { return 1; }
int (*C)() = &foobar3;
#pragma omp declare target
int (*D)() = C;
#pragma omp end declare target

Since we logged all such variable declarations in DeclsToCheckForDeferredDiags, we only need to check variables decls in DeclsToCheckForDeferredDiags and do not need to check variable decls in DeclRefExpr.

[rjmccall]

In D76937#1950764, @yaxunl wrote:

In D76937#1950077, @rjmccall wrote:

Can you explain what exactly the emission/semantic model is for variables? Normal code-generation absolutely triggers the emission of many variables lazily (e.g. internal-linkage globals, C++ inline variables); and any variable that's *not* being emitted lazily actually needs to be treated as a potential root into the delayed-diagnostic graph.

Currently only in the following case a global variable can change the emission state of a function:

int foobar3() { return 1; }
#pragma omp declare target
int (*C)() = &foobar3;
#pragma omp end declare target

The global variable needs to be in a omp declare target directive. Its initializer references a host function.

This will cause the function emitted in device compilation.

This is not transitive for variable declaration/references, e.g. the following case will not cause foobar3 to be emitted in device compilation, unless variable C itself is enclosed in omp declare target directive.

int foobar3() { return 1; }
int (*C)() = &foobar3;
#pragma omp declare target
int (*D)() = C;
#pragma omp end declare target

Since we logged all such variable declarations in DeclsToCheckForDeferredDiags, we only need to check variables decls in DeclsToCheckForDeferredDiags and do not need to check variable decls in DeclRefExpr.

Okay. There are other AST nodes besides DeclRefExpr that can cause an ODR-use of a VarDecl; you could add special cases for them, too, but part of the point of this design is so that you can stop worrying about that kind of thing. I think I've said this before, but you're making things harder for yourself by re-using visitUsedDecl as both the "callback" of the UsedDeclVisitor and the original method that you call for things from DeclsToCheckForDeferredDiags. Just use visitUsedDecl for the former, make a new method for the latter, and extract any common logic you want (e.g. for FunctionDecls) into helper methods that both of them can use. That way you can just unconditionally ignore non-functions in visitUsedDecl, assert that DeclsToCheckForDeferredDiags only contains the kinds of declarations you think it should, and so on.

[yaxunl]

Revised by John's comments. Also only check file scope variables.

[yaxunl]

fix assert message

[rjmccall]

Thanks, LGTM.

clang/lib/Sema/Sema.cpp

public:
 
   void VisitOMPTargetDirective(OMPTargetDirective *Node) {
     ++InOMPDeviceContext;
     Inherited::VisitOMPTargetDirective(Node);
     --InOMPDeviceContext;
   }
 
   void visitUsedDecl(SourceLocation Loc, Decl *D) {
-    if (auto *FD = dyn_cast<FunctionDecl>(D)) {
-      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(D))
-        return;
-      // Finalize analysis of OpenMP-specific constructs.
-      if (Caller && S.LangOpts.OpenMP && UseStack.size() == 1)
-        S.finalizeOpenMPDelayedAnalysis(Caller, FD, Loc);
-      if (Caller)
-        S.DeviceKnownEmittedFns[FD] = {Caller, Loc};
-      if (ShouldEmit || InOMPDeviceContext)
-        S.emitDeferredDiags(FD, Caller);
-      Visited.insert(D);
-      UseStack.push_back(FD);
-      if (auto *S = FD->getBody()) {
-        this->Visit(S);
-      }
-      UseStack.pop_back();
-      Visited.erase(D);
-    } else if (auto *VD = dyn_cast<VarDecl>(D)) {
-      if (auto *Init = VD->getInit()) {
-        auto DevTy = OMPDeclareTargetDeclAttr::getDeviceType(VD);
-        bool IsDev = DevTy && (*DevTy == OMPDeclareTargetDeclAttr::DT_NoHost ||
-                               *DevTy == OMPDeclareTargetDeclAttr::DT_Any);
-        if (IsDev)
-          ++InOMPDeviceContext;
-        this->Visit(Init);
-        if (IsDev)
-          --InOMPDeviceContext;
-      }
-    } else
-      Inherited::visitUsedDecl(Loc, D);
+    if (isa<VarDecl>(D))
+      return;
+    if (auto *FD = dyn_cast<FunctionDecl>(D))
+      checkFunc(Loc, FD);
+    else
+      Inherited::visitUsedDecl(Loc, D);
+  }
+
+  void checkVar(VarDecl *VD) {
+    assert(VD->isFileVarDecl() &&
+           "Should only check file-scope variables");
+    if (auto *Init = VD->getInit()) {
+      auto DevTy = OMPDeclareTargetDeclAttr::getDeviceType(VD);
+      bool IsDev = DevTy && (*DevTy == OMPDeclareTargetDeclAttr::DT_NoHost ||
+                             *DevTy == OMPDeclareTargetDeclAttr::DT_Any);
+      if (IsDev)
+        ++InOMPDeviceContext;
+      this->Visit(Init);
+      if (IsDev)
+        --InOMPDeviceContext;
+    }
+  }
+
+  void checkFunc(SourceLocation Loc, FunctionDecl *FD) {
+        S.shouldIgnoreInHostDeviceCheck(FD) || Visited.count(FD))
+    Visited.insert(FD);
+    Visited.erase(FD);
+
+  void checkRecordedDecl(Decl *D) {
+    if (auto *FD = dyn_cast<FunctionDecl>(D))
+      checkFunc(SourceLocation(), FD);
+    else
+      checkVar(cast<VarDecl>(D));
   }
 };
 } // namespace
 
 void Sema::emitDeferredDiags() {
   if (ExternalSource)
     ExternalSource->ReadDeclsToCheckForDeferredDiags(
         DeclsToCheckForDeferredDiags);
 
   if ((DeviceDeferredDiags.empty() && !LangOpts.OpenMP) ||
       DeclsToCheckForDeferredDiags.empty())
     return;
 
   DeferredDiagnosticsEmitter DDE(*this);
   for (auto D : DeclsToCheckForDeferredDiags)
-    DDE.visitUsedDecl(SourceLocation(), D);
+    DDE.checkRecordedDecl(D);
 }
 
 // In CUDA, there are some constructs which may appear in semantically-valid
 // code, but trigger errors if we ever generate code for the function in which
 // they appear.  Essentially every construct you're not allowed to use on the
 // device falls into this category, because you are allowed to use these
 // constructs in a __host__ __device__ function, but only if that function is
 // never codegen'ed on the device.

clang/lib/Sema/SemaDecl.cpp

void Sema::AddInitializerToDecl(Decl *RealDecl, Expr *Init, bool DirectInit) {
   if (CXXDirectInit) {
     assert(DirectInit && "Call-style initializer must be direct init.");
     VDecl->setInitStyle(VarDecl::CallInit);
   } else if (DirectInit) {
     // This must be list-initialization. No other way is direct-initialization.
     VDecl->setInitStyle(VarDecl::ListInit);
   }
 
-  if (LangOpts.OpenMP && VDecl->hasGlobalStorage())
+  if (LangOpts.OpenMP && VDecl->isFileVarDecl())
     DeclsToCheckForDeferredDiags.push_back(VDecl);
   CheckCompleteVariableDeclaration(VDecl);
 }
 
 /// ActOnInitializerError - Given that there was an error parsing an
 /// initializer for the given declaration, try to return to some form
 /// of sanity.
 void Sema::ActOnInitializerError(Decl *D) {

clang/test/OpenMP/nvptx_target_exceptions_messages.cpp

-// RUN: %clang_cc1 -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc -fexceptions -fcxx-exceptions
-// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple nvptx64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - -fexceptions -fcxx-exceptions -ferror-limit 100
+// RUN: %clang_cc1 -fopenmp -x c++ -triple powerpc64le-unknown-unknown \
+// RUN:   -verify=host -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc \
+// RUN:   %s -o %t-ppc-host.bc -fexceptions -fcxx-exceptions
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple nvptx64-unknown-unknown \
+// RUN:   -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s \
+// RUN:   -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - \
+// RUN:   -fexceptions -fcxx-exceptions -ferror-limit 100
 
 #ifndef HEADER
 #define HEADER
 
 template <typename T>
 class TemplateClass {
   T a;
 public:

 int (*A)() = &foobar1;
 #pragma omp declare target
 int (*B)() = &foobar2;
 #pragma omp end declare target
 
 int foobar1() { throw 1; }
 int foobar2() { throw 1; } // expected-error {{cannot use 'throw' with exceptions disabled}}
 
+
+int foobar3();
+int (*C)() = &foobar3; // expected-warning {{declaration is not declared in any declare target region}}
+                       // host-warning@-1 {{declaration is not declared in any declare target region}}
+#pragma omp declare target
+int (*D)() = C; // expected-note {{used here}}
+                // host-note@-1 {{used here}}
+#pragma omp end declare target
+int foobar3() { throw 1; }
+
+// Check no infinite recursion in deferred diagnostic emitter.
+long E = (long)&E;
+
 #endif // HEADER