diff options
author | Jakub Jelinek <jakub@redhat.com> | 2024-03-04 12:55:27 +0100 |
---|---|---|
committer | Tobias Burnus <tburnus@baylibre.com> | 2024-03-04 12:55:27 +0100 |
commit | 2d20f690921a82ee6db0c2fbac7dd5f13d4a0882 (patch) | |
tree | 1fd804b54953ffb2beb4c17573833d95680a1472 /libgomp/testsuite/libgomp.c++/use_device_ptr-1.C | |
parent | 0e7bc3eaa36b81004b799124d2fe00137401a43b (diff) | |
download | gcc-2d20f690921a82ee6db0c2fbac7dd5f13d4a0882.zip gcc-2d20f690921a82ee6db0c2fbac7dd5f13d4a0882.tar.gz gcc-2d20f690921a82ee6db0c2fbac7dd5f13d4a0882.tar.bz2 |
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 <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)
Diffstat (limited to 'libgomp/testsuite/libgomp.c++/use_device_ptr-1.C')
-rw-r--r-- | libgomp/testsuite/libgomp.c++/use_device_ptr-1.C | 126 |
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 (); +} |