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

[Alexey Bataev via cfe-commits]

Author: Alexey Bataev
Date: 2020-08-27T17:07:53-04:00
New Revision: ba1de5f2f7b078f69d5f6b0fe3af4911f76bb8fd

URL: https://github.com/llvm/llvm-project/commit/ba1de5f2f7b078f69d5f6b0fe3af4911f76bb8fd
DIFF: https://github.com/llvm/llvm-project/commit/ba1de5f2f7b078f69d5f6b0fe3af4911f76bb8fd.diff

LOG: [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.

Added: 
    

Modified: 
    clang/lib/Sema/SemaOpenMP.cpp
    clang/test/OpenMP/target_codegen.cpp

Removed: 
    


################################################################################
diff  --git a/clang/lib/Sema/SemaOpenMP.cpp b/clang/lib/Sema/SemaOpenMP.cpp
index cd20b6bfdb98..7b62c841b48a 100644
--- 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 55b03aae00a5..c8570bdf6b65 100644
--- 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 {