diff options
author | Julian Brown <julian@codesourcery.com> | 2019-11-20 17:51:09 +0000 |
---|---|---|
committer | Julian Brown <jules@gcc.gnu.org> | 2019-11-20 17:51:09 +0000 |
commit | e307b05f4377a22811395f6a051d9db864b8785f (patch) | |
tree | 3b7a12b5b4cd395997af29e9da4c99400e961d52 /libgomp | |
parent | 070e3489678380354027144c6b0cef02e7601a37 (diff) | |
download | gcc-e307b05f4377a22811395f6a051d9db864b8785f.zip gcc-e307b05f4377a22811395f6a051d9db864b8785f.tar.gz gcc-e307b05f4377a22811395f6a051d9db864b8785f.tar.bz2 |
OpenACC "present" subarrays: runtime API return value and unmapping fixes
PR libgomp/92511
libgomp/
* oacc-mem.c (present_create_copy): Fix device pointer return value in
case of "present" subarray. Use tgt->tgt_start instead of tgt->to_free
in non-present/create case.
(delete_copyout): Change error condition to fail only on copies outside
of mapped block. Adjust error message accordingly.
* testsuite/libgomp.oacc-c-c++-common/copyin-devptr-1.c: New test.
* testsuite/libgomp.oacc-c-c++-common/copyin-devptr-2.c: New test.
* testsuite/libgomp.oacc-c-c++-common/lib-20.c: Adjust expected error
message.
* testsuite/libgomp.oacc-c-c++-common/lib-23.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/lib-22.c: Allow test to pass now.
* testsuite/libgomp.oacc-c-c++-common/lib-30.c: Likewise.
Reviewed-by: Thomas Schwinge <thomas@codesourcery.com>
From-SVN: r278514
Diffstat (limited to 'libgomp')
-rw-r--r-- | libgomp/ChangeLog | 17 | ||||
-rw-r--r-- | libgomp/oacc-mem.c | 12 | ||||
-rw-r--r-- | libgomp/testsuite/libgomp.oacc-c-c++-common/copyin-devptr-1.c | 28 | ||||
-rw-r--r-- | libgomp/testsuite/libgomp.oacc-c-c++-common/copyin-devptr-2.c | 35 | ||||
-rw-r--r-- | libgomp/testsuite/libgomp.oacc-c-c++-common/lib-20.c | 2 | ||||
-rw-r--r-- | libgomp/testsuite/libgomp.oacc-c-c++-common/lib-22.c | 2 | ||||
-rw-r--r-- | libgomp/testsuite/libgomp.oacc-c-c++-common/lib-23.c | 2 | ||||
-rw-r--r-- | libgomp/testsuite/libgomp.oacc-c-c++-common/lib-30.c | 2 |
8 files changed, 87 insertions, 13 deletions
diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog index a7e6870..9d4a34c 100644 --- a/libgomp/ChangeLog +++ b/libgomp/ChangeLog @@ -1,3 +1,20 @@ +2019-11-20 Julian Brown <julian@codesourcery.com> + + PR libgomp/92511 + + * oacc-mem.c (present_create_copy): Fix device pointer return value in + case of "present" subarray. Use tgt->tgt_start instead of tgt->to_free + in non-present/create case. + (delete_copyout): Change error condition to fail only on copies outside + of mapped block. Adjust error message accordingly. + * testsuite/libgomp.oacc-c-c++-common/copyin-devptr-1.c: New test. + * testsuite/libgomp.oacc-c-c++-common/copyin-devptr-2.c: New test. + * testsuite/libgomp.oacc-c-c++-common/lib-20.c: Adjust expected error + message. + * testsuite/libgomp.oacc-c-c++-common/lib-23.c: Likewise. + * testsuite/libgomp.oacc-c-c++-common/lib-22.c: Allow test to pass now. + * testsuite/libgomp.oacc-c-c++-common/lib-30.c: Likewise. + 2019-11-20 Maciej W. Rozycki <macro@wdc.com> * testsuite/lib/libgomp.exp (libgomp_init): Add flags to find diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c index 2f27100..aafe88d 100644 --- a/libgomp/oacc-mem.c +++ b/libgomp/oacc-mem.c @@ -535,7 +535,7 @@ present_create_copy (unsigned f, void *h, size_t s, int async) if (n) { /* Present. */ - d = (void *) (n->tgt->tgt_start + n->tgt_offset); + d = (void *) (n->tgt->tgt_start + n->tgt_offset + h - n->host_start); if (!(f & FLAG_PRESENT)) { @@ -669,7 +669,6 @@ acc_pcopyin (void *h, size_t s) static void delete_copyout (unsigned f, void *h, size_t s, int async, const char *libfnname) { - size_t host_size; splay_tree_key n; void *d; struct goacc_thread *thr = goacc_thread (); @@ -703,13 +702,12 @@ delete_copyout (unsigned f, void *h, size_t s, int async, const char *libfnname) d = (void *) (n->tgt->tgt_start + n->tgt_offset + (uintptr_t) h - n->host_start); - host_size = n->host_end - n->host_start; - - if (n->host_start != (uintptr_t) h || host_size != s) + if ((uintptr_t) h < n->host_start || (uintptr_t) h + s > n->host_end) { + size_t host_size = n->host_end - n->host_start; gomp_mutex_unlock (&acc_dev->lock); - gomp_fatal ("[%p,%d] surrounds2 [%p,+%d]", - (void *) n->host_start, (int) host_size, (void *) h, (int) s); + gomp_fatal ("[%p,+%d] outside mapped block [%p,+%d]", + (void *) h, (int) s, (void *) n->host_start, (int) host_size); } if (n->refcount == REFCOUNT_INFINITY) diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/copyin-devptr-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/copyin-devptr-1.c new file mode 100644 index 0000000..7e50f3b --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/copyin-devptr-1.c @@ -0,0 +1,28 @@ +/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */ + +#include <openacc.h> +#include <stdlib.h> +#include <assert.h> +#include <stdint.h> + +int main (int argc, char* argv[]) +{ + char *myblock = (char *) malloc (1024); + int i; + void *dst; + for (i = 0; i < 1024; i++) + myblock[i] = i; + dst = acc_copyin (myblock, 1024); + for (i = 0; i < 1024; i += 256) + { + void *partdst = acc_pcopyin (&myblock[i], 256); + assert ((uintptr_t) partdst == (uintptr_t) dst + i); + } + for (i = 0; i < 1024; i += 256) + acc_delete (&myblock[i], 256); + assert (acc_is_present (myblock, 1024)); + acc_delete (myblock, 1024); + assert (!acc_is_present (myblock, 1024)); + free (myblock); + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/copyin-devptr-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/copyin-devptr-2.c new file mode 100644 index 0000000..00e7da1 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/copyin-devptr-2.c @@ -0,0 +1,35 @@ +/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */ + +#include <openacc.h> +#include <stdlib.h> +#include <assert.h> +#include <stdint.h> + +int main (int argc, char* argv[]) +{ + char *block1 = (char *) malloc (1024); + char *block2 = (char *) malloc (1024); + char *block3 = (char *) malloc (1024); + int i; + void *dst; + for (i = 0; i < 1024; i++) + block1[i] = block2[i] = block3[i] = i; + #pragma acc data copyin(block1[0:1024]) copyin(block2[0:1024]) \ + copyin(block3[0:1024]) + { + dst = acc_deviceptr (block2); + for (i = 0; i < 1024; i += 256) + { + void *partdst = acc_pcopyin (&block2[i], 256); + assert ((uintptr_t) partdst == (uintptr_t) dst + i); + } + } + assert (acc_is_present (block2, 1024)); + for (i = 0; i < 1024; i += 256) + acc_delete (&block2[i], 256); + assert (!acc_is_present (block2, 1024)); + free (block1); + free (block2); + free (block3); + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-20.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-20.c index 25ceb3a..10d3cbc 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-20.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-20.c @@ -31,5 +31,5 @@ main (int argc, char **argv) } /* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */ -/* { dg-output "\\\[\[0-9a-fA-FxX\]+,256\\\] surrounds2 \\\[\[0-9a-fA-FxX\]+,\\\+257\\\]" } */ +/* { dg-output "\\\[\[0-9a-fA-FxX\]+,\\\+257\\\] outside mapped block \\\[\[0-9a-fA-FxX\]+,\\\+256\\\]" } */ /* { dg-shouldfail "" } */ diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-22.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-22.c index 65ff440..cb32bbc 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-22.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-22.c @@ -31,5 +31,3 @@ main (int argc, char **argv) } /* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */ -/* { dg-output "\\\[\[0-9a-fA-FxX\]+,256\\\] surrounds2 \\\[\[0-9a-fA-FxX\]+,\\\+255\\\]" } */ -/* { dg-shouldfail "" } */ diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-23.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-23.c index fd4dc59..b1f3e71 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-23.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-23.c @@ -41,5 +41,5 @@ main (int argc, char **argv) } /* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */ -/* { dg-output "\\\[\[0-9a-fA-FxX\]+,256\\\] surrounds2 \\\[\[0-9a-fA-FxX\]+,\\\+512\\\]" } */ +/* { dg-output "\\\[\[0-9a-fA-FxX\]+,\\\+512\\\] outside mapped block \\\[\[0-9a-fA-FxX\]+,\\\+256\\\]" } */ /* { dg-shouldfail "" } */ diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-30.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-30.c index 9bc9ecc..d0e5ffb 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-30.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-30.c @@ -28,5 +28,3 @@ main (int argc, char **argv) } /* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */ -/* { dg-output "\\\[\[0-9a-fA-FxX\]+,256\\\] surrounds2 \\\[\[0-9a-fA-FxX\]+,\\\+254\\\]" } */ -/* { dg-shouldfail "" } */ |