[OPENMP]Fix PR44578: crash in target construct with captured global.

Target regions have implicit outer region which may erroneously capture
some globals when it should not. It may lead to a compiler crash at the
compile time.

clang/include/clang/Sema/Sema.h

@@ -9681,7 +9681,8 @@ public:
   /// Check if the specified variable is captured  by 'target' directive.
   /// \param Level Relative level of nested OpenMP construct for that the check
   /// is performed.
-  bool isOpenMPTargetCapturedDecl(const ValueDecl *D, unsigned Level) const;
+  bool isOpenMPTargetCapturedDecl(const ValueDecl *D, unsigned Level,
+                                  unsigned CaptureLevel) const;
 
   ExprResult PerformOpenMPImplicitIntegerConversion(SourceLocation OpLoc,
                                                     Expr *Op);

clang/lib/Sema/SemaExpr.cpp

@@ -16280,8 +16280,10 @@ bool Sema::tryCaptureVariable(
               captureVariablyModifiedType(Context, QTy, OuterRSI);
             }
           }
-          bool IsTargetCap = !IsOpenMPPrivateDecl &&
+          bool IsTargetCap =
-                             isOpenMPTargetCapturedDecl(Var, RSI->OpenMPLevel);
+              !IsOpenMPPrivateDecl &&
+              isOpenMPTargetCapturedDecl(Var, RSI->OpenMPLevel,
+                                         RSI->OpenMPCaptureLevel);
           // When we detect target captures we are looking from inside the
           // target region, therefore we need to propagate the capture from the
           // enclosing region. Therefore, the capture is not initially nested.

clang/lib/Sema/SemaOpenMP.cpp

@@ -31,6 +31,7 @@
 #include "clang/Sema/SemaInternal.h"
 #include "llvm/ADT/IndexedMap.h"
 #include "llvm/ADT/PointerEmbeddedInt.h"
+#include "llvm/ADT/STLExtras.h"
 #include "llvm/Frontend/OpenMP/OMPConstants.h"
 using namespace clang;
 using namespace llvm::omp;
@@ -2010,7 +2011,23 @@ VarDecl *Sema::isOpenMPCapturedDecl(ValueDecl *D, bool CheckScopeInfo,
       //
       if (OMPDeclareTargetDeclAttr::isDeclareTargetDeclaration(VD))
         return nullptr;
-      return VD;
+      CapturedRegionScopeInfo *CSI = nullptr;
+      for (FunctionScopeInfo *FSI : llvm::drop_begin(
+               llvm::reverse(FunctionScopes),
+               CheckScopeInfo ? (FunctionScopes.size() - (StopAt + 1)) : 0)) {
+        if (!isa<CapturingScopeInfo>(FSI))
+          return nullptr;
+        if (auto *RSI = dyn_cast<CapturedRegionScopeInfo>(FSI))
+          if (RSI->CapRegionKind == CR_OpenMP) {
+            CSI = RSI;
+            break;
+          }
+      }
+      SmallVector<OpenMPDirectiveKind, 4> Regions;
+      getOpenMPCaptureRegions(Regions,
+                              DSAStack->getDirective(CSI->OpenMPLevel));
+      if (Regions[CSI->OpenMPCaptureLevel] != OMPD_task)
+        return VD;
     }
   }
 
@@ -2151,15 +2168,18 @@ void Sema::setOpenMPCaptureKind(FieldDecl *FD, const ValueDecl *D,
     FD->addAttr(OMPCaptureKindAttr::CreateImplicit(Context, OMPC));
 }
 
-bool Sema::isOpenMPTargetCapturedDecl(const ValueDecl *D,
+bool Sema::isOpenMPTargetCapturedDecl(const ValueDecl *D, unsigned Level,
-                                      unsigned Level) const {
+                                      unsigned CaptureLevel) const {
   assert(LangOpts.OpenMP && "OpenMP is not allowed");
   // Return true if the current level is no longer enclosed in a target region.
 
+  SmallVector<OpenMPDirectiveKind, 4> Regions;
+  getOpenMPCaptureRegions(Regions, DSAStack->getDirective(Level));
   const auto *VD = dyn_cast<VarDecl>(D);
   return VD && !VD->hasLocalStorage() &&
          DSAStack->hasExplicitDirective(isOpenMPTargetExecutionDirective,
-                                        Level);
+                                        Level) &&
+         Regions[CaptureLevel] != OMPD_task;
 }
 
 void Sema::DestroyDataSharingAttributesStack() { delete DSAStack; }

clang/test/OpenMP/target_depend_codegen.cpp

