[PATCH 3/3] OpenACC dynamic data lifetimes ending within structured blocks
Julian Brown
julian@codesourcery.com
Fri Jan 17 21:31:00 GMT 2020
This patch adds a new function to logically decrement the "dynamic
reference counter" for a mapped OpenACC variable, and handles some cases
in which that counter drops to zero inside a structured data
block. Previously, it's likely that at least in some cases, ending a
dynamic data lifetime in this way could behave unpredictably.
Several new test cases are included.
This patch is strongly related to the previous two, but is somewhat of
a separate change, and those two patches can stand alone if this one
gets deferred.
Tested alongside the previous patches in the series with offloading to NVPTX.
OK?
Thanks,
Julian
ChangeLog
libgomp/
* oacc-mem.c (decr_dynamic_refcount): New function.
(goacc_exit_datum): Call above function.
(goacc_exit_data_internal): Call above function.
* testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-1.c: New
test.
* testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-1-lib.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-6.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-6-lib.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-7.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-7-lib.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-8.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-8-lib.c:
Likewise.
---
libgomp/oacc-mem.c | 128 ++++++++++----
.../static-dynamic-lifetimes-1-lib.c | 3 +
.../static-dynamic-lifetimes-1.c | 160 ++++++++++++++++++
.../static-dynamic-lifetimes-6-lib.c | 5 +
.../static-dynamic-lifetimes-6.c | 46 +++++
.../static-dynamic-lifetimes-7-lib.c | 5 +
.../static-dynamic-lifetimes-7.c | 45 +++++
.../static-dynamic-lifetimes-8-lib.c | 5 +
.../static-dynamic-lifetimes-8.c | 50 ++++++
9 files changed, 412 insertions(+), 35 deletions(-)
create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-1-lib.c
create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-1.c
create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-6-lib.c
create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-6.c
create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-7-lib.c
create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-7.c
create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-8-lib.c
create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-8.c
diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c
index 783e7f363fb..f34ffa67079 100644
--- a/libgomp/oacc-mem.c
+++ b/libgomp/oacc-mem.c
@@ -725,6 +725,92 @@ acc_pcopyin (void *h, size_t s)
#endif
+/* Perform actions necessary to decrement the dynamic reference counter for
+ splay tree key N. Returns TRUE on success, or FALSE on failure (e.g. if we
+ hit a case we can't presently handle inside a data region). */
+
+static bool
+decr_dynamic_refcount (splay_tree_key n, bool finalize)
+{
+ if (finalize)
+ {
+ if (n->refcount != REFCOUNT_INFINITY)
+ n->refcount -= n->virtual_refcount;
+ n->virtual_refcount = 0;
+ }
+
+ if (n->virtual_refcount > 0)
+ {
+ if (n->refcount != REFCOUNT_INFINITY)
+ n->refcount--;
+ n->virtual_refcount--;
+ }
+ /* An initial "enter data" mapping might create a target_mem_desc (in
+ gomp_map_vars_async via goacc_enter_datum or
+ goacc_enter_data_internal). In that case we have a structural
+ reference count but a zero virtual reference count: we nevertheless
+ want to do the "exit data" operation here. Detect the special case
+ using a sentinel value stored in the "prev" field, which is otherwise
+ unused for dynamic data mappings. */
+ else if (n->refcount > 0
+ && n->refcount != REFCOUNT_INFINITY
+ && n->tgt->prev == &dyn_tgt_sentinel)
+ {
+ n->refcount--;
+ /* We know n->virtual_refcount is zero here, so if we still have a
+ non-zero n->refcount we are ending a dynamically-scoped variable
+ lifetime in the middle of a static lifetime for the same variable.
+ If we're not careful this results in a dangling reference. Attempt
+ to handle this here, if only in simple cases. E.g.:
+
+ #pragma acc enter data copyin(var)
+ #pragma acc data copy(var{, ...})
+ {
+ #pragma acc exit data copyout(var)
+ }
+
+ Here (the "exit data"), we reattach the relevant fields of the
+ previously dynamically-scoped target_mem_desc to the static data
+ region's target_mem_desc, hence merging the former into the latter.
+ The old dynamic target_mem_desc can then be freed.
+
+ We can't deal with static data regions that refer to existing dynamic
+ data mappings or that introduce new static lifetimes of their own. */
+ if (n->refcount > 0
+ && n->tgt->list_count == 1
+ && n->tgt->refcount == 1)
+ {
+ struct goacc_thread *thr = goacc_thread ();
+ struct target_mem_desc *tgt, *static_tgt = NULL;
+ for (tgt = thr->mapped_data;
+ tgt != NULL && static_tgt == NULL;
+ tgt = tgt->prev)
+ for (int j = 0; j < tgt->list_count; j++)
+ if (tgt->list[j].key == n)
+ {
+ static_tgt = tgt;
+ break;
+ }
+ if (!static_tgt
+ || static_tgt->to_free != NULL
+ || static_tgt->array != NULL)
+ return false;
+ static_tgt->to_free = n->tgt->to_free;
+ static_tgt->array = n->tgt->array;
+ static_tgt->tgt_start = n->tgt->tgt_start;
+ static_tgt->tgt_end = n->tgt->tgt_end;
+ static_tgt->to_free = n->tgt->to_free;
+ static_tgt->refcount++;
+ free (n->tgt);
+ n->tgt = static_tgt;
+ }
+ else if (n->refcount > 0)
+ return false;
+ }
+
+ return true;
+}
+
/* Exit a dynamic mapping for a single variable. */
static void
@@ -767,29 +853,12 @@ goacc_exit_datum (void *h, size_t s, unsigned short kind, int async)
bool finalize = (kind == GOMP_MAP_DELETE
|| kind == GOMP_MAP_FORCE_FROM);
- if (finalize)
- {
- if (n->refcount != REFCOUNT_INFINITY)
- n->refcount -= n->virtual_refcount;
- n->virtual_refcount = 0;
- }
- if (n->virtual_refcount > 0)
+ if (!decr_dynamic_refcount (n, finalize))
{
- if (n->refcount != REFCOUNT_INFINITY)
- n->refcount--;
- n->virtual_refcount--;
+ gomp_mutex_unlock (&acc_dev->lock);
+ gomp_fatal ("cannot handle delete/copyout within data region");
}
- /* An initial "enter data" mapping might create a target_mem_desc (in
- gomp_map_vars_async via goacc_enter_datum). In that case we have a
- structural reference count but a zero virtual reference count: we
- nevertheless want to do the "exit data" operation here. Detect the
- special case using a sentinel value stored in the "prev" field, which is
- otherwise unused for dynamic data mappings. */
- else if (n->refcount > 0
- && n->refcount != REFCOUNT_INFINITY
- && n->tgt->prev == &dyn_tgt_sentinel)
- n->refcount--;
if (n->refcount == 0)
{
@@ -1216,23 +1285,12 @@ goacc_exit_data_internal (struct gomp_device_descr *acc_dev, size_t mapnum,
if (n == NULL)
continue;
- if (finalize)
- {
- if (n->refcount != REFCOUNT_INFINITY)
- n->refcount -= n->virtual_refcount;
- n->virtual_refcount = 0;
- }
-
- if (n->virtual_refcount > 0)
+ if (!decr_dynamic_refcount (n, finalize))
{
- if (n->refcount != REFCOUNT_INFINITY)
- n->refcount--;
- n->virtual_refcount--;
+ /* The user is trying to do something too tricky for us. */
+ gomp_mutex_unlock (&acc_dev->lock);
+ gomp_fatal ("cannot handle 'exit data' within data region");
}
- else if (n->refcount > 0
- && n->refcount != REFCOUNT_INFINITY
- && n->tgt->prev == &dyn_tgt_sentinel)
- n->refcount--;
if (copyfrom
&& n->refcount != REFCOUNT_INFINITY
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-1-lib.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-1-lib.c
new file mode 100644
index 00000000000..23c20d4fab7
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-1-lib.c
@@ -0,0 +1,3 @@
+/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
+/* { dg-additional-options "-DOPENACC_API" } */
+#include "static-dynamic-lifetimes-1.c"
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-1.c
new file mode 100644
index 00000000000..a743660f53e
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-1.c
@@ -0,0 +1,160 @@
+/* Test transitioning of data lifetimes between static and dynamic. */
+
+/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
+
+#include <openacc.h>
+#include <assert.h>
+#include <stdlib.h>
+
+#define SIZE 1024
+
+void
+f1 (void)
+{
+ char *block1 = (char *) malloc (SIZE);
+
+#ifdef OPENACC_API
+ acc_copyin (block1, SIZE);
+ acc_copyin (block1, SIZE);
+#else
+#pragma acc enter data copyin(block1[0:SIZE])
+#pragma acc enter data copyin(block1[0:SIZE])
+#endif
+
+#pragma acc data copy(block1[0:SIZE])
+ {
+#ifdef OPENACC_API
+ acc_copyin (block1, SIZE);
+#else
+#pragma acc enter data copyin(block1[0:SIZE])
+#endif
+ }
+
+ assert (acc_is_present (block1, SIZE));
+
+#ifdef OPENACC_API
+ acc_copyout (block1, SIZE);
+ assert (acc_is_present (block1, SIZE));
+ acc_copyout (block1, SIZE);
+ assert (acc_is_present (block1, SIZE));
+ acc_copyout (block1, SIZE);
+ assert (!acc_is_present (block1, SIZE));
+#else
+#pragma acc exit data copyout(block1[0:SIZE])
+ assert (acc_is_present (block1, SIZE));
+#pragma acc exit data copyout(block1[0:SIZE])
+ assert (acc_is_present (block1, SIZE));
+#pragma acc exit data copyout(block1[0:SIZE])
+ assert (!acc_is_present (block1, SIZE));
+#endif
+
+ free (block1);
+}
+
+void
+f2 (void)
+{
+ char *block1 = (char *) malloc (SIZE);
+
+#ifdef OPENACC_API
+ acc_copyin (block1, SIZE);
+#else
+#pragma acc enter data copyin(block1[0:SIZE])
+#endif
+
+#pragma acc data copy(block1[0:SIZE])
+ {
+#ifdef OPENACC_API
+ acc_copyout (block1, SIZE);
+#else
+#pragma acc exit data copyout(block1[0:SIZE])
+#endif
+ /* This should stay present until the end of the static data lifetime. */
+ assert (acc_is_present (block1, SIZE));
+ }
+
+ assert (!acc_is_present (block1, SIZE));
+
+ free (block1);
+}
+
+void
+f3 (void)
+{
+ char *block1 = (char *) malloc (SIZE);
+
+#ifdef OPENACC_API
+ acc_copyin (block1, SIZE);
+#else
+#pragma acc enter data copyin(block1[0:SIZE])
+#endif
+
+#pragma acc data copy(block1[0:SIZE])
+ {
+#ifdef OPENACC_API
+ acc_copyout (block1, SIZE);
+ acc_copyin (block1, SIZE);
+#else
+#pragma acc exit data copyout(block1[0:SIZE])
+#pragma acc enter data copyin(block1[0:SIZE])
+#endif
+ assert (acc_is_present (block1, SIZE));
+ }
+
+ assert (acc_is_present (block1, SIZE));
+#ifdef OPENACC_API
+ acc_copyout (block1, SIZE);
+#else
+#pragma acc exit data copyout(block1[0:SIZE])
+#endif
+ assert (!acc_is_present (block1, SIZE));
+
+ free (block1);
+}
+
+void
+f4 (void)
+{
+ char *block1 = (char *) malloc (SIZE);
+ char *block2 = (char *) malloc (SIZE);
+ char *block3 = (char *) malloc (SIZE);
+
+#pragma acc data copy(block1[0:SIZE], block2[0:SIZE], block3[0:SIZE])
+ {
+ /* The first copyin of block2 is the enclosing data region. This
+ "enter data" should make it live beyond the end of this region.
+ This works, though the on-target copies of block1, block2 and block3
+ will stay allocated until block2 is unmapped because they are bound
+ together in a single target_mem_desc. */
+#ifdef OPENACC_API
+ acc_copyin (block2, SIZE);
+#else
+#pragma acc enter data copyin(block2[0:SIZE])
+#endif
+ }
+
+ assert (!acc_is_present (block1, SIZE));
+ assert (acc_is_present (block2, SIZE));
+ assert (!acc_is_present (block3, SIZE));
+
+#ifdef OPENACC_API
+ acc_copyout (block2, SIZE);
+#else
+#pragma acc exit data copyout(block2[0:SIZE])
+#endif
+ assert (!acc_is_present (block2, SIZE));
+
+ free (block1);
+ free (block2);
+ free (block3);
+}
+
+int
+main (int argc, char *argv[])
+{
+ f1 ();
+ f2 ();
+ f3 ();
+ f4 ();
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-6-lib.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-6-lib.c
new file mode 100644
index 00000000000..8507a0586a5
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-6-lib.c
@@ -0,0 +1,5 @@
+/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
+/* { dg-additional-options "-DOPENACC_API" } */
+#include "static-dynamic-lifetimes-6.c"
+/* { dg-output "libgomp: cannot handle delete/copyout within data region" } */
+/* { dg-shouldfail "" } */
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-6.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-6.c
new file mode 100644
index 00000000000..ca3b385fbcc
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-6.c
@@ -0,0 +1,46 @@
+/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
+
+#include <openacc.h>
+#include <assert.h>
+#include <stdlib.h>
+
+#define SIZE 1024
+
+int
+main (int argc, char *argv[])
+{
+ char *block1 = (char *) malloc (SIZE);
+ char *block2 = (char *) malloc (SIZE);
+
+#ifdef OPENACC_API
+ acc_copyin (block1, SIZE);
+ acc_copyin (block2, SIZE);
+#else
+#pragma acc enter data copyin(block1[0:SIZE], block2[0:SIZE])
+#endif
+
+#pragma acc data copy(block1[0:SIZE], block2[0:SIZE])
+ {
+#ifdef OPENACC_API
+ acc_copyout (block1, SIZE);
+ acc_copyout (block2, SIZE);
+ /* Error output checked in static-dynamic-lifetimes-6-lib.c. */
+#else
+#pragma acc exit data copyout(block1[0:SIZE], block2[0:SIZE])
+/* We can only do this for a single dynamic data mapping at present. */
+/* { dg-output "libgomp: cannot handle .exit data. within data region" } */
+/* { dg-shouldfail "" } */
+#endif
+ /* These should stay present until the end of the static data lifetime. */
+ assert (acc_is_present (block1, SIZE));
+ assert (acc_is_present (block2, SIZE));
+ }
+
+ assert (!acc_is_present (block1, SIZE));
+ assert (!acc_is_present (block2, SIZE));
+
+ free (block1);
+ free (block2);
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-7-lib.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-7-lib.c
new file mode 100644
index 00000000000..962b5926f79
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-7-lib.c
@@ -0,0 +1,5 @@
+/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
+/* { dg-additional-options "-DOPENACC_API" } */
+#include "static-dynamic-lifetimes-7.c"
+/* { dg-output "libgomp: cannot handle delete/copyout within data region" } */
+/* { dg-shouldfail "" } */
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-7.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-7.c
new file mode 100644
index 00000000000..dfcc7cae961
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-7.c
@@ -0,0 +1,45 @@
+/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
+
+#include <openacc.h>
+#include <assert.h>
+#include <stdlib.h>
+
+#define SIZE 1024
+
+int
+main (int argc, char *argv[])
+{
+ char *block1 = (char *) malloc (SIZE);
+ char *block2 = (char *) malloc (SIZE);
+
+#ifdef OPENACC_API
+ acc_copyin (block1, SIZE);
+#else
+#pragma acc enter data copyin(block1[0:SIZE])
+#endif
+
+#pragma acc data copy(block1[0:SIZE], block2[0:SIZE])
+ {
+/* We can't attach the dynamic data mapping's (block1) target_mem_desc to the
+ enclosing static data region here, because that region maps block2 also. */
+#ifdef OPENACC_API
+ acc_copyout (block1, SIZE);
+ /* Error output checked in static-dynamic-lifetimes-7-lib.c. */
+#else
+#pragma acc exit data copyout(block1[0:SIZE])
+/* { dg-output "libgomp: cannot handle .exit data. within data region" } */
+/* { dg-shouldfail "" } */
+#endif
+ /* These should stay present until the end of the static data lifetime. */
+ assert (acc_is_present (block1, SIZE));
+ assert (acc_is_present (block2, SIZE));
+ }
+
+ assert (!acc_is_present (block1, SIZE));
+ assert (!acc_is_present (block2, SIZE));
+
+ free (block1);
+ free (block2);
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-8-lib.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-8-lib.c
new file mode 100644
index 00000000000..2581d7e2559
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-8-lib.c
@@ -0,0 +1,5 @@
+/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
+/* { dg-additional-options "-DOPENACC_API" } */
+#include "static-dynamic-lifetimes-8.c"
+/* { dg-output "libgomp: cannot handle delete/copyout within data region" } */
+/* { dg-shouldfail "" } */
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-8.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-8.c
new file mode 100644
index 00000000000..e3a64399fe9
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-8.c
@@ -0,0 +1,50 @@
+/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
+
+#include <openacc.h>
+#include <assert.h>
+#include <stdlib.h>
+
+#define SIZE 1024
+
+int
+main (int argc, char *argv[])
+{
+ char *block1 = (char *) malloc (SIZE);
+ char *block2 = (char *) malloc (SIZE);
+
+#ifdef OPENACC_API
+ acc_copyin (block1, SIZE);
+#else
+#pragma acc enter data copyin(block1[0:SIZE])
+#endif
+
+#pragma acc data copy(block1[0:SIZE], block2[0:SIZE])
+ {
+#ifdef OPENACC_API
+ acc_copyout (block1, SIZE);
+ acc_copyin (block2, SIZE);
+ /* Error output checked in static-dynamic-lifetimes-8-lib.c. */
+#else
+#pragma acc exit data copyout(block1[0:SIZE])
+/* { dg-output "libgomp: cannot handle .exit data. within data region" } */
+/* { dg-shouldfail "" } */
+#pragma acc enter data copyin(block2[0:SIZE])
+#endif
+ assert (acc_is_present (block1, SIZE));
+ assert (acc_is_present (block2, SIZE));
+ }
+
+ assert (!acc_is_present (block1, SIZE));
+ assert (acc_is_present (block2, SIZE));
+#ifdef OPENACC_API
+ acc_copyout (block2, SIZE);
+#else
+#pragma acc exit data copyout(block2[0:SIZE])
+#endif
+ assert (!acc_is_present (block2, SIZE));
+
+ free (block1);
+ free (block2);
+
+ return 0;
+}
--
2.23.0
More information about the Gcc-patches
mailing list