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