[PR92843] [OpenACC] Fix dynamic reference counting for structured 'REFCOUNT_INFINITY'

	libgomp/
	PR libgomp/92843
	* oacc-mem.c (present_create_copy, delete_copyout): Fix dynamic
	reference counting for structured 'REFCOUNT_INFINITY'.  Add some
	assertions.
	(goacc_insert_pointer, goacc_remove_pointer): Adjust accordingly.
	* testsuite/libgomp.oacc-c-c++-common/pr92843-1.c: New file.
	* testsuite/libgomp.oacc-c-c++-common/clauses-1.c: Fix OpenACC.
	* testsuite/libgomp.oacc-c-c++-common/lib-82.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/nested-1.c: Likewise.

From-SVN: r279234

diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog
index 0a5650ed4384796e6c4b721efc95c5b4a71a44df..e5fb05aea6dd72c44c1851ccfffea5386b0d0653 100644 (file)
--- a/libgomp/ChangeLog
+++ b/libgomp/ChangeLog
@@ -1,5 +1,15 @@
 2019-12-11  Thomas Schwinge  <thomas@codesourcery.com>
 
+       PR libgomp/92843
+       * oacc-mem.c (present_create_copy, delete_copyout): Fix dynamic
+       reference counting for structured 'REFCOUNT_INFINITY'.  Add some
+       assertions.
+       (goacc_insert_pointer, goacc_remove_pointer): Adjust accordingly.
+       * testsuite/libgomp.oacc-c-c++-common/pr92843-1.c: New file.
+       * testsuite/libgomp.oacc-c-c++-common/clauses-1.c: Fix OpenACC.
+       * testsuite/libgomp.oacc-c-c++-common/lib-82.c: Likewise.
+       * testsuite/libgomp.oacc-c-c++-common/nested-1.c: Likewise.
+
        * oacc-parallel.c (find_pointer, GOACC_enter_exit_data): Move...
        * oacc-mem.c: ... here.
        (gomp_acc_insert_pointer, gomp_acc_remove_pointer): Rename to

diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c
index 571e0606ac82cab2aa12c65e1f0a5a4ca89aa05c..a809d0495a6725c9640431421b8484bed7810ca5 100644 (file)
--- a/libgomp/oacc-mem.c
+++ b/libgomp/oacc-mem.c
@@ -543,11 +543,11 @@ present_create_copy (unsigned f, void *h, size_t s, int async)
          gomp_fatal ("[%p,+%d] not mapped", (void *)h, (int)s);
        }
 
+      assert (n->refcount != REFCOUNT_LINK);
       if (n->refcount != REFCOUNT_INFINITY)
-       {
-         n->refcount++;
-         n->dynamic_refcount++;
-       }
+       n->refcount++;
+      n->dynamic_refcount++;
+
       gomp_mutex_unlock (&acc_dev->lock);
     }
   else if (!(f & FLAG_CREATE))
@@ -573,8 +573,10 @@ present_create_copy (unsigned f, void *h, size_t s, int async)
 
       tgt = gomp_map_vars_async (acc_dev, aq, mapnum, &hostaddrs, NULL, &s,
                                 &kinds, true, GOMP_MAP_VARS_OPENACC);
-      /* Initialize dynamic refcount.  */
-      tgt->list[0].key->dynamic_refcount = 1;
+      n = tgt->list[0].key;
+      assert (n->refcount == 1);
+      assert (n->dynamic_refcount == 0);
+      n->dynamic_refcount++;
 
       d = tgt->to_free;
     }
@@ -698,12 +700,9 @@ delete_copyout (unsigned f, void *h, size_t s, int async, const char *libfnname)
                  (void *) h, (int) s, (void *) n->host_start, (int) host_size);
     }
 
