From 4f82d5a95a244d0aa4f8b2541b47a21bce8a191b Mon Sep 17 00:00:00 2001 From: Jakub Jelinek Date: Fri, 1 Mar 2024 17:26:42 +0100 Subject: 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 Tobias Burnus 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 --- gcc/gimplify.cc | 20 +++--- gcc/testsuite/g++.dg/gomp/target-lambda-1.C | 94 ----------------------------- 2 files changed, 13 insertions(+), 101 deletions(-) delete mode 100644 gcc/testsuite/g++.dg/gomp/target-lambda-1.C (limited to 'gcc') 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 -#include - -template -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" } } */ -- cgit v1.1