From bea573cb7ea13cece9c51ca9eb1cc9c34005dedf Mon Sep 17 00:00:00 2001
From: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
Date: Mon, 9 Dec 2019 22:52:36 +0000
Subject: [PATCH] [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.
git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/trunk@279145 138bc75d-0d04-0410-961f-82ee72b054a4
---
libgomp/ChangeLog | 12 +++++++++
libgomp/oacc-mem.c | 18 ++++++++++++-
.../acc_unmap_data-pr92840-1.c | 27 +++++++++++++++++++
.../acc_unmap_data-pr92840-2.c | 25 +++++++++++++++++
.../acc_unmap_data-pr92840-3.c | 26 ++++++++++++++++++
.../libgomp.oacc-c-c++-common/clauses-1.c | 21 ++++++++-------
.../libgomp.oacc-c-c++-common/nested-1.c | 14 +++++-----
7 files changed, 126 insertions(+), 17 deletions(-)
create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/acc_unmap_data-pr92840-1.c
create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/acc_unmap_data-pr92840-2.c
create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/acc_unmap_data-pr92840-3.c
@@ -1,5 +1,17 @@
2019-12-09 Thomas Schwinge <thomas@codesourcery.com>
+ 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...
@@ -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;
new file mode 100644
@@ -0,0 +1,27 @@
+/* Verify that we refuse 'acc_unmap_data', after 'acc_create'. */
+
+/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */
+
+#include <stdio.h>
+#include <stdlib.h>
+#include <openacc.h>
+
+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 "" } */
new file mode 100644
@@ -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 <stdio.h>
+#include <stdlib.h>
+#include <openacc.h>
+
+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 "" } */
new file mode 100644
@@ -0,0 +1,26 @@
+/* Verify that we refuse 'acc_unmap_data', inside 'data'. */
+
+/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */
+
+#include <stdio.h>
+#include <stdlib.h>
+#include <openacc.h>
+
+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 "" } */
@@ -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);
@@ -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++)
--
2.17.1