From fe570ff8d4347ea98108da3cf0d4f3800294d5c5 Mon Sep 17 00:00:00 2001 From: Cesar Philippidis Date: Fri, 30 Nov 2018 12:39:49 -0800 Subject: [PATCH] [PR88288, OpenACC, libgomp] Adjust offsets for present data clauses Make libgomp respect the on device offset of subarrays which may arise in present data clauses. libgomp/ PR libgomp/88288 * oacc-parallel.c (GOACC_parallel_keyed): Add offset to devaddrs. * testsuite/libgomp.oacc-c-c++-common/pr88288.c: New test. Reviewed-by: Thomas Schwinge From-SVN: r266688 --- libgomp/ChangeLog | 6 ++++ libgomp/oacc-parallel.c | 3 +- .../testsuite/libgomp.oacc-c-c++-common/pr88288.c | 41 ++++++++++++++++++++++ 3 files changed, 49 insertions(+), 1 deletion(-) create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/pr88288.c diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog index a9dcbd8..d095a19 100644 --- a/libgomp/ChangeLog +++ b/libgomp/ChangeLog @@ -1,3 +1,9 @@ +2018-11-30 Cesar Philippidis + + PR libgomp/88288 + * oacc-parallel.c (GOACC_parallel_keyed): Add offset to devaddrs. + * testsuite/libgomp.oacc-c-c++-common/pr88288.c: New test. + 2018-11-30 Thomas Schwinge * testsuite/libgomp.oacc-fortran/lib-16-2.f90: New file. diff --git a/libgomp/oacc-parallel.c b/libgomp/oacc-parallel.c index b80ace5..1e08af7 100644 --- a/libgomp/oacc-parallel.c +++ b/libgomp/oacc-parallel.c @@ -232,7 +232,8 @@ GOACC_parallel_keyed (int device, void (*fn) (void *), devaddrs = gomp_alloca (sizeof (void *) * mapnum); for (i = 0; i < mapnum; i++) devaddrs[i] = (void *) (tgt->list[i].key->tgt->tgt_start - + tgt->list[i].key->tgt_offset); + + tgt->list[i].key->tgt_offset + + tgt->list[i].offset); acc_dev->openacc.exec_func (tgt_fn, mapnum, hostaddrs, devaddrs, async, dims, tgt); diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr88288.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr88288.c new file mode 100644 index 0000000..d13e335 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr88288.c @@ -0,0 +1,41 @@ +/* Test present data clauses in acc offloaded regions when the + subarray inside the present clause does not have the same base + offset value as the subarray in the enclosing acc data or acc enter + data variable. */ + +#include + +void +offset (int *data, int n) +{ + int i; + +#pragma acc parallel loop present (data[0:n]) + for (i = 0; i < n; i++) + data[i] = n; +} + +int +main () +{ + const int n = 30; + int data[n], i; + + for (i = 0; i < n; i++) + data[i] = -1; + +#pragma acc data copy(data[0:n]) + { + offset (data + 10, 10); + } + + for (i = 0; i < n; i++) + { + if (i < 10 || i >= 20) + assert (data[i] == -1); + else + assert (data[i] == 10); + } + + return 0; +} -- 2.7.4