From 7c8ffaf54af2c8acb77f82349aac4dd68d47ad9d Mon Sep 17 00:00:00 2001
From: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
Date: Wed, 11 Dec 2019 16:49:27 +0000
Subject: [PATCH] [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.
git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/trunk@279234 138bc75d-0d04-0410-961f-82ee72b054a4
---
libgomp/ChangeLog | 10 +
libgomp/oacc-mem.c | 42 ++--
.../libgomp.oacc-c-c++-common/clauses-1.c | 16 +-
.../libgomp.oacc-c-c++-common/lib-82.c | 6 +-
.../libgomp.oacc-c-c++-common/nested-1.c | 10 +-
.../libgomp.oacc-c-c++-common/pr92843-1.c | 179 ++++++++++++++++++
6 files changed, 242 insertions(+), 21 deletions(-)
create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/pr92843-1.c
@@ -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
@@ -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);
@@ -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++)
@@ -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);
@@ -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);
new file mode 100644
@@ -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;
+}
--
2.17.1