OpenMP/C++: Fix (first)private clause with member variables [PR110347]devel/omp/gcc-13

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>
(cherry picked from commit 4f82d5a95a244d0aa4f8b2541b47a21bce8a191b)

1 files changed, 126 insertions, 0 deletions

diff --git a/libgomp/testsuite/libgomp.c++/use_device_ptr-1.C b/libgomp/testsuite/libgomp.c++/use_device_ptr-1.C
new file mode 100644
index 0000000..bc3cc8f
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/use_device_ptr-1.C
@@ -0,0 +1,126 @@
+/* PR c++/110347 */
+
+#include <omp.h>
+
+#define N 30
+
+struct t {
+  int *A;
+  void f (int dev);
+};
+
+void
+t::f (int dev)
+{
+  int *ptr;
+  int B[N];
+  for (int i = 0; i < N; i++)
+    B[i] = 1 + i;
+  ptr = A = (int *) omp_target_alloc (sizeof (int) * N, dev);
+  omp_target_memcpy (A, B, sizeof (int) * N, 0, 0, dev, omp_initial_device);
+
+  #pragma omp target is_device_ptr (A) device(dev)
+  {
+    for (int i = 0; i < N; i++)
+      if (A[i] != 1 + i)
+	__builtin_abort ();
+    for (int i = 0; i < N; i++)
+      A[i] = (-2-i)*10;
+    A = (int *) 0x12345;
+  }
+  if (ptr != A)
+    __builtin_abort ();
+
+  #pragma omp target is_device_ptr (A) device(dev)
+  {
+    for (int i = 0; i < N; i++)
+      if (A[i] != (-2-i)*10)
+	__builtin_abort ();
+    for (int i = 0; i < N; i++)
+      A[i] = (3+i)*11;
+    A = (int *) 0x12345;
+  }
+  if (ptr != A)
+    __builtin_abort ();
+
+  int *C = (int *) __builtin_malloc (sizeof(int)*N);
+  omp_target_memcpy (C, A, sizeof (int) * N, 0, 0, omp_initial_device, dev);
+  for (int i = 0; i < N; i++)
+    if (C[i] != (3+i)*11)
+      __builtin_abort ();
+  __builtin_free (C);
+  omp_target_free (A, dev);
+}
+
+template <typename T>
+struct tt {
+  T *D;
+  void g (int dev);
+};
+
+template <typename T>
+void
+tt<T>::g (int dev)
+{
+  T *ptr;
+  T E[N];
+  for (int i = 0; i < N; i++)
+    E[i] = 1 + i;
+  ptr = D = (T *) omp_target_alloc (sizeof (T) * N, dev);
+  omp_target_memcpy (D, E, sizeof (T) * N, 0, 0, dev, omp_initial_device);
+
+  #pragma omp target is_device_ptr (D) device(dev)
+  {
+    for (int i = 0; i < N; i++)
+      if (D[i] != 1 + i)
+	__builtin_abort ();
+    for (int i = 0; i < N; i++)
+      D[i] = (-2-i)*10;
+    D = (T *) 0x12345;
+  }
+  if (ptr != D)
+    __builtin_abort ();
+
+  #pragma omp target is_device_ptr (D) device(dev)
+  {
+    for (int i = 0; i < N; i++)
+      if (D[i] != (-2-i)*10)
+	__builtin_abort ();
+    for (int i = 0; i < N; i++)
+      D[i] = (3+i)*11;
+    D = (T *) 0x12345;
+  }
+  if (ptr != D)
+    __builtin_abort ();
+
+  T *F = (T *) __builtin_malloc (sizeof(T)*N);
+  omp_target_memcpy (F, D, sizeof (T) * N, 0, 0, omp_initial_device, dev);
+  for (int i = 0; i < N; i++)
+    if (F[i] != (3+i)*11)
+      __builtin_abort ();
+  __builtin_free (F);
+  omp_target_free (D, dev);
+}
+
+void
+foo ()
+{
+  struct t x;
+  for (int dev = 0; dev <= omp_get_num_devices (); dev++)
+    x.f (dev);
+}
+
+void
+bar ()
+{
+  struct tt<int> y;
+  for (int dev = 0; dev <= omp_get_num_devices (); dev++)
+    y.g (dev);
+}
+
+int
+main ()
+{
+  foo ();
+  bar ();
+}