-  if (n->refcount == REFCOUNT_INFINITY)
-    {
-      n->refcount = 0;
-      n->dynamic_refcount = 0;
-    }
-  if (n->refcount < n->dynamic_refcount)
+  assert (n->refcount != REFCOUNT_LINK);
+  if (n->refcount != REFCOUNT_INFINITY
+      && n->refcount < n->dynamic_refcount)
     {
       gomp_mutex_unlock (&acc_dev->lock);
       gomp_fatal ("Dynamic reference counting assert fail\n");
@@ -711,13 +710,15 @@ delete_copyout (unsigned f, void *h, size_t s, int async, const char *libfnname)
 
   if (f & FLAG_FINALIZE)
     {
-      n->refcount -= n->dynamic_refcount;
+      if (n->refcount != REFCOUNT_INFINITY)
+       n->refcount -= n->dynamic_refcount;
       n->dynamic_refcount = 0;
     }
   else if (n->dynamic_refcount)
     {
+      if (n->refcount != REFCOUNT_INFINITY)
+       n->refcount--;
       n->dynamic_refcount--;
-      n->refcount--;
     }
 
   if (n->refcount == 0)
@@ -895,6 +896,8 @@ goacc_insert_pointer (size_t mapnum, void **hostaddrs, size_t *sizes,
       splay_tree_key n;
       gomp_mutex_lock (&acc_dev->lock);
       n = lookup_host (acc_dev, *hostaddrs, *sizes);
+      assert (n->refcount != REFCOUNT_INFINITY
+             && n->refcount != REFCOUNT_LINK);
       gomp_mutex_unlock (&acc_dev->lock);
 
       tgt = n->tgt;
@@ -917,10 +920,11 @@ goacc_insert_pointer (size_t mapnum, void **hostaddrs, size_t *sizes,
   goacc_aq aq = get_goacc_asyncqueue (async);
   tgt = gomp_map_vars_async (acc_dev, aq, mapnum, hostaddrs,
                             NULL, sizes, kinds, true, GOMP_MAP_VARS_OPENACC);
+  splay_tree_key n = tgt->list[0].key;
+  assert (n->refcount == 1);
+  assert (n->dynamic_refcount == 0);
+  n->dynamic_refcount++;
   gomp_debug (0, "  %s: mappings prepared\n", __FUNCTION__);
-
-  /* Initialize dynamic refcount.  */
-  tgt->list[0].key->dynamic_refcount = 1;
 }
 
 static void
@@ -950,6 +954,8 @@ goacc_remove_pointer (void *h, size_t s, bool force_copyfrom, int async,
 
   t = n->tgt;
 
+  assert (n->refcount != REFCOUNT_INFINITY
+         && n->refcount != REFCOUNT_LINK);
   if (n->refcount < n->dynamic_refcount)
     {
       gomp_mutex_unlock (&acc_dev->lock);
@@ -963,8 +969,8 @@ goacc_remove_pointer (void *h, size_t s, bool force_copyfrom, int async,
     }
   else if (n->dynamic_refcount)
     {
-      n->dynamic_refcount--;
       n->refcount--;
+      n->dynamic_refcount--;
     }
 
   gomp_mutex_unlock (&acc_dev->lock);

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 b0a96348c3a2085e18030071160482e246764ee2..8f01d5f32f83579d4281e5cbd2fbf139c1a8b9ff 100644 (file)
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/clauses-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/clauses-1.c
@@ -469,7 +469,9 @@ main (int argc, char **argv)
     if (!acc_is_present (c, (N * sizeof (float))))
       abort ();
 
-    acc_copyout (b, N * sizeof (float));
+    d = (float *) acc_deviceptr (b);
+
+    acc_memcpy_from_device (b, d, N * sizeof (float));
 
     for (i = 0; i < N; i++)
     {
@@ -485,10 +487,22 @@ main (int argc, char **argv)
     if (acc_is_present (a, N * sizeof (float)))
       abort ();
 
+    d = (float *) acc_deviceptr (b);
+
+    acc_unmap_data (b);
+
+    if (acc_is_present (b, N * sizeof (float)))
+      abort ();
+
+    acc_free (d);
+
     d = (float *) acc_deviceptr (c);
 
     acc_unmap_data (c);
 
+    if (acc_is_present (c, N * sizeof (float)))
+      abort ();
+
     acc_free (d);
 
     for (i = 0; i < N; i++)

diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-82.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-82.c
index be30a7f28ac074d0f914e627852ae7e22c3852ce..9cf73b31964c315c3f2a67e90b9d7166ef69b287 100644 (file)
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-82.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-82.c
@@ -120,9 +120,13 @@ main (int argc, char **argv)
 
   for (i = 0; i < N; i++)
     {
-      acc_copyout (a[i], nbytes);
+      acc_memcpy_from_device (a[i], d_a[i], nbytes);
       if (*a[i] != i)
        abort ();
+
+      acc_unmap_data (a[i]);
+
+      acc_free (d_a[i]);
     }
 
   free (streams);

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 4c599cda4b38fb81b2d3f18c66a076c9bb1357a1..3e8b42527cf9b3ac5f44aebbef80380dc1471145 100644 (file)
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/nested-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/nested-1.c
@@ -517,7 +517,9 @@ main (int argc, char **argv)
     if (!acc_is_present (c, (N * sizeof (float))))
       abort ();
 
-    acc_copyout (b, N * sizeof (float));
+    d = (float *) acc_deviceptr (b);
+
+    acc_memcpy_from_device (b, d, N * sizeof (float));
 
     for (i = 0; i < N; i++)
     {
@@ -534,6 +536,12 @@ main (int argc, char **argv)
 
     acc_free (d);
 
+    d = (float *) acc_deviceptr (b);
+
+    acc_unmap_data (b);
+
+    acc_free (d);
+
     d = (float *) acc_deviceptr (c);
 
     acc_unmap_data (c);

diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr92843-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr92843-1.c
new file mode 100644 (file)
index 0000000..db5b35b
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr92843-1.c
@@ -0,0 +1,179 @@
+/* Verify that 'acc_copyout' etc. is a no-op if there's still a structured
+   reference count.  */
+
+/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */
+
+#include <assert.h>
+#include <stdlib.h>
+#include <openacc.h>
+
+
+const int c0 = 58;
+const int c1 = 81;
+
+static void
+assign_array (char *array, size_t size, char value)
+{
+  for (size_t i = 0; i < size; ++i)
+    array[i] = value;
+}
+
+static void
+verify_array (const char *array, size_t size, char value)
+{
+  for (size_t i = 0; i < size; ++i)
+    assert (array[i] == value);
+}
+
+
+float global_var;
+#pragma acc declare create (global_var)
+
+static void
+test_acc_declare ()
+{
+  assert (acc_is_present (&global_var, sizeof global_var));
+
+  global_var = c0;
+#pragma acc update device (global_var)
+
+  global_var = c1;
+  acc_copyout (&global_var, sizeof global_var);
+  assert (acc_is_present (&global_var, sizeof global_var));
+  assert (global_var == c1);
+
+  global_var = c1;
+  acc_copyout_finalize (&global_var, sizeof global_var);
+  assert (acc_is_present (&global_var, sizeof global_var));
+  assert (global_var == c1);
+
+  void *global_var_d_p = acc_deviceptr (&global_var);
+  assert (global_var_d_p);
+
+  void *d_p = acc_copyin (&global_var, sizeof global_var);
+  assert (d_p == global_var_d_p);
+
+  acc_copyout (&global_var, sizeof global_var);
+  assert (acc_is_present (&global_var, sizeof global_var));
+
+  d_p = acc_copyin (&global_var, sizeof global_var);
+  assert (d_p == global_var_d_p);
+
+  d_p = acc_copyin (&global_var, sizeof global_var);
+  assert (d_p == global_var_d_p);
+
+  global_var = c1;
+  acc_copyout_finalize (&global_var, sizeof global_var);
+  assert (acc_is_present (&global_var, sizeof global_var));
+  assert (global_var == c1);
+
+  global_var = c1;
+  acc_copyout (&global_var, sizeof global_var);
+  assert (acc_is_present (&global_var, sizeof global_var));
+  assert (global_var == c1);
+}
+
+
+static void
+test_acc_map_data ()
+{
+  const int N = 801;
+
+  char *h = (char *) malloc (N);
+  assert (h);
+  void *d = acc_malloc (N);
+  assert (d);
+  acc_map_data (h, d, N);
+  assert (acc_is_present (h, N));
+
+  assign_array (h, N, c0);
+#pragma acc update device (h[0:N])
+
+  assign_array (h, N, c1);
+#pragma acc exit data copyout (h[0:N])
+  assert (acc_is_present (h, N));
+  verify_array (h, N, c1);
+
+  assign_array (h, N, c1);
+#pragma acc exit data copyout (h[0:N]) finalize
+  assert (acc_is_present (h, N));
+  verify_array (h, N, c1);
+
+#pragma acc enter data copyin (h[0:N])
+
+  assign_array (h, N, c1);
+#pragma acc exit data copyout (h[0:N])
+  assert (acc_is_present (h, N));
+  verify_array (h, N, c1);
+
+#pragma acc enter data copyin (h[0:N])
+
+#pragma acc enter data copyin (h[0:N])
+
+  assign_array (h, N, c1);
+#pragma acc exit data copyout (h[0:N]) finalize
+  assert (acc_is_present (h, N));
+  verify_array (h, N, c1);
+
+  assign_array (h, N, c1);
+#pragma acc exit data copyout (h[0:N])
+  assert (acc_is_present (h, N));
+  verify_array (h, N, c1);
+}
+
+
+static void
+test_acc_data ()
+{
+#define N 23
+  char h[N];
+
+  assign_array (h, N, c0);
+#pragma acc data copyin (h)
+  {
+    assert (acc_is_present (h, sizeof h));
+
+    assign_array (h, N, c1);
+    acc_copyout_finalize (h, sizeof h);
+    assert (acc_is_present (h, sizeof h));
+    verify_array (h, N, c1);
+
+    assign_array (h, N, c1);
+    acc_copyout (h, sizeof h);
+    assert (acc_is_present (h, sizeof h));
+    verify_array (h, N, c1);
+
+    acc_copyin (h, sizeof h);
+
+    assign_array (h, N, c1);
+    acc_copyout (h, sizeof h);
+    assert (acc_is_present (h, sizeof h));
+    verify_array (h, N, c1);
+
+    acc_copyin (h, sizeof h);
+
+    acc_copyin (h, sizeof h);
+
+    assign_array (h, N, c1);
+    acc_copyout_finalize (h, sizeof h);
+    assert (acc_is_present (h, sizeof h));
+    verify_array (h, N, c1);
+
+    assign_array (h, N, c1);
+    acc_copyout (h, sizeof h);
+    assert (acc_is_present (h, sizeof h));
+    verify_array (h, N, c1);
+  }
+#undef N
+}
+
+
+int
+main ()
+{
+  test_acc_declare ();
+  test_acc_map_data ();
+  test_acc_data ();
+
+  return 0;
+}