From: Thomas Schwinge Date: Mon, 9 Dec 2019 22:52:36 +0000 (+0100) Subject: [PR92840] [OpenACC] Refuse 'acc_unmap_data' unless mapped by 'acc_map_data' X-Git-Url: https://git.libre-soc.org/?a=commitdiff_plain;h=e103542bc8606e7b5033631e33bdfb9e29191b24;p=gcc.git [PR92840] [OpenACC] Refuse 'acc_unmap_data' unless mapped by 'acc_map_data' libgomp/ PR libgomp/92840 * oacc-mem.c (acc_map_data): Clarify reference counting behavior. (acc_unmap_data): Add error case for 'REFCOUNT_INFINITY'. * testsuite/libgomp.oacc-c-c++-common/acc_unmap_data-pr92840-1.c: New file. * testsuite/libgomp.oacc-c-c++-common/acc_unmap_data-pr92840-2.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/acc_unmap_data-pr92840-3.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/clauses-1.c: Adjust. * testsuite/libgomp.oacc-c-c++-common/nested-1.c: Adjust. From-SVN: r279145 --- diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog index 739a76d48ac..7606f17825d 100644 --- a/libgomp/ChangeLog +++ b/libgomp/ChangeLog @@ -1,5 +1,17 @@ 2019-12-09 Thomas Schwinge + PR libgomp/92840 + * oacc-mem.c (acc_map_data): Clarify reference counting behavior. + (acc_unmap_data): Add error case for 'REFCOUNT_INFINITY'. + * testsuite/libgomp.oacc-c-c++-common/acc_unmap_data-pr92840-1.c: + New file. + * testsuite/libgomp.oacc-c-c++-common/acc_unmap_data-pr92840-2.c: + Likewise. + * testsuite/libgomp.oacc-c-c++-common/acc_unmap_data-pr92840-3.c: + Likewise. + * testsuite/libgomp.oacc-c-c++-common/clauses-1.c: Adjust. + * testsuite/libgomp.oacc-c-c++-common/nested-1.c: Adjust. + PR libgomp/92511 * testsuite/libgomp.oacc-c-c++-common/copyin-devptr-1.c: Remove this file... diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c index 55c195bd819..480b9fbb71b 100644 --- a/libgomp/oacc-mem.c +++ b/libgomp/oacc-mem.c @@ -407,7 +407,11 @@ acc_map_data (void *h, void *d, size_t s) tgt = gomp_map_vars (acc_dev, mapnum, &hostaddrs, &devaddrs, &sizes, &kinds, true, GOMP_MAP_VARS_OPENACC); - tgt->list[0].key->refcount = REFCOUNT_INFINITY; + splay_tree_key n = tgt->list[0].key; + assert (n->refcount == 1); + assert (n->dynamic_refcount == 0); + /* Special reference counting behavior. */ + n->refcount = REFCOUNT_INFINITY; if (profiling_p) { @@ -459,6 +463,18 @@ acc_unmap_data (void *h) gomp_fatal ("[%p,%d] surrounds %p", (void *) n->host_start, (int) host_size, (void *) h); } + /* TODO This currently doesn't catch 'REFCOUNT_INFINITY' usage different from + 'acc_map_data'. Maybe 'dynamic_refcount' can be used for disambiguating + the different 'REFCOUNT_INFINITY' cases, or simply separate + 'REFCOUNT_INFINITY' values per different usage ('REFCOUNT_ACC_MAP_DATA' + etc.)? */ + else if (n->refcount != REFCOUNT_INFINITY) + { + gomp_mutex_unlock (&acc_dev->lock); + gomp_fatal ("refusing to unmap block [%p,+%d] that has not been mapped" + " by 'acc_map_data'", + (void *) h, (int) host_size); + } /* Mark for removal. */ n->refcount = 1; diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_unmap_data-pr92840-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_unmap_data-pr92840-1.c new file mode 100644 index 00000000000..d7ae59dd548 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_unmap_data-pr92840-1.c @@ -0,0 +1,27 @@ +/* Verify that we refuse 'acc_unmap_data', after 'acc_create'. */ + +/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */ + +#include +#include +#include + +int +main () +{ + const int N = 101; + + char *h = (char *) malloc (N); + void *d = acc_create (h, N - 3); + if (!d) + abort (); + + fprintf (stderr, "CheCKpOInT\n"); + acc_unmap_data (h); + + return 0; +} + +/* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */ +/* { dg-output "refusing to unmap block \\\[\[0-9a-fA-FxX\]+,\\\+98\\\] that has not been mapped by 'acc_map_data'" } */ +/* { dg-shouldfail "" } */ diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_unmap_data-pr92840-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_unmap_data-pr92840-2.c new file mode 100644 index 00000000000..751e6ebc33f --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_unmap_data-pr92840-2.c @@ -0,0 +1,25 @@ +/* Verify that we refuse 'acc_unmap_data', after '#pragma acc enter data create'. */ + +/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */ + +#include +#include +#include + +int +main () +{ + const int N = 101; + + char *h = (char *) malloc (N); +#pragma acc enter data create (h[0:N - 77]) + + fprintf (stderr, "CheCKpOInT\n"); + acc_unmap_data (h); + + return 0; +} + +/* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */ +/* { dg-output "refusing to unmap block \\\[\[0-9a-fA-FxX\]+,\\\+24\\\] that has not been mapped by 'acc_map_data'" } */ +/* { dg-shouldfail "" } */ diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_unmap_data-pr92840-3.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_unmap_data-pr92840-3.c new file mode 100644 index 00000000000..915a2705ea9 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_unmap_data-pr92840-3.c @@ -0,0 +1,26 @@ +/* Verify that we refuse 'acc_unmap_data', inside 'data'. */ + +/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */ + +#include +#include +#include + +int +main () +{ + const int N = 101; + + char *h = (char *) malloc (N); +#pragma acc data create (h[0:N - 55]) + { + fprintf (stderr, "CheCKpOInT\n"); + acc_unmap_data (h); + } + + return 0; +} + +/* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */ +/* { dg-output "refusing to unmap block \\\[\[0-9a-fA-FxX\]+,\\\+46\\\] that has not been mapped by 'acc_map_data'" } */ +/* { dg-shouldfail "" } */ diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/clauses-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/clauses-1.c index 410c46cc651..d36a2f1c304 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/clauses-1.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/clauses-1.c @@ -266,13 +266,15 @@ main (int argc, char **argv) if (!acc_is_present (&b[0], (N * sizeof (float)))) abort (); - d = (float *) acc_deviceptr (&a[0]); - acc_unmap_data (&a[0]); - acc_free (d); + acc_delete (&a[0], N * sizeof (float)); - d = (float *) acc_deviceptr (&b[0]); - acc_unmap_data (&b[0]); - acc_free (d); + if (acc_is_present (&a[0], N * sizeof (float))) + abort (); + + acc_delete (&b[0], N * sizeof (float)); + + if (acc_is_present (&b[0], N * sizeof (float))) + abort (); for (i = 0; i < N; i++) { @@ -475,11 +477,10 @@ main (int argc, char **argv) abort (); } - d = (float *) acc_deviceptr (a); - - acc_unmap_data (a); + acc_delete (a, N * sizeof (float)); - acc_free (d); + if (acc_is_present (a, N * sizeof (float))) + abort (); d = (float *) acc_deviceptr (c); diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/nested-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/nested-1.c index ededf2ba5b7..7ebfb8a562b 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/nested-1.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/nested-1.c @@ -290,13 +290,15 @@ main (int argc, char **argv) if (!acc_is_present (&b[0], (N * sizeof (float)))) abort (); - d = (float *) acc_deviceptr (&a[0]); - acc_unmap_data (&a[0]); - acc_free (d); + acc_delete (&a[0], N * sizeof (float)); - d = (float *) acc_deviceptr (&b[0]); - acc_unmap_data (&b[0]); - acc_free (d); + if (acc_is_present (&a[0], N * sizeof (float))) + abort (); + + acc_delete (&b[0], N * sizeof (float)); + + if (acc_is_present (&b[0], N * sizeof (float))) + abort (); for (i = 0; i < N; i++)