[OPENMP]Fix processing of the clauses on target combined directives.

For constants with the predefined data-sharing clauses we may had
troubles with the target combined directives. It may cause compiler
crash in some corner cases.

git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@350127 91177308-0d34-0410-b5e6-96231b3b80d8

diff --git a/lib/Sema/SemaOpenMP.cpp b/lib/Sema/SemaOpenMP.cpp
index 24e9ef4aaa2241310c6194b3e8043590739f5c72..b4eb4664763cfc45453268d4901881e3184ad73b 100644 (file)
--- a/lib/Sema/SemaOpenMP.cpp
+++ b/lib/Sema/SemaOpenMP.cpp
@@ -1179,10 +1179,13 @@ const DSAStackTy::DSAVarData DSAStackTy::getTopDSA(ValueDecl *D,
         RD->hasMutableFields())) {
     // Variables with const-qualified type having no mutable member may be
     // listed in a firstprivate clause, even if they are static data members.
-    DSAVarData DVarTemp =
-        hasDSA(D, [](OpenMPClauseKind C) { return C == OMPC_firstprivate; },
-               MatchesAlways, FromParent);
-    if (DVarTemp.CKind == OMPC_firstprivate && DVarTemp.RefExpr)
+    DSAVarData DVarTemp = hasInnermostDSA(
+        D,
+        [](OpenMPClauseKind C) {
+          return C == OMPC_firstprivate || C == OMPC_shared;
+        },
+        MatchesAlways, FromParent);
+    if (DVarTemp.CKind != OMPC_unknown && DVarTemp.RefExpr)
       return DVarTemp;
 
     DVar.CKind = OMPC_shared;

diff --git a/test/OpenMP/target_teams_codegen.cpp b/test/OpenMP/target_teams_codegen.cpp
index bd6eb8cb1c149edd870c5088db47751c42ee3ca2..044893311186fb2960aa1bf74090f1af0743b0c1 100644 (file)
--- a/test/OpenMP/target_teams_codegen.cpp
+++ b/test/OpenMP/target_teams_codegen.cpp
@@ -74,6 +74,8 @@
 // CHECK-DAG: @{{.*}} = weak constant i8 0
 // CHECK-DAG: @{{.*}} = weak constant i8 0
 // CHECK-DAG: @{{.*}} = weak constant i8 0
+// CHECK-DAG: @{{.*}} = weak constant i8 0
+// CHECK-DAG: @{{.*}} = weak constant i8 0
 
 // TCHECK: @{{.+}} = weak constant [[ENTTY]]
 // TCHECK: @{{.+}} = weak constant [[ENTTY]]
@@ -82,6 +84,8 @@
 // TCHECK: @{{.+}} = weak constant [[ENTTY]]
 // TCHECK: @{{.+}} = weak constant [[ENTTY]]
 // TCHECK: @{{.+}} = weak constant [[ENTTY]]
+// TCHECK: @{{.+}} = weak constant [[ENTTY]]
+// TCHECK: @{{.+}} = weak constant [[ENTTY]]
 // TCHECK-NOT: @{{.+}} = weak constant [[ENTTY]]
 
 // Check if offloading descriptor is created.
@@ -836,4 +840,21 @@ int bar(int n){
 // CHECK:       define internal {{.*}}void [[OMP_OUTLINED7]](i32* noalias %.global_tid., i32* noalias %.bound_tid., i[[SZ]] %{{.+}}, i[[SZ]] %{{.+}}, [10 x i32]* {{.+}})
 // To reduce complexity, we're only going as far as validating the signature of the outlined parallel function.
 
+void foo1() {
+  const int n = 0;
+  #pragma omp target teams shared(n)
+  #pragma omp parallel firstprivate(n)
+  (void)n;
+}
+void foo() {
+  const int n = 0;
+  #pragma omp target teams firstprivate(n)
+  #pragma omp parallel shared(n)
+  (void)n;
+}
+
+// define {{.*}}void @__omp_offloading_{{.*}}foo1{{.*}}_l841(i[[SZ]] %{{.+}})
+// define internal void {{@.+}}(i32* {{.+}}, i32* {{.+}}, i[[SZ]] %{{.+}})
+// define {{.*}}void @__omp_offloading_{{.*}}foo1{{.*}}_l847(i[[SZ]] %{{.+}})
+// define internal void {{@.+}}(i32* {{.+}}, i32* {{.+}}, i32* dereferenceable{{.+}})
 #endif