8860d2706d
* gimplify.c (omp_add_variable): Use GOVD_PRIVATE | GOVD_EXPLICIT for VLA helper variables on target data even if not GOVD_FIRSTPRIVATE. (gimplify_scan_omp_clauses): For OMP_CLAUSE_USE_DEVICE_* use just GOVD_EXPLICIT flags. (gimplify_omp_workshare): For OMP_TARGET_DATA move all OMP_CLAUSE_USE_DEVICE_* clauses to the end of clauses chain. * omp-low.c (scan_sharing_clauses): For OMP_CLAUSE_USE_DEVICE_* call install_var_field with mask 11 instead of 3. (lower_omp_target): For OMP_CLAUSE_USE_DEVICE_* use pass (splay_tree_key) &DECL_UID (var) to build_sender_ref instead of var. gcc/c/ * c-typeck.c (c_finish_omp_clauses): For C_ORT_OMP OMP_CLAUSE_USE_DEVICE_* clauses use oacc_reduction_head bitmap instead of generic_head to track duplicates. gcc/cp/ * semantics.c (finish_omp_clauses): For C_ORT_OMP OMP_CLAUSE_USE_DEVICE_* clauses use oacc_reduction_head bitmap instead of generic_head to track duplicates. libgomp/ * target.c (gomp_map_vars_internal): For GOMP_MAP_USE_DEVICE_PTR perform the lookup in the first loop only if !not_found_cnt, otherwise perform lookups for it in the second loop guarded with if (not_found_cnt || has_firstprivate). * testsuite/libgomp.c/target-37.c: New test. * testsuite/libgomp.c++/target-22.C: New test. From-SVN: r274206
72 lines
1.4 KiB
C
72 lines
1.4 KiB
C
extern void abort (void);
|
|
struct S { int e, f; };
|
|
|
|
void
|
|
foo (int n)
|
|
{
|
|
int a[4] = { 0, 1, 2, 3 }, b[n], c = 4;
|
|
struct S d = { 5, 6 };
|
|
int *p = a + 1, i, err;
|
|
for (i = 0; i < n; i++)
|
|
b[i] = 9 + i;
|
|
#pragma omp target data use_device_ptr(p) map(from:err) map(to:a)
|
|
#pragma omp target is_device_ptr(p) private(i) map(from:err)
|
|
{
|
|
err = 0;
|
|
for (i = 0; i < 4; i++)
|
|
if (p[i - 1] != i)
|
|
err = 1;
|
|
}
|
|
if (err)
|
|
abort ();
|
|
for (i = 0; i < 4; i++)
|
|
a[i] = 23 + i;
|
|
#pragma omp target data map(to:a) use_device_addr(a) map(from:err)
|
|
#pragma omp target is_device_ptr(a) private(i) map(from:err)
|
|
{
|
|
err = 0;
|
|
for (i = 0; i < 4; i++)
|
|
if (a[i] != 23 + i)
|
|
err = 1;
|
|
}
|
|
if (err)
|
|
abort ();
|
|
#pragma omp target data use_device_addr(b) map(from:err) map(to:b)
|
|
#pragma omp target is_device_ptr(b) private(i) map(from:err)
|
|
{
|
|
err = 0;
|
|
for (i = 0; i < 4; i++)
|
|
if (b[i] != 9 + i)
|
|
err = 1;
|
|
}
|
|
if (err)
|
|
abort ();
|
|
#pragma omp target data map(to:c) use_device_addr(c) map(from:err)
|
|
{
|
|
int *q = &c;
|
|
#pragma omp target is_device_ptr(q) map(from:err)
|
|
{
|
|
err = *q != 4;
|
|
}
|
|
}
|
|
if (err)
|
|
abort ();
|
|
#pragma omp target data use_device_addr(d) map(to:d) map(from:err)
|
|
{
|
|
struct S *r = &d;
|
|
#pragma omp target is_device_ptr(r) map(from:err)
|
|
{
|
|
err = r->e != 5 || r->f != 6;
|
|
}
|
|
}
|
|
if (err)
|
|
abort ();
|
|
}
|
|
|
|
int
|
|
main ()
|
|
{
|
|
foo (9);
|
|
return 0;
|
|
}
|