[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

7 files changed, 126 insertions, 17 deletions

diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog
index 739a76d..7606f17 100644
--- a/libgomp/ChangeLog
+++ b/libgomp/ChangeLog
@@ -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...
diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c
index 55c195b..480b9fb 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 0000000..d7ae59d
--- /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 <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 "" } */
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 0000000..751e6eb
--- /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 <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 "" } */
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 0000000..915a270
--- /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 <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 "" } */
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 410c46c..d36a2f1 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 ededf2b..7ebfb8a 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++)