[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.

From-SVN: r279146

18 files changed, 242 insertions, 30 deletions

diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog
index 7606f17..62092a2 100644
--- a/libgomp/ChangeLog
+++ b/libgomp/ChangeLog
@@ -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'.
diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c
index 480b9fb..81ebddf 100644
--- a/libgomp/oacc-mem.c
+++ b/libgomp/oacc-mem.c
@@ -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);
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-1.c
new file mode 100644
index 0000000..4fc6068
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-1.c
@@ -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 "" } */
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-2.c
new file mode 100644
index 0000000..3f6a8e5
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-2.c
@@ -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 "" } */
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-3-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-3-2.c
new file mode 100644
index 0000000..9f45048
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-3-2.c
@@ -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 "" } */
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-3.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-3.c
new file mode 100644
index 0000000..1620830
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-3.c
@@ -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 "" } */
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-4-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-4-2.c
new file mode 100644
index 0000000..bbf4431
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-4-2.c
@@ -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 "" } */
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-4.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-4.c
new file mode 100644
index 0000000..6212f9e
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-4.c
@@ -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 "" } */
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 d36a2f1..b0a9634 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/clauses-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/clauses-1.c
@@ -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++)
     {
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/context-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/context-1.c
index dabc706..2e3b62e 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/context-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/context-1.c
@@ -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);
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/context-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/context-2.c
index 6a52f74..6bdcfe7 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/context-2.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/context-2.c
@@ -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);
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/context-3.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/context-3.c
index ccd276c..8f14560 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/context-3.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/context-3.c
@@ -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);
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/context-4.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/context-4.c
index 71365e8..b403a5c 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/context-4.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/context-4.c
@@ -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);
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-13.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-13.c
index d665533..aca4c25 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-13.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-13.c
@@ -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 ();
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-14.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-14.c
index ee21257..de6d38b 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-14.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-14.c
@@ -48,7 +48,7 @@ main (int argc, char **argv)
 	abort ();
     }
 
-  acc_free (d);
+  acc_delete (h, N);
 
   for (i = 0; i < N; i++)
     {
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-18.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-18.c
index b686cc9..93bfb99 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-18.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-18.c
@@ -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);
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-91.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-91.c
index e00ef4f..36fff08 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-91.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-91.c
@@ -72,6 +72,8 @@ main (int argc, char **argv)
   if (async > (sync * 1.5))
     abort ();
 
+  acc_unmap_data (h);
+
   acc_free (d);
 
   free (h);
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 7ebfb8a..4c599cd 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/nested-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/nested-1.c
@@ -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++)
     {