Jakub Jelinek wrote:
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.
...
I have now added Jakub's updated the gimplify.cc patch, renamed the test
files, added the proposed lambda test case as well, did add a missing
line break, and updated the target-lambda-1.C to also work with shared
memory.
I think the patch should be good, having testing it with offloading here
and Jakub also testing it on his side.
Final comments, suggestions, remarks?
Tobias
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 <ja...@redhat.com>
Tobias Burnus <tbur...@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 <tbur...@baylibre.com>
gcc/gimplify.cc | 20 +-
gcc/testsuite/g++.dg/gomp/target-lambda-1.C | 94 -------
libgomp/testsuite/libgomp.c++/firstprivate-1.C | 305 +++++++++++++++++++++++
libgomp/testsuite/libgomp.c++/firstprivate-2.C | 125 ++++++++++
libgomp/testsuite/libgomp.c++/private-1.C | 247 ++++++++++++++++++
libgomp/testsuite/libgomp.c++/private-2.C | 117 +++++++++
libgomp/testsuite/libgomp.c++/target-lambda-1.C | 15 +-
libgomp/testsuite/libgomp.c++/target-lambda-3.C | 104 ++++++++
libgomp/testsuite/libgomp.c++/target-lambda-4.C | 41 +++
libgomp/testsuite/libgomp.c++/use_device_ptr-1.C | 126 ++++++++++
10 files changed, 1089 insertions(+), 105 deletions(-)
diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc
index 7f79b3cc7e6..6ebca964cb2 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 5ce8ceadb19..00000000000
--- 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 <cstdlib>
-#include <cstring>
-
-template <typename L>
-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" } } */
diff --git a/libgomp/testsuite/libgomp.c++/firstprivate-1.C b/libgomp/testsuite/libgomp.c++/firstprivate-1.C
new file mode 100644
index 00000000000..ae5d4fbe1bf
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/firstprivate-1.C
@@ -0,0 +1,305 @@
+/* PR c++/110347 */
+
+#include <omp.h>
+#include <stdint.h>
+#include <stdlib.h>
+
+struct S {
+ int A, B[10], *C;
+ void f (int dev);
+ void g (int dev);
+};
+
+template<typename T>
+struct St {
+ T A, B[10], *C;
+ void ft (int dev);
+ void gt (int dev);
+};
+
+
+void
+S::f (int dev)
+{
+ A = 5;
+ C = (int *) malloc (sizeof (int) * 10);
+ uintptr_t c_saved = (uintptr_t) C;
+ for (int i = 0; i < 10; i++)
+ B[i] = C[i] = i+5;
+
+ #pragma omp target firstprivate(A) firstprivate(B) firstprivate(C) \
+ firstprivate(c_saved) device(dev)
+ {
+ if (A != 5)
+ abort ();
+ for (int i = 0; i < 10; i++)
+ if (B[i] != i + 5)
+ abort ();
+ if (c_saved != (uintptr_t) C)
+ abort ();
+ A = 99;
+ for (int i = 0; i < 10; i++)
+ B[i] = -i-23;
+ C = &A;
+ }
+
+ if (A != 5)
+ abort ();
+ if (c_saved != (uintptr_t) C)
+ abort ();
+ for (int i = 0; i < 10; i++)
+ if (B[i] != i + 5 || C[i] != i+5)
+ abort ();
+
+ #pragma omp parallel if (0) firstprivate(A) firstprivate(B) firstprivate(C)
+ {
+ if (A != 5)
+ abort ();
+ for (int i = 0; i < 10; i++)
+ if (B[i] != i + 5)
+ abort ();
+ if (c_saved != (uintptr_t) C)
+ abort ();
+ A = 99;
+ for (int i = 0; i < 10; i++)
+ B[i] = -i-23;
+ C = &A;
+ }
+
+ if (A != 5)
+ abort ();
+ if (c_saved != (uintptr_t) C)
+ abort ();
+ for (int i = 0; i < 10; i++)
+ if (B[i] != i + 5 || C[i] != i+5)
+ abort ();
+
+ free (C);
+}
+
+void
+S::g (int dev)
+{
+ A = 5;
+ C = (int *) malloc (sizeof (int) * 10);
+ uintptr_t c_saved = (uintptr_t) C;
+ for (int i = 0; i < 10; i++)
+ B[i] = C[i] = i+5;
+
+ #pragma omp target firstprivate(A) firstprivate(B) firstprivate(C) \
+ allocate(allocator(omp_low_lat_mem_alloc), align(128): A, B, C) \
+ device(dev)
+ {
+#if 0 /* FIXME: The following is disabled because of PR middle-end/113436. */
+ if (((uintptr_t) &A) % 128 != 0)
+ abort ();
+ if (((uintptr_t) &B) % 128 != 0)
+ abort ();
+ if (((uintptr_t) &C) % 128 != 0)
+ abort ();
+#endif
+ if (A != 5)
+ abort ();
+ for (int i = 0; i < 10; i++)
+ if (B[i] != i + 5)
+ abort ();
+ if (c_saved != (uintptr_t) C)
+ abort ();
+ A = 99;
+ for (int i = 0; i < 10; i++)
+ B[i] = -i-23;
+ C = &A;
+ }
+
+ if (A != 5)
+ abort ();
+ if (c_saved != (uintptr_t) C)
+ abort ();
+ for (int i = 0; i < 10; i++)
+ if (B[i] != i + 5 || C[i] != i+5)
+ abort ();
+
+ #pragma omp parallel if (0) firstprivate(A) firstprivate(B) firstprivate(C) \
+ allocate(allocator(omp_low_lat_mem_alloc), align(128): A, B, C)
+ {
+ if (A != 5)
+ abort ();
+ for (int i = 0; i < 10; i++)
+ if (B[i] != i + 5)
+ abort ();
+ if (c_saved != (uintptr_t) C)
+ abort ();
+ if (((uintptr_t) &A) % 128 != 0)
+ abort ();
+ if (((uintptr_t) &B) % 128 != 0)
+ abort ();
+ if (((uintptr_t) &C) % 128 != 0)
+ abort ();
+ A = 99;
+ for (int i = 0; i < 10; i++)
+ B[i] = -i-23;
+ C = &A;
+ }
+
+ if (A != 5)
+ abort ();
+ if (c_saved != (uintptr_t) C)
+ abort ();
+ for (int i = 0; i < 10; i++)
+ if (B[i] != i + 5 || C[i] != i+5)
+ abort ();
+
+ free (C);
+}
+
+
+template<typename T>
+void
+St<T>::ft (int dev)
+{
+ A = 5;
+ C = (T *) malloc (sizeof (T) * 10);
+ uintptr_t c_saved = (uintptr_t) C;
+ for (int i = 0; i < 10; i++)
+ B[i] = C[i] = i+5;
+
+ #pragma omp target firstprivate(A) firstprivate(B) firstprivate(C) \
+ firstprivate(c_saved) device(dev)
+ {
+ if (A != 5)
+ abort ();
+ for (int i = 0; i < 10; i++)
+ if (B[i] != i + 5)
+ abort ();
+ if (c_saved != (uintptr_t) C)
+ abort ();
+ A = 99;
+ for (int i = 0; i < 10; i++)
+ B[i] = -i-23;
+ C = &A;
+ }
+
+ if (A != 5)
+ abort ();
+ if (c_saved != (uintptr_t) C)
+ abort ();
+ for (int i = 0; i < 10; i++)
+ if (B[i] != i + 5 || C[i] != i+5)
+ abort ();
+
+ #pragma omp parallel if (0) firstprivate(A) firstprivate(B) firstprivate(C)
+ {
+ if (A != 5)
+ abort ();
+ for (int i = 0; i < 10; i++)
+ if (B[i] != i + 5)
+ abort ();
+ if (c_saved != (uintptr_t) C)
+ abort ();
+ A = 99;
+ for (int i = 0; i < 10; i++)
+ B[i] = -i-23;
+ C = &A;
+ }
+
+ if (A != 5)
+ abort ();
+ if (c_saved != (uintptr_t) C)
+ abort ();
+ for (int i = 0; i < 10; i++)
+ if (B[i] != i + 5 || C[i] != i+5)
+ abort ();
+
+ free (C);
+}
+
+template<typename T>
+void
+St<T>::gt (int dev)
+{
+ A = 5;
+ C = (T *) malloc (sizeof (T) * 10);
+ uintptr_t c_saved = (uintptr_t) C;
+ for (int i = 0; i < 10; i++)
+ B[i] = C[i] = i+5;
+
+ #pragma omp target firstprivate(A) firstprivate(B) firstprivate(C) \
+ allocate(allocator(omp_low_lat_mem_alloc), align(128): A, B, C) \
+ device(dev)
+ {
+#if 0 /* FIXME: The following is disabled because of PR middle-end/113436. */
+ if (((uintptr_t) &A) % 128 != 0)
+ abort ();
+ if (((uintptr_t) &B) % 128 != 0)
+ abort ();
+ if (((uintptr_t) &C) % 128 != 0)
+ abort ();
+#endif
+ if (A != 5)
+ abort ();
+ for (int i = 0; i < 10; i++)
+ if (B[i] != i + 5)
+ abort ();
+ if (c_saved != (uintptr_t) C)
+ abort ();
+ A = 99;
+ for (int i = 0; i < 10; i++)
+ B[i] = -i-23;
+ C = &A;
+ }
+
+ if (A != 5)
+ abort ();
+ if (c_saved != (uintptr_t) C)
+ abort ();
+ for (int i = 0; i < 10; i++)
+ if (B[i] != i + 5 || C[i] != i+5)
+ abort ();
+
+ #pragma omp parallel if (0) firstprivate(A) firstprivate(B) firstprivate(C) \
+ allocate(allocator(omp_low_lat_mem_alloc), align(128): A, B, C)
+ {
+ if (A != 5)
+ abort ();
+ for (int i = 0; i < 10; i++)
+ if (B[i] != i + 5)
+ abort ();
+ if (c_saved != (uintptr_t) C)
+ abort ();
+ if (((uintptr_t) &A) % 128 != 0)
+ abort ();
+ if (((uintptr_t) &B) % 128 != 0)
+ abort ();
+ if (((uintptr_t) &C) % 128 != 0)
+ abort ();
+ A = 99;
+ for (int i = 0; i < 10; i++)
+ B[i] = -i-23;
+ C = &A;
+ }
+
+ if (A != 5)
+ abort ();
+ if (c_saved != (uintptr_t) C)
+ abort ();
+ for (int i = 0; i < 10; i++)
+ if (B[i] != i + 5 || C[i] != i+5)
+ abort ();
+
+ free (C);
+}
+
+int
+main ()
+{
+ struct S s;
+ struct St<int> st;
+ for (int dev = 0; dev <= omp_get_num_devices(); dev++)
+ {
+ s.f (dev);
+ st.ft (dev);
+ s.g (dev);
+ st.gt (dev);
+ }
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c++/firstprivate-2.C b/libgomp/testsuite/libgomp.c++/firstprivate-2.C
new file mode 100644
index 00000000000..a4f2514b591
--- /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 ();
+}
diff --git a/libgomp/testsuite/libgomp.c++/private-1.C b/libgomp/testsuite/libgomp.c++/private-1.C
new file mode 100644
index 00000000000..19ee726a222
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/private-1.C
@@ -0,0 +1,247 @@
+/* PR c++/110347 */
+
+#include <omp.h>
+#include <stdint.h>
+#include <stdlib.h>
+
+struct S {
+ int A, B[10], *C;
+ void f (int dev);
+ void g (int dev);
+};
+
+template<typename T>
+struct St {
+ T A, B[10], *C;
+ void ft (int dev);
+ void gt (int dev);
+};
+
+
+void
+S::f (int dev)
+{
+ A = 5;
+ C = (int *) malloc (sizeof (int) * 10);
+ uintptr_t c_saved = (uintptr_t) C;
+ for (int i = 0; i < 10; i++)
+ B[i] = C[i] = i+5;
+
+ #pragma omp target private(A) private(B) private(C) device(dev)
+ {
+ A = 99;
+ for (int i = 0; i < 10; i++)
+ B[i] = -i-23;
+ C = &A;
+ }
+
+ if (A != 5)
+ abort ();
+ if (c_saved != (uintptr_t) C)
+ abort ();
+ for (int i = 0; i < 10; i++)
+ if (B[i] != i + 5 || C[i] != i+5)
+ abort ();
+
+ #pragma omp parallel if (0) private(A) private(B) private(C)
+ {
+ A = 99;
+ for (int i = 0; i < 10; i++)
+ B[i] = -i-23;
+ C = &A;
+ }
+
+ if (A != 5)
+ abort ();
+ if (c_saved != (uintptr_t) C)
+ abort ();
+ for (int i = 0; i < 10; i++)
+ if (B[i] != i + 5 || C[i] != i+5)
+ abort ();
+
+ free (C);
+}
+
+void
+S::g (int dev)
+{
+ A = 5;
+ C = (int *) malloc (sizeof (int) * 10);
+ uintptr_t c_saved = (uintptr_t) C;
+ for (int i = 0; i < 10; i++)
+ B[i] = C[i] = i+5;
+
+ #pragma omp target private(A) private(B) private(C) \
+ allocate(allocator(omp_low_lat_mem_alloc), align(128): A, B, C) \
+ device(dev)
+ {
+#if 0 /* FIXME: The following is disabled because of PR middle-end/113436. */
+ if (((uintptr_t) &A) % 128 != 0)
+ abort ();
+ if (((uintptr_t) &B) % 128 != 0)
+ abort ();
+ if (((uintptr_t) &C) % 128 != 0)
+ abort ();
+#endif
+ A = 99;
+ for (int i = 0; i < 10; i++)
+ B[i] = -i-23;
+ C = &A;
+ }
+
+ if (A != 5)
+ abort ();
+ if (c_saved != (uintptr_t) C)
+ abort ();
+ for (int i = 0; i < 10; i++)
+ if (B[i] != i + 5 || C[i] != i+5)
+ abort ();
+
+ #pragma omp parallel if (0) private(A) private(B) private(C) \
+ allocate(allocator(omp_low_lat_mem_alloc), align(128): A, B, C)
+ {
+ if (((uintptr_t) &A) % 128 != 0)
+ abort ();
+ if (((uintptr_t) &B) % 128 != 0)
+ abort ();
+ if (((uintptr_t) &C) % 128 != 0)
+ abort ();
+ A = 99;
+ for (int i = 0; i < 10; i++)
+ B[i] = -i-23;
+ C = &A;
+ }
+
+ if (A != 5)
+ abort ();
+ if (c_saved != (uintptr_t) C)
+ abort ();
+ for (int i = 0; i < 10; i++)
+ if (B[i] != i + 5 || C[i] != i+5)
+ abort ();
+
+ free (C);
+}
+
+
+template<typename T>
+void
+St<T>::ft (int dev)
+{
+ A = 5;
+ C = (T *) malloc (sizeof (T) * 10);
+ uintptr_t c_saved = (uintptr_t) C;
+ for (int i = 0; i < 10; i++)
+ B[i] = C[i] = i+5;
+
+ #pragma omp target private(A) private(B) private(C) device(dev)
+ {
+ A = 99;
+ for (int i = 0; i < 10; i++)
+ B[i] = -i-23;
+ C = &A;
+ }
+
+ if (A != 5)
+ abort ();
+ if (c_saved != (uintptr_t) C)
+ abort ();
+ for (int i = 0; i < 10; i++)
+ if (B[i] != i + 5 || C[i] != i+5)
+ abort ();
+
+ #pragma omp parallel if (0) private(A) private(B) private(C)
+ {
+ A = 99;
+ for (int i = 0; i < 10; i++)
+ B[i] = -i-23;
+ C = &A;
+ }
+
+ if (A != 5)
+ abort ();
+ if (c_saved != (uintptr_t) C)
+ abort ();
+ for (int i = 0; i < 10; i++)
+ if (B[i] != i + 5 || C[i] != i+5)
+ abort ();
+
+ free (C);
+}
+
+template<typename T>
+void
+St<T>::gt (int dev)
+{
+ A = 5;
+ C = (T *) malloc (sizeof (T) * 10);
+ uintptr_t c_saved = (uintptr_t) C;
+ for (int i = 0; i < 10; i++)
+ B[i] = C[i] = i+5;
+
+ #pragma omp target private(A) private(B) private(C) \
+ allocate(allocator(omp_low_lat_mem_alloc), align(128): A, B, C) \
+ device(dev)
+ {
+#if 0 /* FIXME: The following is disabled because of PR middle-end/113436. */
+ if (((uintptr_t) &A) % 128 != 0)
+ abort ();
+ if (((uintptr_t) &B) % 128 != 0)
+ abort ();
+ if (((uintptr_t) &C) % 128 != 0)
+ abort ();
+#endif
+ A = 99;
+ for (int i = 0; i < 10; i++)
+ B[i] = -i-23;
+ C = &A;
+ }
+
+ if (A != 5)
+ abort ();
+ if (c_saved != (uintptr_t) C)
+ abort ();
+ for (int i = 0; i < 10; i++)
+ if (B[i] != i + 5 || C[i] != i+5)
+ abort ();
+
+ #pragma omp parallel if (0) private(A) private(B) private(C) \
+ allocate(allocator(omp_low_lat_mem_alloc), align(128): A, B, C)
+ {
+ if (((uintptr_t) &A) % 128 != 0)
+ abort ();
+ if (((uintptr_t) &B) % 128 != 0)
+ abort ();
+ if (((uintptr_t) &C) % 128 != 0)
+ abort ();
+ A = 99;
+ for (int i = 0; i < 10; i++)
+ B[i] = -i-23;
+ C = &A;
+ }
+
+ if (A != 5)
+ abort ();
+ if (c_saved != (uintptr_t) C)
+ abort ();
+ for (int i = 0; i < 10; i++)
+ if (B[i] != i + 5 || C[i] != i+5)
+ abort ();
+
+ free (C);
+}
+
+int
+main ()
+{
+ struct S s;
+ struct St<int> st;
+ for (int dev = 0; dev <= omp_get_num_devices(); dev++)
+ {
+ s.f (dev);
+ st.ft (dev);
+ s.g (dev);
+ st.gt (dev);
+ }
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c++/private-2.C b/libgomp/testsuite/libgomp.c++/private-2.C
new file mode 100644
index 00000000000..aa472cb62ee
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/private-2.C
@@ -0,0 +1,117 @@
+/* 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 private(A) if(0) shared(B) default(none)
+ {
+ 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 private(A)if(0) shared(B) default(none)
+ {
+ 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 private(A) map(from:B) device(dev)
+ {
+ 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 private(A) map(from:B) device(dev)
+ {
+ 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 private(C) if(0) shared(D) default(none)
+ {
+ 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 private(C)if(0) shared(D) default(none)
+ {
+ 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 private(C) map(from:D) defaultmap(none) device(dev)
+ {
+ 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 private(C) map(from:D) defaultmap(none) device(dev)
+ {
+ 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 ();
+}
diff --git a/libgomp/testsuite/libgomp.c++/target-lambda-1.C b/libgomp/testsuite/libgomp.c++/target-lambda-1.C
index fa882d09800..6eb0d0bb1db 100644
--- a/libgomp/testsuite/libgomp.c++/target-lambda-1.C
+++ b/libgomp/testsuite/libgomp.c++/target-lambda-1.C
@@ -1,4 +1,4 @@
-// { dg-do run { target offload_device_nonshared_as } }
+// { dg-do run }
#include <cstdlib>
#include <cstring>
@@ -48,7 +48,11 @@ int main (void)
int *data1 = new int[N];
int *data2 = new int[N];
memset (data1, 0xab, sizeof (int) * N);
- memset (data1, 0xcd, sizeof (int) * N);
+ memset (data2, 0xcd, sizeof (int) * N);
+
+ bool shared_mem = false;
+ #pragma omp target map(to: shared_mem)
+ shared_mem = true;
int val = 1;
int &valref = val;
@@ -77,13 +81,16 @@ int main (void)
if (f ()) abort ();
#pragma omp target enter data map(to: data2[:N])
- if (!f ()) abort ();
+ if (!f () && !shared_mem) abort ();
#pragma omp target exit data map(from: data1[:N], data2[:N])
+ if (!shared_mem)
for (int i = 0; i < N; i++)
{
- if (data1[i] != 0xf) abort ();
+ /* With shared memory, data1 is not modified inside 'f'
+ as mapped = false. */
+ if (!shared_mem && data1[i] != 0xf) abort ();
if (data2[i] != 2) abort ();
}
diff --git a/libgomp/testsuite/libgomp.c++/target-lambda-3.C b/libgomp/testsuite/libgomp.c++/target-lambda-3.C
new file mode 100644
index 00000000000..6be8426bd3e
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/target-lambda-3.C
@@ -0,0 +1,104 @@
+// We use 'auto' without a function return type, so specify dialect here
+// { dg-additional-options "-std=c++14 -fdump-tree-gimple" }
+#include <cstdlib>
+#include <cstring>
+#include <omp.h>
+
+template <typename L>
+void
+omp_target_loop (int begin, int end, L loop, int dev)
+{
+ #pragma omp target teams distribute parallel for device(dev)
+ for (int i = begin; i < end; i++)
+ loop (i);
+}
+
+struct S
+{
+ int a, len;
+ int *ptr;
+
+ auto merge_data_func (int *iptr, int &b, int dev)
+ {
+ auto fn = [=](void) -> bool
+ {
+ bool mapped = (omp_target_is_present (iptr, dev)
+ && omp_target_is_present (ptr, dev));
+ #pragma omp target device(dev)
+ {
+ if (mapped)
+ {
+ for (int i = 0; i < len; i++)
+ ptr[i] += a + b + iptr[i];
+ }
+ }
+ return mapped;
+ };
+ return fn;
+ }
+};
+
+int x = 1;
+
+void run (int dev)
+{
+ const int N = 10;
+ int *data1 = new int[N];
+ int *data2 = new int[N];
+ memset (data1, 0xab, sizeof (int) * N);
+ memset (data2, 0xcd, sizeof (int) * N);
+
+ bool shared_mem = (omp_target_is_present (data1, dev)
+ && omp_target_is_present (data2, dev));
+ int val = 1;
+ int &valref = val;
+ #pragma omp target enter data map(alloc: data1[:N], data2[:N]) device(dev)
+
+ omp_target_loop (0, N, [=](int i) { data1[i] = val; }, dev);
+ omp_target_loop (0, N, [=](int i) { data2[i] = valref + 1; }, dev);
+
+ #pragma omp target update from(data1[:N], data2[:N]) device(dev)
+
+ 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]) device(dev)
+
+ int b = 8;
+ S s = { 4, N, data1 };
+ auto f = s.merge_data_func (data2, b, dev);
+ if (f () ^ shared_mem) abort ();
+
+ #pragma omp target enter data map(to: data1[:N]) device(dev)
+ if (f () ^ shared_mem) abort ();
+
+ #pragma omp target enter data map(to: data2[:N]) device(dev)
+ if (!f ()) abort ();
+
+ #pragma omp target exit data map(from: data1[:N], data2[:N]) device(dev)
+
+ for (int i = 0; i < N; i++)
+ {
+ if ((!shared_mem && data1[i] != 0xf)
+ || (shared_mem && data1[i] != 0x2b))
+ abort ();
+ if (data2[i] != 2) abort ();
+ }
+ delete [] data1;
+ delete [] data2;
+}
+
+int main ()
+{
+ for (int dev = 0; dev <= omp_get_num_devices (); dev++)
+ run (dev);
+}
+
+/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(b\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:iptr \[pointer assign, bias: 0\]\) firstprivate\(mapped\) 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\(alloc:\*_[0-9]+ \[len: 0\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) device\(_[0-9]+\) 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\]\) device\(dev.[0-9_]+\) 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\]\) device\(dev.[0-9_]+\) map\(attach_zero_length_array_section:loop\.__data2 \[bias: 0\]\)} "gimple" } } */
diff --git a/libgomp/testsuite/libgomp.c++/target-lambda-4.C b/libgomp/testsuite/libgomp.c++/target-lambda-4.C
new file mode 100644
index 00000000000..4830cbce523
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/target-lambda-4.C
@@ -0,0 +1,41 @@
+int
+foo ()
+{
+ int var = 42;
+ [&var] () {
+#pragma omp target firstprivate(var)
+ {
+ var += 26;
+ if (var != 42 + 26)
+ __builtin_abort ();
+ }
+ } ();
+ return var;
+}
+
+
+template <typename T>
+struct A {
+ A () : a(), b()
+ {
+ [&] ()
+ {
+#pragma omp target firstprivate (a) map (from: b)
+ b = ++a;
+ } ();
+ }
+
+ T a, b;
+};
+
+
+int
+main ()
+{
+ if (foo () != 42)
+ __builtin_abort ();
+
+ A<int> x;
+ if (x.a != 0 || x.b != 1)
+ __builtin_abort ();
+}
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 00000000000..bc3cc8f3da2
--- /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 ();
+}
