mirror of git://gcc.gnu.org/git/gcc.git
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 4f82d5a95a)
This commit is contained in:
Jakub Jelinek
Tobias Burnus
parent
0e7bc3eaa3
commit
2d20f69092
|
|
@ -1,3 +1,14 @@
|
||||||
|
2024-03-04 Tobias Burnus <tburnus@baylibre.com>
|
||||||
|
|
||||||
|
Backported from master:
|
||||||
|
2024-03-01 Jakub Jelinek <jakub@redhat.com>
|
||||||
|
Tobias Burnus <tburnus@baylibre.com>
|
||||||
|
|
||||||
|
PR c++/110347
|
||||||
|
* gimplify.cc (omp_notice_variable): Fix 'shared' arg to
|
||||||
|
lang_hooks.decls.omp_disregard_value_expr for
|
||||||
|
(first)private in target regions.
|
||||||
|
|
||||||
2023-11-28 Andrew Stubbs <ams@codesourcery.com>
|
2023-11-28 Andrew Stubbs <ams@codesourcery.com>
|
||||||
|
|
||||||
Backport from mainline:
|
Backport from mainline:
|
||||||
|
|
|
||||||
|
|
@ -8120,13 +8120,6 @@ omp_notice_variable (struct gimplify_omp_ctx *ctx, tree decl, bool in_code)
|
||||||
n = splay_tree_lookup (ctx->variables, (splay_tree_key)decl);
|
n = splay_tree_lookup (ctx->variables, (splay_tree_key)decl);
|
||||||
if ((ctx->region_type & ORT_TARGET) != 0)
|
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)
|
if (n == NULL)
|
||||||
{
|
{
|
||||||
unsigned nflags = flags;
|
unsigned nflags = flags;
|
||||||
|
|
@ -8252,9 +8245,22 @@ omp_notice_variable (struct gimplify_omp_ctx *ctx, tree decl, bool in_code)
|
||||||
}
|
}
|
||||||
found_outer:
|
found_outer:
|
||||||
omp_add_variable (ctx, decl, nflags);
|
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
|
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 nothing changed, there's nothing left to do. */
|
||||||
if ((n->value & flags) == flags)
|
if ((n->value & flags) == flags)
|
||||||
return ret;
|
return ret;
|
||||||
|
|
|
||||||
|
|
@ -1,3 +1,13 @@
|
||||||
|
2024-03-04 Tobias Burnus <tburnus@baylibre.com>
|
||||||
|
|
||||||
|
Backported from master:
|
||||||
|
2024-03-01 Jakub Jelinek <jakub@redhat.com>
|
||||||
|
Tobias Burnus <tburnus@baylibre.com>
|
||||||
|
|
||||||
|
PR c++/110347
|
||||||
|
* g++.dg/gomp/target-lambda-1.C: Moved to become a
|
||||||
|
run-time test under testsuite/libgomp.c++.
|
||||||
|
|
||||||
2023-11-28 Andrew Stubbs <ams@codesourcery.com>
|
2023-11-28 Andrew Stubbs <ams@codesourcery.com>
|
||||||
|
|
||||||
Backport from mainline:
|
Backport from mainline:
|
||||||
|
|
|
||||||
|
|
@ -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" } } */
|
|
||||||
|
|
@ -1,3 +1,21 @@
|
||||||
|
2024-03-04 Tobias Burnus <tburnus@baylibre.com>
|
||||||
|
|
||||||
|
Backported from master:
|
||||||
|
2024-03-01 Jakub Jelinek <jakub@redhat.com>
|
||||||
|
Tobias Burnus <tburnus@baylibre.com>
|
||||||
|
|
||||||
|
PR c++/110347
|
||||||
|
* 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.
|
||||||
|
|
||||||
2023-11-28 Andrew Stubbs <ams@codesourcery.com>
|
2023-11-28 Andrew Stubbs <ams@codesourcery.com>
|
||||||
|
|
||||||
Backport from mainline:
|
Backport from mainline:
|
||||||
|
|
|
||||||
|
|
@ -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;
|
||||||
|
}
|
||||||
|
|
@ -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 ();
|
||||||
|
}
|
||||||
|
|
@ -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;
|
||||||
|
}
|
||||||
|
|
@ -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 ();
|
||||||
|
}
|
||||||
|
|
@ -1,4 +1,4 @@
|
||||||
// { dg-do run { target offload_device_nonshared_as } }
|
// { dg-do run }
|
||||||
|
|
||||||
#include <cstdlib>
|
#include <cstdlib>
|
||||||
#include <cstring>
|
#include <cstring>
|
||||||
|
|
@ -48,7 +48,11 @@ int main (void)
|
||||||
int *data1 = new int[N];
|
int *data1 = new int[N];
|
||||||
int *data2 = new int[N];
|
int *data2 = new int[N];
|
||||||
memset (data1, 0xab, sizeof (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 val = 1;
|
||||||
int &valref = val;
|
int &valref = val;
|
||||||
|
|
@ -77,13 +81,16 @@ int main (void)
|
||||||
if (f ()) abort ();
|
if (f ()) abort ();
|
||||||
|
|
||||||
#pragma omp target enter data map(to: data2[:N])
|
#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])
|
#pragma omp target exit data map(from: data1[:N], data2[:N])
|
||||||
|
|
||||||
|
if (!shared_mem)
|
||||||
for (int i = 0; i < N; i++)
|
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 ();
|
if (data2[i] != 2) abort ();
|
||||||
}
|
}
|
||||||
|
|
||||||
|
|
|
||||||
|
|
@ -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" } } */
|
||||||
|
|
@ -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 ();
|
||||||
|
}
|
||||||
|
|
@ -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 ();
|
||||||
|
}
|
||||||
Loading…
Reference in New Issue