@@ -209,7 +209,7 @@ int foo(int n) {
 // CHECK-NEXT:  br i1 [[ERROR]], label %[[FAIL:[^,]+]], label %[[END:[^,]+]]
 // CHECK:       [[FAIL]]
 // CHECK:       [[BP0:%.+]] = load i[[SZ]]*, i[[SZ]]** %
-// CHECK:       [[BP1_I32:%.+]] = load i32, i32* %
+// CHECK:       [[BP1_I32:%.+]] = load i32, i32* @
 // CHECK-64:    [[BP1_CAST:%.+]] = bitcast i[[SZ]]* [[BP1_PTR:%.+]] to i32*
 // CHECK-64:    store i32 [[BP1_I32]], i32* [[BP1_CAST]],
 // CHECK-32:    store i32 [[BP1_I32]], i32* [[BP1_PTR:%.+]],
@@ -223,7 +223,7 @@ int foo(int n) {
 // CHECK:       call void (i8*, ...) %
 // CHECK:       [[DEVICE_CAP:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* %{{.+}}, i32 0, i32 2
 // CHECK:       [[BP0:%.+]] = load i[[SZ]]*, i[[SZ]]** %
-// CHECK:       [[BP1_I32:%.+]] = load i32, i32* %
+// CHECK:       [[BP1_I32:%.+]] = load i32, i32* @
 // CHECK-64:    [[BP1_CAST:%.+]] = bitcast i[[SZ]]* [[BP1_PTR:%.+]] to i32*
 // CHECK-64:    store i32 [[BP1_I32]], i32* [[BP1_CAST]],
 // CHECK-32:    store i32 [[BP1_I32]], i32* [[BP1_PTR:%.+]],

clang/test/OpenMP/target_messages.cpp

@@ -112,4 +112,12 @@ int main(int argc, char **argv) {
 
   return 0;
 }
+
+template <class> struct a { static bool b; };
+template <class c, bool = a<c>::b> void e(c) { // expected-note {{candidate template ignored: substitution failure [with c = int]: non-type template argument is not a constant expression}}
+#pragma omp target
+  {
+    int d ; e(d); // expected-error {{no matching function for call to 'e'}}
+  }
+}
 #endif

clang/test/OpenMP/target_parallel_depend_codegen.cpp

@@ -209,7 +209,7 @@ int foo(int n) {
 // CHECK-NEXT:  br i1 [[ERROR]], label %[[FAIL:[^,]+]], label %[[END:[^,]+]]
 // CHECK:       [[FAIL]]
 // CHECK:       [[BP0:%.+]] = load i[[SZ]]*, i[[SZ]]** %
-// CHECK:       [[BP1_I32:%.+]] = load i32, i32* %
+// CHECK:       [[BP1_I32:%.+]] = load i32, i32* @
 // CHECK-64:    [[BP1_CAST:%.+]] = bitcast i[[SZ]]* [[BP1_PTR:%.+]] to i32*
 // CHECK-64:    store i32 [[BP1_I32]], i32* [[BP1_CAST]],
 // CHECK-32:    store i32 [[BP1_I32]], i32* [[BP1_PTR:%.+]],
@@ -223,7 +223,7 @@ int foo(int n) {
 // CHECK:       call void (i8*, ...) %
 // CHECK:       [[DEVICE_CAP:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* %{{.+}}, i32 0, i32 2
 // CHECK:       [[BP0:%.+]] = load i[[SZ]]*, i[[SZ]]** %
-// CHECK:       [[BP1_I32:%.+]] = load i32, i32* %
+// CHECK:       [[BP1_I32:%.+]] = load i32, i32* @
 // CHECK-64:    [[BP1_CAST:%.+]] = bitcast i[[SZ]]* [[BP1_PTR:%.+]] to i32*
 // CHECK-64:    store i32 [[BP1_I32]], i32* [[BP1_CAST]],
 // CHECK-32:    store i32 [[BP1_I32]], i32* [[BP1_PTR:%.+]],

clang/test/OpenMP/target_parallel_for_depend_codegen.cpp

@@ -209,7 +209,7 @@ int foo(int n) {
 // CHECK-NEXT:  br i1 [[ERROR]], label %[[FAIL:[^,]+]], label %[[END:[^,]+]]
 // CHECK:       [[FAIL]]
 // CHECK:       [[BP0:%.+]] = load i[[SZ]]*, i[[SZ]]** %
-// CHECK:       [[BP1_I32:%.+]] = load i32, i32* %
+// CHECK:       [[BP1_I32:%.+]] = load i32, i32* @
 // CHECK-64:    [[BP1_CAST:%.+]] = bitcast i[[SZ]]* [[BP1_PTR:%.+]] to i32*
 // CHECK-64:    store i32 [[BP1_I32]], i32* [[BP1_CAST]],
 // CHECK-32:    store i32 [[BP1_I32]], i32* [[BP1_PTR:%.+]],
@@ -223,7 +223,7 @@ int foo(int n) {
 // CHECK:       call void (i8*, ...) %
 // CHECK:       [[DEVICE_CAP:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* %{{.+}}, i32 0, i32 2
 // CHECK:       [[BP0:%.+]] = load i[[SZ]]*, i[[SZ]]** %
-// CHECK:       [[BP1_I32:%.+]] = load i32, i32* %
+// CHECK:       [[BP1_I32:%.+]] = load i32, i32* @
 // CHECK-64:    [[BP1_CAST:%.+]] = bitcast i[[SZ]]* [[BP1_PTR:%.+]] to i32*
 // CHECK-64:    store i32 [[BP1_I32]], i32* [[BP1_CAST]],
 // CHECK-32:    store i32 [[BP1_I32]], i32* [[BP1_PTR:%.+]],

clang/test/OpenMP/target_parallel_for_simd_depend_codegen.cpp

@@ -209,7 +209,7 @@ int foo(int n) {
 // CHECK-NEXT:  br i1 [[ERROR]], label %[[FAIL:[^,]+]], label %[[END:[^,]+]]
 // CHECK:       [[FAIL]]
 // CHECK:       [[BP0:%.+]] = load i[[SZ]]*, i[[SZ]]** %
-// CHECK:       [[BP1_I32:%.+]] = load i32, i32* %
+// CHECK:       [[BP1_I32:%.+]] = load i32, i32* @
 // CHECK-64:    [[BP1_CAST:%.+]] = bitcast i[[SZ]]* [[BP1_PTR:%.+]] to i32*
 // CHECK-64:    store i32 [[BP1_I32]], i32* [[BP1_CAST]],
 // CHECK-32:    store i32 [[BP1_I32]], i32* [[BP1_PTR:%.+]],
@@ -223,7 +223,7 @@ int foo(int n) {
 // CHECK:       call void (i8*, ...) %
 // CHECK:       [[DEVICE_CAP:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* %{{.+}}, i32 0, i32 2
 // CHECK:       [[BP0:%.+]] = load i[[SZ]]*, i[[SZ]]** %
-// CHECK:       [[BP1_I32:%.+]] = load i32, i32* %
+// CHECK:       [[BP1_I32:%.+]] = load i32, i32* @
 // CHECK-64:    [[BP1_CAST:%.+]] = bitcast i[[SZ]]* [[BP1_PTR:%.+]] to i32*
 // CHECK-64:    store i32 [[BP1_I32]], i32* [[BP1_CAST]],
 // CHECK-32:    store i32 [[BP1_I32]], i32* [[BP1_PTR:%.+]],

clang/test/OpenMP/target_simd_depend_codegen.cpp

@@ -209,7 +209,7 @@ int foo(int n) {
 // CHECK-NEXT:  br i1 [[ERROR]], label %[[FAIL:[^,]+]], label %[[END:[^,]+]]
 // CHECK:       [[FAIL]]
 // CHECK:       [[BP0:%.+]] = load i[[SZ]]*, i[[SZ]]** %
-// CHECK:       [[BP1_I32:%.+]] = load i32, i32* %
+// CHECK:       [[BP1_I32:%.+]] = load i32, i32* @
 // CHECK-64:    [[BP1_CAST:%.+]] = bitcast i[[SZ]]* [[BP1_PTR:%.+]] to i32*
 // CHECK-64:    store i32 [[BP1_I32]], i32* [[BP1_CAST]],
 // CHECK-32:    store i32 [[BP1_I32]], i32* [[BP1_PTR:%.+]],
@@ -223,7 +223,7 @@ int foo(int n) {
 // CHECK:       call void (i8*, ...) %
 // CHECK:       [[DEVICE_CAP:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* %{{.+}}, i32 0, i32 2
 // CHECK:       [[BP0:%.+]] = load i[[SZ]]*, i[[SZ]]** %
-// CHECK:       [[BP1_I32:%.+]] = load i32, i32* %
+// CHECK:       [[BP1_I32:%.+]] = load i32, i32* @
 // CHECK-64:    [[BP1_CAST:%.+]] = bitcast i[[SZ]]* [[BP1_PTR:%.+]] to i32*
 // CHECK-64:    store i32 [[BP1_I32]], i32* [[BP1_CAST]],
 // CHECK-32:    store i32 [[BP1_I32]], i32* [[BP1_PTR:%.+]],

clang/test/OpenMP/target_teams_depend_codegen.cpp

@@ -209,7 +209,7 @@ int foo(int n) {
 // CHECK-NEXT:  br i1 [[ERROR]], label %[[FAIL:[^,]+]], label %[[END:[^,]+]]
 // CHECK:       [[FAIL]]
 // CHECK:       [[BP0:%.+]] = load i[[SZ]]*, i[[SZ]]** %
-// CHECK:       [[BP1_I32:%.+]] = load i32, i32* %
+// CHECK:       [[BP1_I32:%.+]] = load i32, i32* @
 // CHECK-64:    [[BP1_CAST:%.+]] = bitcast i[[SZ]]* [[BP1_PTR:%.+]] to i32*
 // CHECK-64:    store i32 [[BP1_I32]], i32* [[BP1_CAST]],
 // CHECK-32:    store i32 [[BP1_I32]], i32* [[BP1_PTR:%.+]],
@@ -223,7 +223,7 @@ int foo(int n) {
 // CHECK:       call void (i8*, ...) %
 // CHECK:       [[DEVICE_CAP:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* %{{.+}}, i32 0, i32 2
 // CHECK:       [[BP0:%.+]] = load i[[SZ]]*, i[[SZ]]** %
-// CHECK:       [[BP1_I32:%.+]] = load i32, i32* %
+// CHECK:       [[BP1_I32:%.+]] = load i32, i32* @
 // CHECK-64:    [[BP1_CAST:%.+]] = bitcast i[[SZ]]* [[BP1_PTR:%.+]] to i32*
 // CHECK-64:    store i32 [[BP1_I32]], i32* [[BP1_CAST]],
 // CHECK-32:    store i32 [[BP1_I32]], i32* [[BP1_PTR:%.+]],

clang/test/OpenMP/target_teams_distribute_depend_codegen.cpp

@@ -209,7 +209,7 @@ int foo(int n) {
 // CHECK-NEXT:  br i1 [[ERROR]], label %[[FAIL:[^,]+]], label %[[END:[^,]+]]
 // CHECK:       [[FAIL]]
 // CHECK:       [[BP0:%.+]] = load i[[SZ]]*, i[[SZ]]** %
-// CHECK:       [[BP1_I32:%.+]] = load i32, i32* %
+// CHECK:       [[BP1_I32:%.+]] = load i32, i32* @
 // CHECK-64:    [[BP1_CAST:%.+]] = bitcast i[[SZ]]* [[BP1_PTR:%.+]] to i32*
 // CHECK-64:    store i32 [[BP1_I32]], i32* [[BP1_CAST]],
 // CHECK-32:    store i32 [[BP1_I32]], i32* [[BP1_PTR:%.+]],
@@ -223,7 +223,7 @@ int foo(int n) {
 // CHECK:       call void (i8*, ...) %
 // CHECK:       [[DEVICE_CAP:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* %{{.+}}, i32 0, i32 2
 // CHECK:       [[BP0:%.+]] = load i[[SZ]]*, i[[SZ]]** %
-// CHECK:       [[BP1_I32:%.+]] = load i32, i32* %
+// CHECK:       [[BP1_I32:%.+]] = load i32, i32* @
 // CHECK-64:    [[BP1_CAST:%.+]] = bitcast i[[SZ]]* [[BP1_PTR:%.+]] to i32*
 // CHECK-64:    store i32 [[BP1_I32]], i32* [[BP1_CAST]],
 // CHECK-32:    store i32 [[BP1_I32]], i32* [[BP1_PTR:%.+]],

clang/test/OpenMP/target_teams_distribute_parallel_for_depend_codegen.cpp

@@ -209,7 +209,7 @@ int foo(int n) {
 // CHECK-NEXT:  br i1 [[ERROR]], label %[[FAIL:[^,]+]], label %[[END:[^,]+]]
 // CHECK:       [[FAIL]]
 // CHECK:       [[BP0:%.+]] = load i[[SZ]]*, i[[SZ]]** %
-// CHECK:       [[BP1_I32:%.+]] = load i32, i32* %
+// CHECK:       [[BP1_I32:%.+]] = load i32, i32* @
 // CHECK-64:    [[BP1_CAST:%.+]] = bitcast i[[SZ]]* [[BP1_PTR:%.+]] to i32*
 // CHECK-64:    store i32 [[BP1_I32]], i32* [[BP1_CAST]],
 // CHECK-32:    store i32 [[BP1_I32]], i32* [[BP1_PTR:%.+]],
@@ -223,7 +223,7 @@ int foo(int n) {
 // CHECK:       call void (i8*, ...) %
 // CHECK:       [[DEVICE_CAP:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* %{{.+}}, i32 0, i32 2
 // CHECK:       [[BP0:%.+]] = load i[[SZ]]*, i[[SZ]]** %
-// CHECK:       [[BP1_I32:%.+]] = load i32, i32* %
+// CHECK:       [[BP1_I32:%.+]] = load i32, i32* @
 // CHECK-64:    [[BP1_CAST:%.+]] = bitcast i[[SZ]]* [[BP1_PTR:%.+]] to i32*
 // CHECK-64:    store i32 [[BP1_I32]], i32* [[BP1_CAST]],
 // CHECK-32:    store i32 [[BP1_I32]], i32* [[BP1_PTR:%.+]],

clang/test/OpenMP/target_teams_distribute_parallel_for_simd_depend_codegen.cpp

@@ -209,7 +209,7 @@ int foo(int n) {
 // CHECK-NEXT:  br i1 [[ERROR]], label %[[FAIL:[^,]+]], label %[[END:[^,]+]]
 // CHECK:       [[FAIL]]
 // CHECK:       [[BP0:%.+]] = load i[[SZ]]*, i[[SZ]]** %
-// CHECK:       [[BP1_I32:%.+]] = load i32, i32* %
+// CHECK:       [[BP1_I32:%.+]] = load i32, i32* @
 // CHECK-64:    [[BP1_CAST:%.+]] = bitcast i[[SZ]]* [[BP1_PTR:%.+]] to i32*
 // CHECK-64:    store i32 [[BP1_I32]], i32* [[BP1_CAST]],
 // CHECK-32:    store i32 [[BP1_I32]], i32* [[BP1_PTR:%.+]],
@@ -223,7 +223,7 @@ int foo(int n) {
 // CHECK:       call void (i8*, ...) %
 // CHECK:       [[DEVICE_CAP:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* %{{.+}}, i32 0, i32 2
 // CHECK:       [[BP0:%.+]] = load i[[SZ]]*, i[[SZ]]** %
-// CHECK:       [[BP1_I32:%.+]] = load i32, i32* %
+// CHECK:       [[BP1_I32:%.+]] = load i32, i32* @
 // CHECK-64:    [[BP1_CAST:%.+]] = bitcast i[[SZ]]* [[BP1_PTR:%.+]] to i32*
 // CHECK-64:    store i32 [[BP1_I32]], i32* [[BP1_CAST]],
 // CHECK-32:    store i32 [[BP1_I32]], i32* [[BP1_PTR:%.+]],

clang/test/OpenMP/target_teams_distribute_simd_depend_codegen.cpp

@@ -209,7 +209,7 @@ int foo(int n) {
 // CHECK-NEXT:  br i1 [[ERROR]], label %[[FAIL:[^,]+]], label %[[END:[^,]+]]
 // CHECK:       [[FAIL]]
 // CHECK:       [[BP0:%.+]] = load i[[SZ]]*, i[[SZ]]** %
-// CHECK:       [[BP1_I32:%.+]] = load i32, i32* %
+// CHECK:       [[BP1_I32:%.+]] = load i32, i32* @
 // CHECK-64:    [[BP1_CAST:%.+]] = bitcast i[[SZ]]* [[BP1_PTR:%.+]] to i32*
 // CHECK-64:    store i32 [[BP1_I32]], i32* [[BP1_CAST]],
 // CHECK-32:    store i32 [[BP1_I32]], i32* [[BP1_PTR:%.+]],
@@ -223,7 +223,7 @@ int foo(int n) {
 // CHECK:       call void (i8*, ...) %
 // CHECK:       [[DEVICE_CAP:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* %{{.+}}, i32 0, i32 2
 // CHECK:       [[BP0:%.+]] = load i[[SZ]]*, i[[SZ]]** %
-// CHECK:       [[BP1_I32:%.+]] = load i32, i32* %
+// CHECK:       [[BP1_I32:%.+]] = load i32, i32* @
 // CHECK-64:    [[BP1_CAST:%.+]] = bitcast i[[SZ]]* [[BP1_PTR:%.+]] to i32*
 // CHECK-64:    store i32 [[BP1_I32]], i32* [[BP1_CAST]],
 // CHECK-32:    store i32 [[BP1_I32]], i32* [[BP1_PTR:%.+]],