OpenMP/C++: Fix (first)private clause with member variables [PR110347]

OpenMP permits '(first)private' for C++ member variables, which GCC handles
by tagging those by DECL_OMP_PRIVATIZED_MEMBER, adding a temporary VAR_DECL
and DECL_VALUE_EXPR pointing to the 'this->member_var' in the C++ front end.

The idea is that in omp-low.cc, the DECL_VALUE_EXPR is used before the
region (for 'firstprivate'; ignored for 'private') while in the region,
the DECL itself is used.

In gimplify, the value expansion is suppressed and deferred if the
  lang_hooks.decls.omp_disregard_value_expr (decl, shared)
returns true - which is never the case if 'shared' is true. In OpenMP 4.5,
only 'map' and 'use_device_ptr' was permitted for the 'target' directive.
And when OpenMP 5.0's 'private'/'firstprivate' clauses was added, the
the update that now 'shared' argument could be false was missed. The
respective check has now been added.

2024-03-01  Jakub Jelinek  <jakub@redhat.com>
	    Tobias Burnus  <tburnus@baylibre.com>

	PR c++/110347

gcc/ChangeLog:

	* gimplify.cc (omp_notice_variable): Fix 'shared' arg to
	lang_hooks.decls.omp_disregard_value_expr for
	(first)private in target regions.

libgomp/ChangeLog:

	* testsuite/libgomp.c++/target-lambda-3.C: Moved from
	gcc/testsuite/g++.dg/gomp/ and fixed is-mapped handling.
	* testsuite/libgomp.c++/target-lambda-1.C: Modify to also
	also work without offloading.
	* testsuite/libgomp.c++/firstprivate-1.C: New test.
	* testsuite/libgomp.c++/firstprivate-2.C: New test.
	* testsuite/libgomp.c++/private-1.C: New test.
	* testsuite/libgomp.c++/private-2.C: New test.
	* testsuite/libgomp.c++/target-lambda-4.C: New test.
	* testsuite/libgomp.c++/use_device_ptr-1.C: New test.

gcc/testsuite/ChangeLog:

	* g++.dg/gomp/target-lambda-1.C: Moved to become a
	run-time test under testsuite/libgomp.c++.

Co-authored-by: Tobias Burnus <tburnus@baylibre.com>

2 files changed, 13 insertions, 101 deletions

diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc
index 7f79b3c..6ebca96 100644
--- a/gcc/gimplify.cc
+++ b/gcc/gimplify.cc
@@ -8144,13 +8144,6 @@ omp_notice_variable (struct gimplify_omp_ctx *ctx, tree decl, bool in_code)
   n = splay_tree_lookup (ctx->variables, (splay_tree_key)decl);
   if ((ctx->region_type & ORT_TARGET) != 0)
     {
-      if (ctx->region_type & ORT_ACC)
-	/* For OpenACC, as remarked above, defer expansion.  */
-	shared = false;
-      else
-	shared = true;
-
-      ret = lang_hooks.decls.omp_disregard_value_expr (decl, shared);
       if (n == NULL)
 	{
 	  unsigned nflags = flags;
@@ -8275,9 +8268,22 @@ omp_notice_variable (struct gimplify_omp_ctx *ctx, tree decl, bool in_code)
 	    }
 	found_outer:
 	  omp_add_variable (ctx, decl, nflags);
+	  if (ctx->region_type & ORT_ACC)
+	    /* For OpenACC, as remarked above, defer expansion.  */
+	    shared = false;
+	  else
+	    shared = (nflags & (GOVD_PRIVATE | GOVD_FIRSTPRIVATE)) == 0;
+	  ret = lang_hooks.decls.omp_disregard_value_expr (decl, shared);
 	}
       else
 	{
+	  if (ctx->region_type & ORT_ACC)
+	    /* For OpenACC, as remarked above, defer expansion.  */
+	    shared = false;
+	  else
+	    shared = ((n->value | flags)
+		      & (GOVD_PRIVATE | GOVD_FIRSTPRIVATE)) == 0;
+	  ret = lang_hooks.decls.omp_disregard_value_expr (decl, shared);
 	  /* If nothing changed, there's nothing left to do.  */
 	  if ((n->value & flags) == flags)
 	    return ret;
diff --git a/gcc/testsuite/g++.dg/gomp/target-lambda-1.C b/gcc/testsuite/g++.dg/gomp/target-lambda-1.C
deleted file mode 100644
index 5ce8cea..0000000
--- a/gcc/testsuite/g++.dg/gomp/target-lambda-1.C
+++ /dev/null
@@ -1,94 +0,0 @@
-// We use 'auto' without a function return type, so specify dialect here
-// { dg-additional-options "-std=c++14 -fdump-tree-gimple" }
-#include <cstdlib>
-#include <cstring>
-
-template <typename L>
-void
-omp_target_loop (int begin, int end, L loop)
-{
-  #pragma omp target teams distribute parallel for
-  for (int i = begin; i < end; i++)
-    loop (i);
-}
-
-struct S
-{
-  int a, len;
-  int *ptr;
-
-  auto merge_data_func (int *iptr, int &b)
-  {
-    auto fn = [=](void) -> bool
-      {
-	bool mapped;
-	#pragma omp target map(from:mapped)
-	{
-	  mapped = (ptr != NULL && iptr != NULL);
-	  if (mapped)
-	    {
-	      for (int i = 0; i < len; i++)
-		ptr[i] += a + b + iptr[i];
-	    }
-	}
-	return mapped;
-      };
-    return fn;
-  }
-};
-
-int x = 1;
-
-int main (void)
-{
-  const int N = 10;
-  int *data1 = new int[N];
-  int *data2 = new int[N];
-  memset (data1, 0xab, sizeof (int) * N);
-  memset (data1, 0xcd, sizeof (int) * N);
-
-  int val = 1;
-  int &valref = val;
-  #pragma omp target enter data map(alloc: data1[:N], data2[:N])
-
-  omp_target_loop (0, N, [=](int i) { data1[i] = val; });
-  omp_target_loop (0, N, [=](int i) { data2[i] = valref + 1; });
-
-  #pragma omp target update from(data1[:N], data2[:N])
-
-  for (int i = 0; i < N; i++)
-    {
-      if (data1[i] != 1) abort ();
-      if (data2[i] != 2) abort ();
-    }
-
-  #pragma omp target exit data map(delete: data1[:N], data2[:N])
-
-  int b = 8;
-  S s = { 4, N, data1 };
-  auto f = s.merge_data_func (data2, b);
-
-  if (f ()) abort ();
-
-  #pragma omp target enter data map(to: data1[:N])
-  if (f ()) abort ();
-
-  #pragma omp target enter data map(to: data2[:N])
-  if (!f ()) abort ();
-
-  #pragma omp target exit data map(from: data1[:N], data2[:N])
-
-  for (int i = 0; i < N; i++)
-    {
-      if (data1[i] != 0xf) abort ();
-      if (data2[i] != 2) abort ();
-    }
-
-  return 0;
-}
-
-/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(b\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:iptr \[pointer assign, bias: 0\]\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\) map\(to:\*__closure \[len: [0-9]+\]\) map\(firstprivate:__closure \[pointer assign, bias: 0\]\) map\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\) map\(always_pointer:__closure->__this \[pointer assign, bias: 0\]\) map\(from:mapped \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) map\(attach_zero_length_array_section:__closure->__iptr \[bias: 0\]\) map\(attach_zero_length_array_section:_[0-9]+->ptr \[bias: 0\]\)} "gimple" } } */
-
-/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(end\) firstprivate\(begin\) map\(to:loop \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) map\(attach_zero_length_array_section:loop\.__data1 \[bias: 0\]\)} "gimple" } } */
-
-/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(end\) firstprivate\(begin\) map\(to:loop \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) map\(attach_zero_length_array_section:loop\.__data2 \[bias: 0\]\)} "gimple" } } */