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;
}
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;
+++ /dev/null
-// 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" } } */
--- /dev/null
+/* 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;
+}
--- /dev/null
+/* 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 ();
+}
--- /dev/null
+/* 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;
+}
--- /dev/null
+/* 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 ();
+}
-// { dg-do run { target offload_device_nonshared_as } }
+// { dg-do run }
#include <cstdlib>
#include <cstring>
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;
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 ();
}
--- /dev/null
+// 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" } } */
--- /dev/null
+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 ();
+}
--- /dev/null
+/* 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 ();
+}
projects / thirdparty / gcc.git / commitdiff
