[OPENMP]Do not crash for globals in inner regions with outer target

region.

If the global variable is used in the target region,it is always
captured, if not marked as declare target.

diff --git a/clang/lib/Sema/SemaOpenMP.cpp b/clang/lib/Sema/SemaOpenMP.cpp
index cd20b6b..7b62c84 100644 (file)
--- a/clang/lib/Sema/SemaOpenMP.cpp
+++ b/clang/lib/Sema/SemaOpenMP.cpp
@@ -2417,17 +2417,20 @@ bool Sema::isOpenMPGlobalCapturedDecl(ValueDecl *D, unsigned Level,
 
   if (const auto *VD = dyn_cast<VarDecl>(D)) {
     if (!VD->hasLocalStorage()) {
+      if (isInOpenMPTargetExecutionDirective())
+        return true;
       DSAStackTy::DSAVarData TopDVar =
           DSAStack->getTopDSA(D, /*FromParent=*/false);
       unsigned NumLevels =
           getOpenMPCaptureLevels(DSAStack->getDirective(Level));
       if (Level == 0)
         return (NumLevels == CaptureLevel + 1) && TopDVar.CKind != OMPC_shared;
-      DSAStackTy::DSAVarData DVar = DSAStack->getImplicitDSA(D, Level - 1);
-      return DVar.CKind != OMPC_shared ||
-             isOpenMPGlobalCapturedDecl(
-                 D, Level - 1,
-                 getOpenMPCaptureLevels(DSAStack->getDirective(Level - 1)) - 1);
+      do {
+        --Level;
+        DSAStackTy::DSAVarData DVar = DSAStack->getImplicitDSA(D, Level);
+        if (DVar.CKind != OMPC_shared)
+          return true;
+      } while (Level >= 0);
     }
   }
   return true;

diff --git a/clang/test/OpenMP/target_codegen.cpp b/clang/test/OpenMP/target_codegen.cpp
index 55b03aa..c8570bd 100644 (file)
--- a/clang/test/OpenMP/target_codegen.cpp
+++ b/clang/test/OpenMP/target_codegen.cpp
@@ -706,6 +706,8 @@ int bar(int n){
 
 // CHECK:       [[IFEND]]
 
+// OMP45: define internal void @__omp_offloading_{{.+}}_{{.+}}bar{{.+}}_l838(i[[SZ]] %{{.+}})
+
 // OMP45: define {{.*}}@{{.*}}zee{{.*}}
 
 // OMP45:       [[LOCAL_THIS:%.+]] = alloca [[S2]]*
@@ -803,6 +805,7 @@ int bar(int n){
 // CHECK-DAG:   load i16, i16* [[REF_AA]]
 // CHECK-DAG:   getelementptr inbounds [10 x i32], [10 x i32]* [[REF_B]], i[[SZ]] 0, i[[SZ]] 2
 
+// OMP50: define internal void @__omp_offloading_{{.+}}_{{.+}}bar{{.+}}_l838(i[[SZ]] %{{.+}})
 
 // OMP50: define {{.*}}@{{.*}}zee{{.*}}
 
@@ -833,7 +836,11 @@ int bar(int n){
 void bar () {
 #define pragma_target _Pragma("omp target")
 pragma_target
-{}
+{
+  global = 0;
+#pragma omp parallel shared(global)
+  global = 1;
+}
 }
 
 class S2 {