From b136b7a78774107943fe94051c42b5a968a3ad3f Mon Sep 17 00:00:00 2001
From: Jakub Jelinek <jakub@redhat.com>
Date: Wed, 21 Jul 2021 09:45:02 +0200
Subject: [PATCH] openmp: Fix up omp_check_private [PR101535]

The target data construct shouldn't affect omp_check_private, unless
the decl there is privatized (use_device_* clauses).  The routine
had some code for that, but it just did continue; in a loop that looped
only if the region type is one of selected 4 kinds, so effectively resulted
in return false; instead of looping again.  And not diagnosing lastprivate
(or reduction etc.) on a variable that is private to containing parallel
results in ICEs later on, as there is no original list item to which store
the last result.
The target construct is unclear as it has an implicit parallel region
and it is not obvious if the data privatization clauses on the construct
shall be treated as data privatization on the implicit parallel or just
on the target.  For now treat those as privatization on the implicit
parallel, but treat map clauses as shared on the implicit parallel.

2021-07-21  Jakub Jelinek  <jakub@redhat.com>

	PR middle-end/101535
	* gimplify.c (omp_check_private): Properly skip ORT_TARGET_DATA
	contexts in which decl isn't privatized and for ORT_TARGET return
	false if decl is mapped.

	* c-c++-common/gomp/pr101535-1.c: New test.
	* c-c++-common/gomp/pr101535-2.c: New test.
---
 gcc/gimplify.c                               | 21 +++++++++----
 gcc/testsuite/c-c++-common/gomp/pr101535-1.c | 31 ++++++++++++++++++++
 gcc/testsuite/c-c++-common/gomp/pr101535-2.c | 11 +++++++
 3 files changed, 58 insertions(+), 5 deletions(-)
 create mode 100644 gcc/testsuite/c-c++-common/gomp/pr101535-1.c
 create mode 100644 gcc/testsuite/c-c++-common/gomp/pr101535-2.c

diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index 93a2121fea1..5d43f76f002 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -7798,7 +7798,13 @@ omp_check_private (struct gimplify_omp_ctx *ctx, tree decl, bool copyprivate)
 
       if ((ctx->region_type & (ORT_TARGET | ORT_TARGET_DATA)) != 0
 	  && (n == NULL || (n->value & GOVD_DATA_SHARE_CLASS) == 0))
-	continue;
+	{
+	  if ((ctx->region_type & ORT_TARGET_DATA) != 0
+	      || n == NULL
+	      || (n->value & GOVD_MAP) == 0)
+	    continue;
+	  return false;
+	}
 
       if (n != NULL)
 	{
@@ -7807,11 +7813,16 @@ omp_check_private (struct gimplify_omp_ctx *ctx, tree decl, bool copyprivate)
 	    return false;
 	  return (n->value & GOVD_SHARED) == 0;
 	}
+
+      if (ctx->region_type == ORT_WORKSHARE
+	  || ctx->region_type == ORT_TASKGROUP
+	  || ctx->region_type == ORT_SIMD
+	  || ctx->region_type == ORT_ACC)
+	continue;
+
+      break;
     }
-  while (ctx->region_type == ORT_WORKSHARE
-	 || ctx->region_type == ORT_TASKGROUP
-	 || ctx->region_type == ORT_SIMD
-	 || ctx->region_type == ORT_ACC);
+  while (1);
   return false;
 }
 
diff --git a/gcc/testsuite/c-c++-common/gomp/pr101535-1.c b/gcc/testsuite/c-c++-common/gomp/pr101535-1.c
new file mode 100644
index 00000000000..8285ce011df
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/pr101535-1.c
@@ -0,0 +1,31 @@
+/* PR middle-end/101535 */
+
+void
+foo (void)
+{
+  int a = 1, i;
+  #pragma omp target data map(to:a)
+  #pragma omp for lastprivate(i)	/* { dg-error "lastprivate variable 'i' is private in outer context" } */
+  for (i = 1; i < 2; i++)
+    ;
+}
+
+void
+bar (void)
+{
+  int a = 1, i;
+  #pragma omp target private(i)
+  #pragma omp for lastprivate(i)	/* { dg-error "lastprivate variable 'i' is private in outer context" } */
+  for (i = 1; i < 2; i++)
+    ;
+}
+
+void
+baz (void)
+{
+  int a = 1, i;
+  #pragma omp target firstprivate(i)
+  #pragma omp for lastprivate(i)	/* { dg-error "lastprivate variable 'i' is private in outer context" } */
+  for (i = 1; i < 2; i++)
+    ;
+}
diff --git a/gcc/testsuite/c-c++-common/gomp/pr101535-2.c b/gcc/testsuite/c-c++-common/gomp/pr101535-2.c
new file mode 100644
index 00000000000..23c84af4879
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/pr101535-2.c
@@ -0,0 +1,11 @@
+/* PR middle-end/101535 */
+
+void
+foo (void)
+{
+  int a = 1, i;
+  #pragma omp target map(tofrom:i)
+  #pragma omp for lastprivate(i)
+  for (i = 1; i < 2; i++)
+    ;
+}