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'.
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)
{
(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);
--- /dev/null
+/* 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 "" } */
--- /dev/null
+/* 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 "" } */
--- /dev/null
+/* 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 "" } */
--- /dev/null
+/* 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 "" } */
--- /dev/null
+/* 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 "" } */
--- /dev/null
+/* 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 "" } */
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++)
{
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 ();
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++)
{
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);
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);
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);
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);
if (acc_is_present (h, 0) != 0)
abort ();
- acc_free (d);
+ acc_delete (h, N);
if (acc_is_present (h, 1) != 0)
abort ();
abort ();
}
- acc_free (d);
+ acc_delete (h, N);
for (i = 0; i < N; i++)
{
d = acc_copyin (h, N);
- acc_free (d);
+ acc_delete (h, N);
fprintf (stderr, "CheCKpOInT\n");
acc_copyout (h, N);
if (async > (sync * 1.5))
abort ();
+ acc_unmap_data (h);
+
acc_free (d);
free (h);
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++)
{
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 ();
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++)
{
git://gcc.gnu.org / gcc.git / commitdiff