Re: [Patch] OpenMP/C++: Fix (first)private clause with member variables [PR110347] [was: [RFA/RFC] C++/OpenMP: Supporting (first)private for member variables [PR110347] - or VALUE_EXPR and gimplify]

[Jakub Jelinek <jakub@redhat.com>]

On Sat, Feb 17, 2024 at 12:35:48AM +0100, Tobias Burnus wrote:
> Hence, I now use this code, but also pass a flag to distinguish target
> regions (→ map) from shared usage, assuming that it is needed for the
> latter (otherwise, there wouldn't be that code).
> 
> The issue only showed up for a compile-only testcase, which I have now
> turned into a run-time testcase.
> In order to do so, I had to fix a bogus test for is mapped (or at least
> I think it is bogus) - and for sure it didn't handle shared memory.
> 
> I also modified it such that it iterates over devices. Changes to the dump:
> the 'device' clause had to be added (3x) and for the long line: 'this' and
> 'iptr' swapped the order and 'map(from:mapped)' became
> 'firstprivate(mapped)' due to my changes.
> I appended a patch which only shows the test-case differences as "git diff"
> contains all lines as I move it to libgomp/.
> 
> Comments, remarks, suggestions?

As discussed on IRC, I believe not disregarding the capture proxies in
target regions if they shouldn't be shared is always wrong, but also the
gimplify.cc suggestion was incorrect.

The thing is that at the place where the omp_disregard_value_expr call
is done currently for target region flags is always in_code ? GOVD_SEEN : 0
so by testing flags & anything we actually don't differentiate between
privatized vars and mapped vars.  So, it needs to be moved after we
actually compute the flags, similarly how we do it for non-target.
Now, in the patch I've mentioned on IRC last night I had & GOVD_MAP) != 0
checks, but that breaks e.g. the target-lambda-3.C testcase.  The
problem is that gimplification treats declare target functions as having
an implicit target region around the whole body, GOVD_MAP of course at
that point isn't set for anything and so we treated as privatized and
thus the vanilla trunk to the patched one resulted e.g. in the lambda
body
@@ -82,13 +82,11 @@ void run(int)::<lambda(int)>::operator()
   int * const data2 [value-expr: __closure->__data2];
   const int val [value-expr: __closure->__val];
 
-  _1 = __closure->__val;
-  _2 = __closure->__data2;
-  _3 = (long unsigned int) i;
-  _4 = _3 * 4;
-  _5 = _2 + _4;
-  _6 = _1 + 1;
-  *_5 = _6;
+  _1 = (long unsigned int) i;
+  _2 = _1 * 4;
+  _3 = data2 + _2;
+  _4 = val + 1;
+  *_3 = _4;
 }
changes, which uses uninitialized vars and so overwrites random memory.
The following updated patch checks for non-presence of GOVD_PRIVATE
and GOVD_FIRSTPRIVATE flags rather than presence of GOVD_MAP and worked
on the new testcases from the patch (but haven't tested it further).

> 	* testsuite/libgomp.c++/target-lambda-3.C: Moved from
> 	gcc/testsuite/g++.dg/gomp/ and fixed is-mapped handling.
> 	* testsuite/libgomp.c++/firstprivate-c++-1.C: New test.
> 	* testsuite/libgomp.c++/firstprivate-c++-2.C: New test.
> 	* testsuite/libgomp.c++/private-c++-1.C: New test.
> 	* testsuite/libgomp.c++/private-c++-2.C: New test.
> 	* testsuite/libgomp.c++/use_device_ptr-c++-1.C: New test.

As discussed on IRC, please drop the -c++ infixes from the tests
and renumber if there are existing tests with that name already.
This is in libgomp.c++/ directory, all the tests are C++ in there.
> --- /dev/null
> +++ b/libgomp/testsuite/libgomp.c++/use_device_ptr-c++-1.C
...
> +  omp_target_free (D, dev);}

Please add a newline in between ; and }

The patch below is meant to be used together with the testsuite
updates from your patch, but perhaps we want also some runtime testcase
using
int
foo ()
{
  int var = 42;
  [&var] () {
#pragma omp target firstprivate(var)
    {
      var += 26;
      if (var != 42 + 26)
	__builtin_abort ();
    }
  } ();
  return var;
}

int
main ()
{
  if (foo () != 42)
    __builtin_abort ();
}
and
template <typename T>
struct A {
  A () : a(), b()
  {
    [&] ()
    {
#pragma omp target firstprivate (a) map (from: b)
      b = ++a;
    } ();
  }

  T a, b;
};

int
main ()
{
  A<int> x;
  if (x.a != 0 || x.b != 1)
    __builtin_abort ();
}
or so (unless this is already covered somewhere).

--- gcc/gimplify.cc.jj	2024-02-28 22:24:54.859623016 +0100
+++ gcc/gimplify.cc	2024-02-29 18:03:00.744657060 +0100
@@ -8144,13 +8144,6 @@ omp_notice_variable (struct gimplify_omp
   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
 	    }
 	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;


	Jakub