From: Julian Brown Date: Wed, 20 Nov 2019 17:51:09 +0000 (+0000) Subject: OpenACC "present" subarrays: runtime API return value and unmapping fixes X-Git-Url: https://git.libre-soc.org/?a=commitdiff_plain;h=e307b05f4377a22811395f6a051d9db864b8785f;p=gcc.git 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 From-SVN: r278514 --- diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog index a7e6870c3ea..9d4a34c90a3 100644 --- a/libgomp/ChangeLog +++ b/libgomp/ChangeLog @@ -1,3 +1,20 @@ +2019-11-20 Julian Brown + + 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 * testsuite/lib/libgomp.exp (libgomp_init): Add flags to find diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c index 2f271009fb8..aafe88d3a14 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 00000000000..7e50f3b892e --- /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 +#include +#include +#include + +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 00000000000..00e7da1f128 --- /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 +#include +#include +#include + +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 25ceb3a26af..10d3cbc5cc6 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 65ff440a528..cb32bbcb652 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 fd4dc5971a1..b1f3e71f278 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 9bc9ecc1068..d0e5ffb0691 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 "" } */