From 03383a93c7318009ddd0e8d77b1a950c4b2b8f5a Mon Sep 17 00:00:00 2001
From: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
Date: Mon, 9 Dec 2019 22:52:47 +0000
Subject: [PATCH] [PR92503] [OpenACC] Don't silently 'acc_unmap_data' in
'acc_free'
libgomp/
PR libgomp/92503
* oacc-mem.c (acc_free): Error out instead of 'acc_unmap_data'.
* testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-1.c: New
file.
* testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-2.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-3-2.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-3.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-4-2.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-4.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/clauses-1.c: Adjust.
* testsuite/libgomp.oacc-c-c++-common/context-1.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/context-2.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/context-3.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/context-4.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/lib-13.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/lib-14.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/lib-18.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/lib-91.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/nested-1.c: Likewise.
git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/trunk@279146 138bc75d-0d04-0410-961f-82ee72b054a4
---
libgomp/ChangeLog | 25 +++++++++++++++
libgomp/oacc-mem.c | 17 +++++-----
.../acc_free-pr92503-1.c | 28 ++++++++++++++++
.../acc_free-pr92503-2.c | 27 ++++++++++++++++
.../acc_free-pr92503-3-2.c | 28 ++++++++++++++++
.../acc_free-pr92503-3.c | 28 ++++++++++++++++
.../acc_free-pr92503-4-2.c | 31 ++++++++++++++++++
.../acc_free-pr92503-4.c | 32 +++++++++++++++++++
.../libgomp.oacc-c-c++-common/clauses-1.c | 12 +++++--
.../libgomp.oacc-c-c++-common/context-1.c | 6 ++--
.../libgomp.oacc-c-c++-common/context-2.c | 6 ++--
.../libgomp.oacc-c-c++-common/context-3.c | 6 ++--
.../libgomp.oacc-c-c++-common/context-4.c | 6 ++--
.../libgomp.oacc-c-c++-common/lib-13.c | 2 +-
.../libgomp.oacc-c-c++-common/lib-14.c | 2 +-
.../libgomp.oacc-c-c++-common/lib-18.c | 2 +-
.../libgomp.oacc-c-c++-common/lib-91.c | 2 ++
.../libgomp.oacc-c-c++-common/nested-1.c | 12 +++++--
18 files changed, 242 insertions(+), 30 deletions(-)
create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-1.c
create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-2.c
create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-3-2.c
create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-3.c
create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-4-2.c
create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-4.c
@@ -1,5 +1,30 @@
2019-12-09 Thomas Schwinge <thomas@codesourcery.com>
+ PR libgomp/92503
+ * oacc-mem.c (acc_free): Error out instead of 'acc_unmap_data'.
+ * testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-1.c: New
+ file.
+ * testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-2.c:
+ Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-3-2.c:
+ Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-3.c:
+ Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-4-2.c:
+ Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-4.c:
+ Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/clauses-1.c: Adjust.
+ * testsuite/libgomp.oacc-c-c++-common/context-1.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/context-2.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/context-3.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/context-4.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-13.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-14.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-18.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-91.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/nested-1.c: Likewise.
+
PR libgomp/92840
* oacc-mem.c (acc_map_data): Clarify reference counting behavior.
(acc_unmap_data): Add error case for 'REFCOUNT_INFINITY'.
@@ -121,9 +121,6 @@ acc_malloc (size_t s)
return res;
}
-/* OpenACC 2.0a (3.2.16) doesn't specify what to do in the event
- the device address is mapped. We choose to check if it mapped,
- and if it is, to unmap it. */
void
acc_free (void *d)
{
@@ -152,13 +149,15 @@ acc_free (void *d)
(unless you got that null from acc_malloc). */
if ((k = lookup_dev (acc_dev->openacc.data_environ, d, 1)))
{
- void *offset;
-
- offset = d - k->tgt->tgt_start + k->tgt_offset;
-
+ void *offset = d - k->tgt->tgt_start + k->tgt_offset;
+ void *h = k->host_start + offset;
+ size_t h_size = k->host_end - k->host_start;
gomp_mutex_unlock (&acc_dev->lock);
-
- acc_unmap_data ((void *)(k->host_start + offset));
+ /* PR92503 "[OpenACC] Behavior of 'acc_free' if the memory space is still
+ used in a mapping". */
+ gomp_fatal ("refusing to free device memory space at %p that is still"
+ " mapped at [%p,+%d]",
+ d, h, (int) h_size);
}
else
gomp_mutex_unlock (&acc_dev->lock);
new file mode 100644
@@ -0,0 +1,28 @@
+/* Verify that we refuse 'acc_free', after 'acc_map_data'. */
+
+/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */
+
+#include <stdio.h>
+#include <stdlib.h>
+#include <openacc.h>
+
+int
+main ()
+{
+ const int N = 108;
+
+ char *h = (char *) malloc (N);
+ void *d = acc_malloc (N - 10);
+ if (!d)
+ abort ();
+ acc_map_data (h, d, N - 19);
+
+ fprintf (stderr, "CheCKpOInT\n");
+ acc_free (d);
+
+ return 0;
+}
+
+/* { dg-output "CheCKpOInT(\n|\r\n|\r).*" }
+ { dg-output "refusing to free device memory space at \[0-9a-fA-FxX\]+ that is still mapped at \\\[\[0-9a-fA-FxX\]+,\\\+89\\\]" }
+ { dg-shouldfail "" } */
new file mode 100644
@@ -0,0 +1,27 @@
+/* Verify that we refuse 'acc_free', after 'acc_create'. */
+
+/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */
+
+#include <stdio.h>
+#include <stdlib.h>
+#include <openacc.h>
+
+int
+main ()
+{
+ const int N = 108;
+
+ char *h = (char *) malloc (N);
+ void *d = acc_create (h, N - 1);
+ if (!d)
+ abort ();
+
+ fprintf (stderr, "CheCKpOInT\n");
+ acc_free (d);
+
+ return 0;
+}
+
+/* { dg-output "CheCKpOInT(\n|\r\n|\r).*" }
+ { dg-output "refusing to free device memory space at \[0-9a-fA-FxX\]+ that is still mapped at \\\[\[0-9a-fA-FxX\]+,\\\+107\\\]" }
+ { dg-shouldfail "" } */
new file mode 100644
@@ -0,0 +1,28 @@
+/* Verify that we refuse 'acc_free', inside 'host_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 = 108;
+
+ char *h = (char *) malloc (N);
+#pragma acc enter data create (h[0:N - 2])
+
+#pragma acc host_data use_device (h)
+ {
+ fprintf (stderr, "CheCKpOInT\n");
+ acc_free (h);
+ }
+
+ return 0;
+}
+
+/* { dg-output "CheCKpOInT(\n|\r\n|\r).*" }
+ { dg-output "refusing to free device memory space at \[0-9a-fA-FxX\]+ that is still mapped at \\\[\[0-9a-fA-FxX\]+,\\\+106\\\]" }
+ { dg-shouldfail "" } */
new file mode 100644
@@ -0,0 +1,28 @@
+/* Verify that we refuse 'acc_free', 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 = 108;
+
+ char *h = (char *) malloc (N);
+#pragma acc enter data create (h[0:N - 3])
+ void *d = acc_deviceptr (h);
+ if (!d)
+ abort ();
+
+ fprintf (stderr, "CheCKpOInT\n");
+ acc_free (d);
+
+ return 0;
+}
+
+/* { dg-output "CheCKpOInT(\n|\r\n|\r).*" }
+ { dg-output "refusing to free device memory space at \[0-9a-fA-FxX\]+ that is still mapped at \\\[\[0-9a-fA-FxX\]+,\\\+105\\\]" }
+ { dg-shouldfail "" } */
new file mode 100644
@@ -0,0 +1,31 @@
+/* Verify that we refuse 'acc_free', inside 'host_data', inside 'data'. */
+
+/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */
+
+#include <stdio.h>
+#include <stdlib.h>
+#include <openacc.h>
+
+int
+main ()
+{
+ const int N = 108;
+
+ char *h = (char *) malloc (N);
+#pragma acc data create (h[0:N - 44])
+ {
+#pragma acc host_data use_device (h)
+ {
+ fprintf (stderr, "CheCKpOInT\n");
+ acc_free (h);
+ }
+ }
+
+ return 0;
+}
+
+/* { dg-output "CheCKpOInT(\n|\r\n|\r).*" }
+ TODO PR92877
+ { dg-output "libgomp: cuMemGetAddressRange_v2 error: named symbol not found" { target openacc_nvidia_accel_selected } }
+ { dg-output "refusing to free device memory space at \[0-9a-fA-FxX\]+ that is still mapped at \\\[\[0-9a-fA-FxX\]+,\\\+64\\\]" { xfail *-*-* } }
+ { dg-shouldfail "" } */
new file mode 100644
@@ -0,0 +1,32 @@
+/* Verify that we refuse 'acc_free', inside 'data'. */
+
+/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */
+
+#include <stdio.h>
+#include <stdlib.h>
+#include <openacc.h>
+
+int
+main ()
+{
+ const int N = 108;
+
+ char *h = (char *) malloc (N);
+#pragma acc data create (h[0:N - 21])
+ {
+ void *d = acc_deviceptr (h);
+ if (!d)
+ abort ();
+
+ fprintf (stderr, "CheCKpOInT\n");
+ acc_free (d);
+ }
+
+ return 0;
+}
+
+/* { dg-output "CheCKpOInT(\n|\r\n|\r).*" }
+ TODO PR92877
+ { dg-output "libgomp: cuMemGetAddressRange_v2 error: named symbol not found" { target openacc_nvidia_accel_selected } }
+ { dg-output "refusing to free device memory space at \[0-9a-fA-FxX\]+ that is still mapped at \\\[\[0-9a-fA-FxX\]+,\\\+87\\\]" { xfail *-*-* } }
+ { dg-shouldfail "" } */
@@ -103,7 +103,10 @@ main (int argc, char **argv)
if (acc_is_present (&b[0], (N * sizeof (float))))
abort ();
- acc_free (d);
+ acc_delete (&a[0], N * sizeof (float));
+
+ if (acc_is_present (&a[0], N * sizeof (float)))
+ abort ();
for (i = 0; i < N; i++)
{
@@ -162,7 +165,7 @@ main (int argc, char **argv)
if (!acc_is_present (&b[0], (N * sizeof (float))))
abort ();
- acc_free (d);
+ acc_delete (&b[0], N * sizeof (float));
if (acc_is_present (&b[0], (N * sizeof (float))))
abort ();
@@ -557,7 +560,10 @@ main (int argc, char **argv)
if (acc_is_present (&b[0], (N * sizeof (float))))
abort ();
- acc_free (d);
+ acc_delete (&a[0], N * sizeof (float));
+
+ if (acc_is_present (&a[0], N * sizeof (float)))
+ abort ();
for (i = 0; i < N; i++)
{
@@ -172,13 +172,13 @@ main (int argc, char **argv)
exit (EXIT_FAILURE);
}
+ acc_delete (&h_X[0], N * sizeof (float));
+ acc_delete (&h_Y1[0], N * sizeof (float));
+
free (h_X);
free (h_Y1);
free (h_Y2);
- acc_free (d_X);
- acc_free (d_Y);
-
context_check (pctx);
s = cublasDestroy (h);
@@ -182,13 +182,13 @@ main (int argc, char **argv)
exit (EXIT_FAILURE);
}
+ acc_delete (&h_X[0], N * sizeof (float));
+ acc_delete (&h_Y1[0], N * sizeof (float));
+
free (h_X);
free (h_Y1);
free (h_Y2);
- acc_free (d_X);
- acc_free (d_Y);
-
context_check (pctx);
s = cublasDestroy (h);
@@ -163,13 +163,13 @@ main (int argc, char **argv)
exit (EXIT_FAILURE);
}
+ acc_delete (&h_X[0], N * sizeof (float));
+ acc_delete (&h_Y1[0], N * sizeof (float));
+
free (h_X);
free (h_Y1);
free (h_Y2);
- acc_free (d_X);
- acc_free (d_Y);
-
context_check (pctx);
s = cublasDestroy (h);
@@ -176,13 +176,13 @@ main (int argc, char **argv)
exit (EXIT_FAILURE);
}
+ acc_delete (&h_X[0], N * sizeof (float));
+ acc_delete (&h_Y1[0], N * sizeof (float));
+
free (h_X);
free (h_Y1);
free (h_Y2);
- acc_free (d_X);
- acc_free (d_Y);
-
context_check (pctx);
s = cublasDestroy (h);
@@ -51,7 +51,7 @@ main (int argc, char **argv)
if (acc_is_present (h, 0) != 0)
abort ();
- acc_free (d);
+ acc_delete (h, N);
if (acc_is_present (h, 1) != 0)
abort ();
@@ -48,7 +48,7 @@ main (int argc, char **argv)
abort ();
}
- acc_free (d);
+ acc_delete (h, N);
for (i = 0; i < N; i++)
{
@@ -23,7 +23,7 @@ main (int argc, char **argv)
d = acc_copyin (h, N);
- acc_free (d);
+ acc_delete (h, N);
fprintf (stderr, "CheCKpOInT\n");
acc_copyout (h, N);
@@ -72,6 +72,8 @@ main (int argc, char **argv)
if (async > (sync * 1.5))
abort ();
+ acc_unmap_data (h);
+
acc_free (d);
free (h);
@@ -112,7 +112,10 @@ main (int argc, char **argv)
if (acc_is_present (&b[0], (N * sizeof (float))))
abort ();
- acc_free (d);
+ acc_delete (&a[0], N * sizeof (float));
+
+ if (acc_is_present (&a[0], N * sizeof (float)))
+ abort ();
for (i = 0; i < N; i++)
{
@@ -177,7 +180,7 @@ main (int argc, char **argv)
if (!acc_is_present (&b[0], (N * sizeof (float))))
abort ();
- acc_free (d);
+ acc_delete (&b[0], N * sizeof (float));
if (acc_is_present (&b[0], (N * sizeof (float))))
abort ();
@@ -609,7 +612,10 @@ main (int argc, char **argv)
if (acc_is_present (&b[0], (N * sizeof (float))))
abort ();
- acc_free (d);
+ acc_delete (&a[0], N * sizeof (float));
+
+ if (acc_is_present (&a[0], N * sizeof (float)))
+ abort ();
for (i = 0; i < N; i++)
{
--
2.17.1