diff options
| author | Cesar Philippidis <cesar@codesourcery.com> | 2018-11-30 12:39:49 -0800 |
|---|---|---|
| committer | Thomas Schwinge <tschwinge@gcc.gnu.org> | 2018-11-30 21:39:49 +0100 |
| commit | fe570ff8d4347ea98108da3cf0d4f3800294d5c5 (patch) | |
| tree | bed3444f5ac3e46f1889acb012e868bac81da4ef | |
| parent | 344b0fdf2ee426f81cf8ceafa069c42cfdc9f5fc (diff) | |
| download | gcc-fe570ff8d4347ea98108da3cf0d4f3800294d5c5.zip gcc-fe570ff8d4347ea98108da3cf0d4f3800294d5c5.tar.gz gcc-fe570ff8d4347ea98108da3cf0d4f3800294d5c5.tar.bz2 | |
[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 <thomas@codesourcery.com>
From-SVN: r266688
| -rw-r--r-- | libgomp/ChangeLog | 6 | ||||
| -rw-r--r-- | libgomp/oacc-parallel.c | 3 | ||||
| -rw-r--r-- | libgomp/testsuite/libgomp.oacc-c-c++-common/pr88288.c | 41 |
3 files changed, 49 insertions, 1 deletions
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 <cesar@codesourcery.com> + + 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 <thomas@codesourcery.com> * 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 <assert.h> + +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; +} |