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, 125 insertions, 0 deletions

diff --git a/libgomp/testsuite/libgomp.c++/firstprivate-2.C b/libgomp/testsuite/libgomp.c++/firstprivate-2.C
new file mode 100644
index 0000000..a4f2514
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/firstprivate-2.C
@@ -0,0 +1,125 @@
+/* PR c++/110347 */
+
+#include <omp.h>
+
+struct t {
+  int A;
+  void f (int dev);
+};
+
+void
+t::f (int dev)
+{
+  int B = 49;
+
+  A = 7;
+  #pragma omp parallel firstprivate(A) if(0) shared(B) default(none)
+  {
+    if (A != 7) { __builtin_printf("ERROR 1b: %d (!= 7) inside omp parallel\n", A); __builtin_abort (); }
+    A = 5;
+    B = A;
+  }
+  if (A != 7) { __builtin_printf("ERROR 1: %d (!= 7) omp parallel\n", A); __builtin_abort (); }
+  if (B != 5) { __builtin_printf("ERROR 1a: %d\n", B); __builtin_abort (); }
+  A = 8; B = 49;
+  #pragma omp parallel firstprivate(A)if(0) shared(B) default(none)
+  {
+    if (A != 8) { __builtin_printf("ERROR 1b: %d (!= 8) inside omp parallel\n", A); __builtin_abort (); }
+    A = 6;
+    B = A;
+  }
+  if (A != 8) { __builtin_printf("ERROR 2: %d (!= 8) omp parallel\n", A); __builtin_abort (); }
+  if (B != 6) { __builtin_printf("ERROR 2a: %d\n", B); __builtin_abort (); }
+  A = 8; B = 49;
+
+  #pragma omp target firstprivate(A) map(from:B) device(dev)
+  {
+    if (A != 8) { __builtin_printf("ERROR 2b: %d (!= 8) inside omp target\n", A); __builtin_abort (); }
+    A = 7;
+    B = A;
+  }
+  if (A != 8) { __builtin_printf("ERROR 3: %d (!= 8) omp target\n", A); __builtin_abort (); }
+  if (B != 7) { __builtin_printf("ERROR 3a: %d\n", B); __builtin_abort (); }
+  A = 9; B = 49;
+  #pragma omp target firstprivate(A) map(from:B) device(dev)
+  {
+    if (A != 9) { __builtin_printf("ERROR 3b: %d (!= 9) inside omp target\n", A); __builtin_abort (); }
+    A = 8;
+    B = A;
+  }
+  if (A != 9) { __builtin_printf("ERROR 4: %d (!= 9) omp target\n", A); __builtin_abort (); }
+  if (B != 8) { __builtin_printf("ERROR 4a: %d\n", B); __builtin_abort (); }
+}
+
+
+template <typename T>
+struct tt {
+  T C;
+  void g (int dev);
+};
+
+template <typename T>
+void
+tt<T>::g (int dev)
+{
+  T D = 49;
+  C = 7;
+  #pragma omp parallel firstprivate(C) if(0) shared(D) default(none)
+  {
+    if (C != 7) { __builtin_printf("ERROR 1b: %d (!= 7) inside omp parallel\n", C);__builtin_abort (); }
+    C = 5;
+    D = C;
+  }
+  if (C != 7) { __builtin_printf("ERROR 1: %d (!= 7) omp parallel\n", C);__builtin_abort (); }
+  if (D != 5) { __builtin_printf("ERROR 1a: %d\n", D);__builtin_abort (); }
+  C = 8; D = 49;
+  #pragma omp parallel firstprivate(C)if(0) shared(D) default(none)
+  {
+    if (C != 8) { __builtin_printf("ERROR 1b: %d (!= 8) inside omp parallel\n", C);__builtin_abort (); }
+    C = 6;
+    D = C;
+  }
+  if (C != 8) { __builtin_printf("ERROR 2: %d (!= 8) omp parallel\n", C);__builtin_abort (); }
+  if (D != 6) { __builtin_printf("ERROR 2a: %d\n", D);__builtin_abort (); }
+  C = 8; D = 49;
+  #pragma omp target firstprivate(C) map(from:D) defaultmap(none) device(dev)
+  {
+    if (C != 8) { __builtin_printf("ERROR 2b: %d (!= 8) inside omp target\n", C);__builtin_abort (); }
+    C = 7;
+    D = C;
+  }
+  if (C != 8) { __builtin_printf("ERROR 3: %d (!= 8) omp target\n", C);__builtin_abort (); }
+  if (D != 7) { __builtin_printf("ERROR 3a: %d\n", D);__builtin_abort (); }
+  C = 9; D = 49;
+  #pragma omp target firstprivate(C) map(from:D) defaultmap(none) device(dev)
+  {
+    if (C != 9) { __builtin_printf("ERROR 3b: %d (!= 9) inside omp target\n", C);__builtin_abort (); }
+    C = 8;
+    D = C;
+  }
+  if (C != 9) { __builtin_printf("ERROR 4: %d (!= 9) omp target\n", C); __builtin_abort (); }
+  if (D != 8) { __builtin_printf("ERROR 4a: %d\n", D); }
+}
+
+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 ();
+}