* [PATCH 0/3] Mixed static/dynamic data lifetimes with OpenACC (PR92843)
@ 2020-01-17 21:18 Julian Brown
2020-01-17 21:20 ` [PATCH 2/3] Don't copy back vars mapped with acc_map_data Julian Brown
` (3 more replies)
0 siblings, 4 replies; 8+ messages in thread
From: Julian Brown @ 2020-01-17 21:18 UTC (permalink / raw)
To: gcc-patches; +Cc: Thomas Schwinge, Jakub Jelinek
This patch series provides fixes for some cases of mixing static and
dynamic data lifetimes in OpenACC, hopefully addressing some of Thomas's
concerns in PR92843 -- in particular that an "exit data"-type operation on
a given variable inside a structured block (also mapping that variable)
should be a no-op.
On further investigation of related patterns, other cases of mixing static
and dynamic lifetimes also turn out to problematic at present. Some of
these cases are handled by this patch series, and others are diagnosed
as errors (rather than allowing silent and hard-to-diagnose failures
later during runtime).
The first two patches are sufficient to fix the test case introduced
for PR92843, and the third patch provides further support/diagnostics
for dynamic unmapping operations taking place within structured blocks.
Further commentary provided alongside individual patches. Tested with
offloading to NVPTX, also with a version of my refcount-verification patch
(not currently on trunk).
I believe this (at least the first two parts) fixes a regression (for
the pr92843-1.c test case), so OK for stage 4?
Thanks,
Julian
Julian Brown (3):
Introduce dynamic data mapping sentinel for OpenACC
Don't copy back vars mapped with acc_map_data
OpenACC dynamic data lifetimes ending within structured blocks
libgomp/libgomp.h | 3 +-
libgomp/oacc-int.h | 2 +
libgomp/oacc-mem.c | 158 ++++++++++++---
libgomp/target.c | 56 +++++-
.../libgomp.oacc-c-c++-common/pr92843-1.c | 4 +-
.../static-dynamic-lifetimes-1-lib.c | 3 +
.../static-dynamic-lifetimes-1.c | 160 +++++++++++++++
.../static-dynamic-lifetimes-2-lib.c | 3 +
.../static-dynamic-lifetimes-2.c | 166 ++++++++++++++++
.../static-dynamic-lifetimes-3-lib.c | 3 +
.../static-dynamic-lifetimes-3.c | 183 ++++++++++++++++++
.../static-dynamic-lifetimes-4-lib.c | 6 +
.../static-dynamic-lifetimes-4.c | 71 +++++++
.../static-dynamic-lifetimes-5-lib.c | 6 +
.../static-dynamic-lifetimes-5.c | 63 ++++++
.../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 +++++
21 files changed, 1000 insertions(+), 43 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-2-lib.c
create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-2.c
create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-3-lib.c
create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-3.c
create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-4-lib.c
create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-4.c
create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-5-lib.c
create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-5.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
--
2.23.0
^ permalink raw reply [flat|nested] 8+ messages in thread
* [PATCH 2/3] Don't copy back vars mapped with acc_map_data
2020-01-17 21:18 [PATCH 0/3] Mixed static/dynamic data lifetimes with OpenACC (PR92843) Julian Brown
@ 2020-01-17 21:20 ` Julian Brown
2020-07-03 15:53 ` Thomas Schwinge
2020-01-17 21:31 ` [PATCH 3/3] OpenACC dynamic data lifetimes ending within structured blocks Julian Brown
` (2 subsequent siblings)
3 siblings, 1 reply; 8+ messages in thread
From: Julian Brown @ 2020-01-17 21:20 UTC (permalink / raw)
To: gcc-patches; +Cc: Thomas Schwinge, Jakub Jelinek
This patch prevents "exit data" directives from copying back data that
was mapped with an acc_map_data API call. This matches the behaviour
expected by the pr92843-1.c test, and together with the previous patch
in this series, allows that test to pass (with no other regressions).
Tested alongside other patches in this series with offloading to NVPTX
(with and without the third & final patch).
OK?
Thanks,
Julian
ChangeLog
PR libgomp/92843
libgomp/
* oacc-mem.c (goacc_exit_data_internal): Don't copy-back data mapped
with acc_map_data on an "exit data" directive.
* testsuite/libgomp.oacc-c-c++-common/pr92843-1.c: Remove XFAIL. Add
explanatory comment.
---
libgomp/oacc-mem.c | 1 +
libgomp/testsuite/libgomp.oacc-c-c++-common/pr92843-1.c | 4 +++-
2 files changed, 4 insertions(+), 1 deletion(-)
diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c
index 45ab2b169d7..783e7f363fb 100644
--- a/libgomp/oacc-mem.c
+++ b/libgomp/oacc-mem.c
@@ -1235,6 +1235,7 @@ goacc_exit_data_internal (struct gomp_device_descr *acc_dev, size_t mapnum,
n->refcount--;
if (copyfrom
+ && n->refcount != REFCOUNT_INFINITY
&& (kind != GOMP_MAP_FROM || n->refcount == 0))
gomp_copy_dev2host (acc_dev, aq, (void *) cur_node.host_start,
(void *) (n->tgt->tgt_start + n->tgt_offset
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr92843-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr92843-1.c
index f16c46a37bf..786a12a8504 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr92843-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr92843-1.c
@@ -1,7 +1,6 @@
/* Verify that 'acc_copyout' etc. is a no-op if there's still a structured
reference count. */
-/* { dg-xfail-run-if "TODO PR92843" { *-*-* } } */
/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */
#include <assert.h>
@@ -96,6 +95,9 @@ test_acc_map_data ()
verify_array (h, N, c1);
assign_array (h, N, c1);
+ /* Note that we're not expecting this (nor the copyouts below) to perform
+ an actual "finalize" or copyout since the data was mapped with
+ acc_map_data. */
#pragma acc exit data copyout (h[0:N]) finalize
assert (acc_is_present (h, N));
verify_array (h, N, c1);
--
2.23.0
^ permalink raw reply [flat|nested] 8+ messages in thread
* [PATCH 3/3] OpenACC dynamic data lifetimes ending within structured blocks
2020-01-17 21:18 [PATCH 0/3] Mixed static/dynamic data lifetimes with OpenACC (PR92843) Julian Brown
2020-01-17 21:20 ` [PATCH 2/3] Don't copy back vars mapped with acc_map_data Julian Brown
@ 2020-01-17 21:31 ` Julian Brown
2020-05-11 15:05 ` Thomas Schwinge
2020-01-17 21:31 ` [PATCH 1/3] Introduce dynamic data mapping sentinel for OpenACC Julian Brown
2020-04-10 14:44 ` Test cases for mixed structured/dynamic data lifetimes with OpenACC [PR92843] (was: [PATCH 0/3] Mixed static/dynamic data lifetimes with OpenACC (PR92843)) Thomas Schwinge
3 siblings, 1 reply; 8+ messages in thread
From: Julian Brown @ 2020-01-17 21:31 UTC (permalink / raw)
To: gcc-patches; +Cc: Thomas Schwinge, Jakub Jelinek
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
^ permalink raw reply [flat|nested] 8+ messages in thread
* [PATCH 1/3] Introduce dynamic data mapping sentinel for OpenACC
2020-01-17 21:18 [PATCH 0/3] Mixed static/dynamic data lifetimes with OpenACC (PR92843) Julian Brown
2020-01-17 21:20 ` [PATCH 2/3] Don't copy back vars mapped with acc_map_data Julian Brown
2020-01-17 21:31 ` [PATCH 3/3] OpenACC dynamic data lifetimes ending within structured blocks Julian Brown
@ 2020-01-17 21:31 ` Julian Brown
2020-04-10 14:44 ` Test cases for mixed structured/dynamic data lifetimes with OpenACC [PR92843] (was: [PATCH 0/3] Mixed static/dynamic data lifetimes with OpenACC (PR92843)) Thomas Schwinge
3 siblings, 0 replies; 8+ messages in thread
From: Julian Brown @ 2020-01-17 21:31 UTC (permalink / raw)
To: gcc-patches; +Cc: Thomas Schwinge, Jakub Jelinek
This patch provides a way to distinguish target_mem_descs that arise from
"enter data" operations from those that arise from structured OpenACC
data blocks. In that way, we can implement the equivalent of the "no-op"
behaviour of decrementing a dynamic reference count that is already zero
for some given variable, as described in the OpenACC 2.6 spec.
We do this by re-using the "prev" field of the target_mem_desc (currently
unused for dynamic data mappings) to store a special sentinel value.
Several new tests are added, both for cases that now work, and for
diagnostics for cases that do not. Tested alongside other patches in
this series with offloading to NVPTX.
OK?
Thanks,
Julian
ChangeLog
PR libgomp/92843
libgomp/
* libgomp.h (target_mem_desc): Update comment for prev field.
* oacc-int.h (goacc_mark_dynamic): Add prototype.
* oacc-mem.c (dyn_tgt_sentinel): New static global.
(goacc_mark_dynamic): New function.
(goacc_enter_datum, goacc_enter_data_internal): Call goacc_mark_dynamic
on non-NULL target_mem_desc return from gomp_map_vars_async.
(goacc_exit_datum, goacc_exit_data_internal): Check target_mem_desc for
sentinel value on structural refcount decrement.
* target.c (gomp_unmap_vars_internal): Re-use target_mem_desc for
"structural" data mapping for extending dynamic mapping beyond the end
of a structured block when possible.
* testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-2.c: New
test.
* testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-2-lib.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-3.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-3-lib.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-4.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-4-lib.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-5.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-5-lib.c:
Likewise.
---
libgomp/libgomp.h | 3 +-
libgomp/oacc-int.h | 2 +
libgomp/oacc-mem.c | 53 ++++-
libgomp/target.c | 56 +++++-
.../static-dynamic-lifetimes-2-lib.c | 3 +
.../static-dynamic-lifetimes-2.c | 166 ++++++++++++++++
.../static-dynamic-lifetimes-3-lib.c | 3 +
.../static-dynamic-lifetimes-3.c | 183 ++++++++++++++++++
.../static-dynamic-lifetimes-4-lib.c | 6 +
.../static-dynamic-lifetimes-4.c | 71 +++++++
.../static-dynamic-lifetimes-5-lib.c | 6 +
.../static-dynamic-lifetimes-5.c | 63 ++++++
12 files changed, 596 insertions(+), 19 deletions(-)
create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-2-lib.c
create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-2.c
create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-3-lib.c
create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-3.c
create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-4-lib.c
create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-4.c
create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-5-lib.c
create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-5.c
diff --git a/libgomp/libgomp.h b/libgomp/libgomp.h
index 15a1394c16d..bbab4f9f34f 100644
--- a/libgomp/libgomp.h
+++ b/libgomp/libgomp.h
@@ -975,7 +975,8 @@ struct target_mem_desc {
uintptr_t tgt_end;
/* Handle to free. */
void *to_free;
- /* Previous target_mem_desc. */
+ /* Previous target_mem_desc. Also used in OpenACC to indicate that this
+ target_mem_desc is used only for an "enter data" mapping. */
struct target_mem_desc *prev;
/* Number of items in following list. */
size_t list_count;
diff --git a/libgomp/oacc-int.h b/libgomp/oacc-int.h
index 3c2c9b84b2f..bb67188c3e9 100644
--- a/libgomp/oacc-int.h
+++ b/libgomp/oacc-int.h
@@ -165,6 +165,8 @@ bool _goacc_profiling_setup_p (struct goacc_thread *,
void goacc_profiling_dispatch (acc_prof_info *, acc_event_info *,
acc_api_info *);
+extern void goacc_mark_dynamic (struct target_mem_desc *);
+
#ifdef HAVE_ATTRIBUTE_VISIBILITY
# pragma GCC visibility pop
#endif
diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c
index bd1a99d9277..45ab2b169d7 100644
--- a/libgomp/oacc-mem.c
+++ b/libgomp/oacc-mem.c
@@ -550,6 +550,24 @@ acc_unmap_data (void *h)
}
}
+/* Indicate (via storing its address in the "prev" field) a target_mem_desc
+ that is used for an "enter data" mapping. */
+const static struct target_mem_desc dyn_tgt_sentinel;
+
+/* Mark TGT as the "initial" target_mem_desc created by a dynamic data mapping
+ (acc_create, acc_copyin or an "enter data" directive). For such mappings,
+ to start with, we have a splay tree key with a reference count of 1 and a
+ virtual reference count of 0 (linking to this target_mem_desc). Without
+ this marking, such a mapping is indistinguishable from a target_mem_desc
+ created by e.g. a lexically-scoped "acc data" region, but the difference is
+ important if acc_copyout, acc_delete (etc.) or an "exit data" directive is
+ used to end the data lifetime. */
+
+void
+goacc_mark_dynamic (struct target_mem_desc *tgt)
+{
+ tgt->prev = (struct target_mem_desc *) &dyn_tgt_sentinel;
+}
/* Enter dynamic mapping for a single datum. Return the device pointer. */
@@ -613,8 +631,14 @@ goacc_enter_datum (void **hostaddrs, size_t *sizes, void *kinds, int async)
goacc_aq aq = get_goacc_asyncqueue (async);
- gomp_map_vars_async (acc_dev, aq, mapnum, hostaddrs, NULL, sizes, kinds,
- true, GOMP_MAP_VARS_OPENACC_ENTER_DATA);
+ struct target_mem_desc *tgt;
+ tgt = gomp_map_vars_async (acc_dev, aq, mapnum, hostaddrs, NULL, sizes,
+ kinds, true, GOMP_MAP_VARS_OPENACC_ENTER_DATA);
+
+ /* Mark non-NULL target_mem_descs returned here specially: see comment in
+ goacc_exit_datum. */
+ if (tgt)
+ goacc_mark_dynamic (tgt);
gomp_mutex_lock (&acc_dev->lock);
n = lookup_host (acc_dev, hostaddrs[0], sizes[0]);
@@ -756,7 +780,15 @@ goacc_exit_datum (void *h, size_t s, unsigned short kind, int async)
n->refcount--;
n->virtual_refcount--;
}
- else if (n->refcount > 0 && n->refcount != REFCOUNT_INFINITY)
+ /* 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)
@@ -1081,11 +1113,12 @@ goacc_enter_data_internal (struct gomp_device_descr *acc_dev, size_t mapnum,
dump_mappings ((group_last - i) + 1, &hostaddrs[i], &sizes[i], &kinds[i]);
#endif
- gomp_map_vars_async (acc_dev, aq,
- (group_last - i) + 1,
- &hostaddrs[i], NULL,
- &sizes[i], &kinds[i], true,
- GOMP_MAP_VARS_OPENACC_ENTER_DATA);
+ struct target_mem_desc *tgt;
+ tgt = gomp_map_vars_async (acc_dev, aq, (group_last - i) + 1,
+ &hostaddrs[i], NULL, &sizes[i], &kinds[i],
+ true, GOMP_MAP_VARS_OPENACC_ENTER_DATA);
+ if (tgt)
+ goacc_mark_dynamic (tgt);
i = group_last;
}
@@ -1196,7 +1229,9 @@ goacc_exit_data_internal (struct gomp_device_descr *acc_dev, size_t mapnum,
n->refcount--;
n->virtual_refcount--;
}
- else if (n->refcount > 0 && n->refcount != REFCOUNT_INFINITY)
+ else if (n->refcount > 0
+ && n->refcount != REFCOUNT_INFINITY
+ && n->tgt->prev == &dyn_tgt_sentinel)
n->refcount--;
if (copyfrom
diff --git a/libgomp/target.c b/libgomp/target.c
index 825213f40ec..fb423ced144 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -1629,6 +1629,8 @@ gomp_unmap_vars_internal (struct target_mem_desc *tgt, bool do_copyfrom,
k->refcount == 1, NULL);
}
+ bool have_virtual_refs = false, all_refs_virtual = true;
+
for (i = 0; i < tgt->list_count; i++)
{
splay_tree_key k = tgt->list[i].key;
@@ -1636,21 +1638,21 @@ gomp_unmap_vars_internal (struct target_mem_desc *tgt, bool do_copyfrom,
continue;
bool do_unmap = false;
- if (k->tgt == tgt
- && k->virtual_refcount > 0
- && k->refcount != REFCOUNT_INFINITY)
- {
- k->virtual_refcount--;
- k->refcount--;
- }
- else if (k->refcount > 1 && k->refcount != REFCOUNT_INFINITY)
+ if (k->refcount > 1 && k->refcount != REFCOUNT_INFINITY)
k->refcount--;
else if (k->refcount == 1)
{
k->refcount--;
- do_unmap = true;
+ if (k->virtual_refcount == 0)
+ do_unmap = true;
}
+ if (k->virtual_refcount > 0 && k->refcount == k->virtual_refcount)
+ have_virtual_refs = true;
+
+ if (k->refcount != k->virtual_refcount)
+ all_refs_virtual = false;
+
if ((do_unmap && do_copyfrom && tgt->list[i].copy_from)
|| tgt->list[i].always_copy_from)
gomp_copy_dev2host (devicep, aq,
@@ -1670,6 +1672,42 @@ gomp_unmap_vars_internal (struct target_mem_desc *tgt, bool do_copyfrom,
}
}
+ if (have_virtual_refs)
+ {
+ /* If we have a construct such as this:
+
+ #pragma acc data copy(var1)
+ {
+ #pragma acc enter data copyin(var1)
+ }
+
+ The dynamic data lifetime entered in the middle of the static
+ data lifetime extends beyond the static lifetime. Adjust
+ references and the target descriptor here (the end of the static
+ region) to make it seem like we did "enter data" on the data to
+ start with.
+
+ We can't do this adjustment if the data construct refers to other
+ variables too. */
+ if (!all_refs_virtual)
+ {
+ gomp_mutex_unlock (&devicep->lock);
+ gomp_fatal ("cannot handle create/copyin/'enter data' within data "
+ "region");
+ }
+
+ for (i = 0; i < tgt->list_count; i++)
+ {
+ splay_tree_key k = tgt->list[i].key;
+ if (k == NULL || k->virtual_refcount == 0)
+ continue;
+
+ if (k->refcount == k->virtual_refcount)
+ k->virtual_refcount--;
+ }
+ goacc_mark_dynamic (tgt);
+ }
+
if (aq)
devicep->openacc.async.queue_callback_func (aq, gomp_unref_tgt_void,
(void *) tgt);
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-2-lib.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-2-lib.c
new file mode 100644
index 00000000000..84f41a49dfd
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-2-lib.c
@@ -0,0 +1,3 @@
+/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
+/* { dg-additional-options "-DOPENACC_API" } */
+#include "static-dynamic-lifetimes-2.c"
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-2.c
new file mode 100644
index 00000000000..d3c6f5192d8
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-2.c
@@ -0,0 +1,166 @@
+/* Test nested dynamic/static data mappings. */
+
+/* { 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);
+
+#pragma acc data copy(block1[0:SIZE])
+ {
+#ifdef OPENACC_API
+ acc_copyin (block1, SIZE);
+ acc_copyout (block1, SIZE);
+#else
+#pragma acc enter data copyin(block1[0:SIZE])
+#pragma acc exit data copyout(block1[0:SIZE])
+#endif
+ }
+
+ assert (!acc_is_present (block1, SIZE));
+
+ 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
+
+ assert (!acc_is_present (block1, SIZE));
+
+ free (block1);
+}
+
+void
+f3 (void)
+{
+ char *block1 = (char *) malloc (SIZE);
+
+#pragma acc data copy(block1[0:SIZE])
+ {
+#ifdef OPENACC_API
+ acc_copyin (block1, SIZE);
+ acc_copyin (block1, SIZE);
+ acc_copyout (block1, SIZE);
+ acc_copyout (block1, SIZE);
+#else
+#pragma acc enter data copyin(block1[0:SIZE])
+#pragma acc enter data copyin(block1[0:SIZE])
+#pragma acc exit data copyout(block1[0:SIZE])
+#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);
+
+#pragma acc data copy(block1[0: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_copyin (block1, SIZE);
+ acc_copyout (block1, SIZE);
+#else
+#pragma acc enter data copyin(block1[0:SIZE])
+#pragma acc exit data copyout(block1[0:SIZE])
+#endif
+ }
+
+#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
+f5 (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_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
+ }
+#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);
+}
+
+int
+main (int argc, char *argv[])
+{
+ f1 ();
+ f2 ();
+ f3 ();
+ f4 ();
+ f5 ();
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-3-lib.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-3-lib.c
new file mode 100644
index 00000000000..d9e76c600f0
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-3-lib.c
@@ -0,0 +1,3 @@
+/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
+/* { dg-additional-options "-DOPENACC_API" } */
+#include "static-dynamic-lifetimes-3.c"
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-3.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-3.c
new file mode 100644
index 00000000000..59501864398
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-3.c
@@ -0,0 +1,183 @@
+/* Test nested dynamic/static data mappings (multiple blocks on data
+ regions). */
+
+/* { 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);
+ char *block2 = (char *) malloc (SIZE);
+
+#pragma acc data copy(block1[0:SIZE], block2[0:SIZE])
+ {
+#ifdef OPENACC_API
+ acc_copyin (block1, SIZE);
+ acc_copyout (block1, SIZE);
+#else
+#pragma acc enter data copyin(block1[0:SIZE])
+#pragma acc exit data copyout(block1[0:SIZE])
+#endif
+ }
+
+ assert (!acc_is_present (block1, SIZE));
+ assert (!acc_is_present (block2, SIZE));
+
+ free (block1);
+ free (block2);
+}
+
+void
+f2 (void)
+{
+ 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);
+#else
+#pragma acc exit data copyout(block1[0:SIZE])
+#endif
+
+ assert (!acc_is_present (block1, SIZE));
+ assert (!acc_is_present (block2, SIZE));
+
+ free (block1);
+ free (block2);
+}
+
+void
+f3 (void)
+{
+ char *block1 = (char *) malloc (SIZE);
+ char *block2 = (char *) malloc (SIZE);
+
+#pragma acc data copy(block1[0:SIZE], block2[0:SIZE])
+ {
+#ifdef OPENACC_API
+ acc_copyin (block1, SIZE);
+ acc_copyin (block2, SIZE);
+ acc_copyout (block2, SIZE);
+ acc_copyout (block1, SIZE);
+#else
+#pragma acc enter data copyin(block1[0:SIZE])
+#pragma acc enter data copyin(block2[0:SIZE])
+#pragma acc exit data copyout(block2[0:SIZE])
+#pragma acc exit data copyout(block1[0:SIZE])
+#endif
+ }
+
+ assert (!acc_is_present (block1, SIZE));
+ assert (!acc_is_present (block2, SIZE));
+
+ free (block1);
+ free (block2);
+}
+
+void
+f4 (void)
+{
+ char *block1 = (char *) malloc (SIZE);
+ char *block2 = (char *) malloc (SIZE);
+
+#pragma acc data copy(block1[0:SIZE], block2[0: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_copyin (block2, SIZE);
+ acc_copyout (block2, SIZE);
+#else
+#pragma acc enter data copyin(block2[0:SIZE])
+#pragma acc exit data copyout(block2[0:SIZE])
+#endif
+ }
+#ifdef OPENACC_API
+ acc_copyout (block1, SIZE);
+#else
+#pragma acc exit data copyout(block1[0:SIZE])
+#endif
+ }
+
+ assert (!acc_is_present (block1, SIZE));
+ assert (!acc_is_present (block2, SIZE));
+
+ free (block1);
+ free (block2);
+}
+
+void
+f5 (void)
+{
+ 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_copyin (block2, SIZE);
+#else
+#pragma acc enter data copyin(block2[0:SIZE])
+#endif
+
+#pragma acc data copy(block1[0:SIZE], block2[0:SIZE])
+ {
+ }
+#ifdef OPENACC_API
+ acc_copyout (block2, SIZE);
+#else
+#pragma acc exit data copyout(block2[0:SIZE])
+#endif
+ }
+
+#ifdef OPENACC_API
+ acc_copyout (block1, SIZE);
+#else
+#pragma acc exit data copyout(block1[0:SIZE])
+#endif
+
+ assert (!acc_is_present (block1, SIZE));
+ assert (!acc_is_present (block2, SIZE));
+
+ free (block1);
+ free (block2);
+}
+
+int
+main (int argc, char *argv[])
+{
+ f1 ();
+ f2 ();
+ f3 ();
+ f4 ();
+ f5 ();
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-4-lib.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-4-lib.c
new file mode 100644
index 00000000000..77bcd9e8dd8
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-4-lib.c
@@ -0,0 +1,6 @@
+/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
+/* { dg-additional-options "-DOPENACC_API" } */
+#include "static-dynamic-lifetimes-4.c"
+
+/* { dg-output "libgomp: cannot handle create/copyin/.enter data. within data region" } */
+/* { dg-shouldfail "" } */
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-4.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-4.c
new file mode 100644
index 00000000000..0d9f52febdb
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-4.c
@@ -0,0 +1,71 @@
+/* { 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);
+ char *block3 = (char *) malloc (SIZE);
+
+ /* Doing this twice ensures that we have a non-zero virtual refcount. Make
+ sure that works too. */
+#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], 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. */
+#ifdef OPENACC_API
+ acc_copyin (block2, SIZE);
+ /* Error output checked in static-dynamic-lifetimes-4-lib.c. */
+#else
+#pragma acc enter data copyin(block2[0:SIZE])
+ /* ...except that doesn't work at present because it would mean the dynamic
+ data region would get entangled with the static data region's
+ target_mem_desc that has mappings for each of block1, block2 and block3.
+ Check for runtime error. */
+ /* { dg-output "libgomp: cannot handle create/copyin/.enter data. within data region" } */
+ /* { dg-shouldfail "" } */
+#endif
+ }
+
+ assert (acc_is_present (block1, SIZE));
+ assert (acc_is_present (block2, SIZE));
+ assert (!acc_is_present (block3, 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 (block2, SIZE);
+ assert (!acc_is_present (block2, 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(block2[0:SIZE])
+ assert (!acc_is_present (block2, SIZE));
+#endif
+
+ free (block1);
+ free (block2);
+ free (block3);
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-5-lib.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-5-lib.c
new file mode 100644
index 00000000000..dcf4da6b660
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-5-lib.c
@@ -0,0 +1,6 @@
+/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
+/* { dg-additional-options "-DOPENACC_API" } */
+#include "static-dynamic-lifetimes-5.c"
+
+/* { dg-output "libgomp: cannot handle create/copyin/.enter data. within data region" } */
+/* { dg-shouldfail "" } */
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-5.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-5.c
new file mode 100644
index 00000000000..062ca74f2ab
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-5.c
@@ -0,0 +1,63 @@
+/* { 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);
+ char *block3 = (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], 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. */
+#ifdef OPENACC_API
+ acc_copyin (block2, SIZE);
+ /* Error output checked in static-dynamic-lifetimes-5-lib.c. */
+#else
+#pragma acc enter data copyin(block2[0:SIZE])
+ /* ...except that doesn't work at present because it would mean the dynamic
+ data region would get entangled with the static data region's
+ target_mem_desc that has mappings for each of block1, block2 and block3.
+ Check for runtime error. */
+ /* { dg-output "libgomp: cannot handle create/copyin/.enter data. within data region" } */
+ /* { dg-shouldfail "" } */
+#endif
+ }
+
+ assert (acc_is_present (block1, SIZE));
+ assert (acc_is_present (block2, SIZE));
+ assert (!acc_is_present (block3, SIZE));
+
+#ifdef OPENACC_API
+ acc_copyout (block1, SIZE);
+ assert (!acc_is_present (block1, SIZE));
+
+ acc_copyout (block2, SIZE);
+ assert (!acc_is_present (block2, SIZE));
+#else
+#pragma acc exit data copyout(block1[0:SIZE])
+ assert (!acc_is_present (block1, SIZE));
+
+#pragma acc exit data copyout(block2[0:SIZE])
+ assert (!acc_is_present (block2, SIZE));
+#endif
+
+ free (block1);
+ free (block2);
+ free (block3);
+
+ return 0;
+}
--
2.23.0
^ permalink raw reply [flat|nested] 8+ messages in thread
* Test cases for mixed structured/dynamic data lifetimes with OpenACC [PR92843] (was: [PATCH 0/3] Mixed static/dynamic data lifetimes with OpenACC (PR92843))
2020-01-17 21:18 [PATCH 0/3] Mixed static/dynamic data lifetimes with OpenACC (PR92843) Julian Brown
` (2 preceding siblings ...)
2020-01-17 21:31 ` [PATCH 1/3] Introduce dynamic data mapping sentinel for OpenACC Julian Brown
@ 2020-04-10 14:44 ` Thomas Schwinge
2020-04-13 7:23 ` Test cases for mixed structured/dynamic data lifetimes with OpenACC [PR92843] Thomas Schwinge
3 siblings, 1 reply; 8+ messages in thread
From: Thomas Schwinge @ 2020-04-10 14:44 UTC (permalink / raw)
To: Julian Brown, gcc-patches; +Cc: Jakub Jelinek
[-- Attachment #1: Type: text/plain, Size: 2339 bytes --]
Hi!
On 2020-01-17T12:18:18-0800, Julian Brown <julian@codesourcery.com> wrote:
> This patch series provides fixes for some cases of mixing static and
(It's "structured", not "static".) ;-)
> dynamic data lifetimes in OpenACC, hopefully addressing some of Thomas's
> concerns in PR92843 -- in particular that an "exit data"-type operation on
> a given variable inside a structured block (also mapping that variable)
> should be a no-op.
>
> On further investigation of related patterns, other cases of mixing static
> and dynamic lifetimes also turn out to problematic at present. Some of
> these cases are handled by this patch series, and others are diagnosed
> as errors (rather than allowing silent and hard-to-diagnose failures
> later during runtime).
>
> The first two patches are sufficient to fix the test case introduced
> for PR92843, and the third patch provides further support/diagnostics
> for dynamic unmapping operations taking place within structured blocks.
>
> Further commentary provided alongside individual patches. Tested with
> offloading to NVPTX, also with a version of my refcount-verification patch
> (not currently on trunk).
>
> I believe this (at least the first two parts) fixes a regression (for
> the pr92843-1.c test case), so OK for stage 4?
Thanks for your continued work here.
The code changes will need further review and discussion, but the thing
is: the test cases you provided already PASS now (on master branch, and
also on releases/gcc-9 branch, which is the first branch to implement
these semantics), so I now already pushed these test cases to master
branch in commit be9862dd96945772ae0692bc95b37ec6dbcabda0 "Test cases for
mixed structured/dynamic data lifetimes with OpenACC [PR92843]", and to
releases/gcc-9 branch in commit 3c7a476c5ad3761cb5373f8c59a92e04525c5638
"Test cases for mixed structured/dynamic data lifetimes with OpenACC
[PR92843]", see attached. (That's with all XFAILs removed, meaning that
the XFAILs will need to be re-instantiated with the code changes that
actually break the respective functionality.)
Grüße
Thomas
-----------------
Mentor Graphics (Deutschland) GmbH, Arnulfstraße 201, 80634 München / Germany
Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Alexander Walter
[-- Warning: decoded text below may be mangled, UTF-8 assumed --]
[-- Attachment #2: 0001-Test-cases-for-mixed-structured-dynamic-data-lifetim.patch --]
[-- Type: text/x-diff, Size: 28743 bytes --]
From be9862dd96945772ae0692bc95b37ec6dbcabda0 Mon Sep 17 00:00:00 2001
From: Julian Brown <julian@codesourcery.com>
Date: Fri, 17 Jan 2020 13:18:18 -0800
Subject: [PATCH] Test cases for mixed structured/dynamic data lifetimes with
OpenACC [PR92843]
libgomp/
PR libgomp/92843
* testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-1-lib.c:
New file.
* testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-1.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-2-lib.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-2.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-3-lib.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-3.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-4-lib.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-4.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-5-lib.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-5.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-6-lib.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-6.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-7-lib.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-7.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-8-lib.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-8.c:
Likewise.
---
libgomp/ChangeLog | 37 ++++
.../static-dynamic-lifetimes-1-lib.c | 3 +
.../static-dynamic-lifetimes-1.c | 160 +++++++++++++++
.../static-dynamic-lifetimes-2-lib.c | 3 +
.../static-dynamic-lifetimes-2.c | 166 ++++++++++++++++
.../static-dynamic-lifetimes-3-lib.c | 3 +
.../static-dynamic-lifetimes-3.c | 183 ++++++++++++++++++
.../static-dynamic-lifetimes-4-lib.c | 3 +
.../static-dynamic-lifetimes-4.c | 64 ++++++
.../static-dynamic-lifetimes-5-lib.c | 3 +
.../static-dynamic-lifetimes-5.c | 56 ++++++
.../static-dynamic-lifetimes-6-lib.c | 3 +
.../static-dynamic-lifetimes-6.c | 42 ++++
.../static-dynamic-lifetimes-7-lib.c | 3 +
.../static-dynamic-lifetimes-7.c | 42 ++++
.../static-dynamic-lifetimes-8-lib.c | 3 +
.../static-dynamic-lifetimes-8.c | 47 +++++
17 files changed, 821 insertions(+)
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-2-lib.c
create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-2.c
create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-3-lib.c
create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-3.c
create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-4-lib.c
create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-4.c
create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-5-lib.c
create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-5.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/ChangeLog b/libgomp/ChangeLog
index beff3d65b44..b0f19845ad2 100644
--- a/libgomp/ChangeLog
+++ b/libgomp/ChangeLog
@@ -1,3 +1,40 @@
+2020-04-10 Julian Brown <julian@codesourcery.com>
+ Thomas Schwinge <thomas@codesourcery.com>
+
+ PR libgomp/92843
+ * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-1-lib.c:
+ New file.
+ * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-1.c:
+ Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-2-lib.c:
+ Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-2.c:
+ Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-3-lib.c:
+ Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-3.c:
+ Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-4-lib.c:
+ Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-4.c:
+ Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-5-lib.c:
+ Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-5.c:
+ Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-6-lib.c:
+ Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-6.c:
+ Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-7-lib.c:
+ Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-7.c:
+ Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-8-lib.c:
+ Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-8.c:
+ Likewise.
+
2020-04-10 Thomas Schwinge <thomas@codesourcery.com>
* testsuite/libgomp.fortran/target-enter-data-1.f90: Add 'dg-do
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-2-lib.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-2-lib.c
new file mode 100644
index 00000000000..84f41a49dfd
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-2-lib.c
@@ -0,0 +1,3 @@
+/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
+/* { dg-additional-options "-DOPENACC_API" } */
+#include "static-dynamic-lifetimes-2.c"
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-2.c
new file mode 100644
index 00000000000..d3c6f5192d8
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-2.c
@@ -0,0 +1,166 @@
+/* Test nested dynamic/static data mappings. */
+
+/* { 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);
+
+#pragma acc data copy(block1[0:SIZE])
+ {
+#ifdef OPENACC_API
+ acc_copyin (block1, SIZE);
+ acc_copyout (block1, SIZE);
+#else
+#pragma acc enter data copyin(block1[0:SIZE])
+#pragma acc exit data copyout(block1[0:SIZE])
+#endif
+ }
+
+ assert (!acc_is_present (block1, SIZE));
+
+ 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
+
+ assert (!acc_is_present (block1, SIZE));
+
+ free (block1);
+}
+
+void
+f3 (void)
+{
+ char *block1 = (char *) malloc (SIZE);
+
+#pragma acc data copy(block1[0:SIZE])
+ {
+#ifdef OPENACC_API
+ acc_copyin (block1, SIZE);
+ acc_copyin (block1, SIZE);
+ acc_copyout (block1, SIZE);
+ acc_copyout (block1, SIZE);
+#else
+#pragma acc enter data copyin(block1[0:SIZE])
+#pragma acc enter data copyin(block1[0:SIZE])
+#pragma acc exit data copyout(block1[0:SIZE])
+#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);
+
+#pragma acc data copy(block1[0: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_copyin (block1, SIZE);
+ acc_copyout (block1, SIZE);
+#else
+#pragma acc enter data copyin(block1[0:SIZE])
+#pragma acc exit data copyout(block1[0:SIZE])
+#endif
+ }
+
+#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
+f5 (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_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
+ }
+#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);
+}
+
+int
+main (int argc, char *argv[])
+{
+ f1 ();
+ f2 ();
+ f3 ();
+ f4 ();
+ f5 ();
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-3-lib.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-3-lib.c
new file mode 100644
index 00000000000..d9e76c600f0
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-3-lib.c
@@ -0,0 +1,3 @@
+/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
+/* { dg-additional-options "-DOPENACC_API" } */
+#include "static-dynamic-lifetimes-3.c"
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-3.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-3.c
new file mode 100644
index 00000000000..59501864398
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-3.c
@@ -0,0 +1,183 @@
+/* Test nested dynamic/static data mappings (multiple blocks on data
+ regions). */
+
+/* { 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);
+ char *block2 = (char *) malloc (SIZE);
+
+#pragma acc data copy(block1[0:SIZE], block2[0:SIZE])
+ {
+#ifdef OPENACC_API
+ acc_copyin (block1, SIZE);
+ acc_copyout (block1, SIZE);
+#else
+#pragma acc enter data copyin(block1[0:SIZE])
+#pragma acc exit data copyout(block1[0:SIZE])
+#endif
+ }
+
+ assert (!acc_is_present (block1, SIZE));
+ assert (!acc_is_present (block2, SIZE));
+
+ free (block1);
+ free (block2);
+}
+
+void
+f2 (void)
+{
+ 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);
+#else
+#pragma acc exit data copyout(block1[0:SIZE])
+#endif
+
+ assert (!acc_is_present (block1, SIZE));
+ assert (!acc_is_present (block2, SIZE));
+
+ free (block1);
+ free (block2);
+}
+
+void
+f3 (void)
+{
+ char *block1 = (char *) malloc (SIZE);
+ char *block2 = (char *) malloc (SIZE);
+
+#pragma acc data copy(block1[0:SIZE], block2[0:SIZE])
+ {
+#ifdef OPENACC_API
+ acc_copyin (block1, SIZE);
+ acc_copyin (block2, SIZE);
+ acc_copyout (block2, SIZE);
+ acc_copyout (block1, SIZE);
+#else
+#pragma acc enter data copyin(block1[0:SIZE])
+#pragma acc enter data copyin(block2[0:SIZE])
+#pragma acc exit data copyout(block2[0:SIZE])
+#pragma acc exit data copyout(block1[0:SIZE])
+#endif
+ }
+
+ assert (!acc_is_present (block1, SIZE));
+ assert (!acc_is_present (block2, SIZE));
+
+ free (block1);
+ free (block2);
+}
+
+void
+f4 (void)
+{
+ char *block1 = (char *) malloc (SIZE);
+ char *block2 = (char *) malloc (SIZE);
+
+#pragma acc data copy(block1[0:SIZE], block2[0: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_copyin (block2, SIZE);
+ acc_copyout (block2, SIZE);
+#else
+#pragma acc enter data copyin(block2[0:SIZE])
+#pragma acc exit data copyout(block2[0:SIZE])
+#endif
+ }
+#ifdef OPENACC_API
+ acc_copyout (block1, SIZE);
+#else
+#pragma acc exit data copyout(block1[0:SIZE])
+#endif
+ }
+
+ assert (!acc_is_present (block1, SIZE));
+ assert (!acc_is_present (block2, SIZE));
+
+ free (block1);
+ free (block2);
+}
+
+void
+f5 (void)
+{
+ 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_copyin (block2, SIZE);
+#else
+#pragma acc enter data copyin(block2[0:SIZE])
+#endif
+
+#pragma acc data copy(block1[0:SIZE], block2[0:SIZE])
+ {
+ }
+#ifdef OPENACC_API
+ acc_copyout (block2, SIZE);
+#else
+#pragma acc exit data copyout(block2[0:SIZE])
+#endif
+ }
+
+#ifdef OPENACC_API
+ acc_copyout (block1, SIZE);
+#else
+#pragma acc exit data copyout(block1[0:SIZE])
+#endif
+
+ assert (!acc_is_present (block1, SIZE));
+ assert (!acc_is_present (block2, SIZE));
+
+ free (block1);
+ free (block2);
+}
+
+int
+main (int argc, char *argv[])
+{
+ f1 ();
+ f2 ();
+ f3 ();
+ f4 ();
+ f5 ();
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-4-lib.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-4-lib.c
new file mode 100644
index 00000000000..e3c1bfb473d
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-4-lib.c
@@ -0,0 +1,3 @@
+/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
+/* { dg-additional-options "-DOPENACC_API" } */
+#include "static-dynamic-lifetimes-4.c"
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-4.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-4.c
new file mode 100644
index 00000000000..e9a6510ace8
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-4.c
@@ -0,0 +1,64 @@
+/* { 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);
+ char *block3 = (char *) malloc (SIZE);
+
+ /* Doing this twice ensures that we have a non-zero virtual refcount. Make
+ sure that works too. */
+#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], 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. */
+#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 (block1, SIZE);
+ assert (acc_is_present (block1, SIZE));
+ acc_copyout (block1, SIZE);
+ assert (!acc_is_present (block1, SIZE));
+
+ acc_copyout (block2, SIZE);
+ assert (!acc_is_present (block2, 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(block2[0:SIZE])
+ assert (!acc_is_present (block2, SIZE));
+#endif
+
+ free (block1);
+ free (block2);
+ free (block3);
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-5-lib.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-5-lib.c
new file mode 100644
index 00000000000..77703122ad6
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-5-lib.c
@@ -0,0 +1,3 @@
+/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
+/* { dg-additional-options "-DOPENACC_API" } */
+#include "static-dynamic-lifetimes-5.c"
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-5.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-5.c
new file mode 100644
index 00000000000..9807076d3f4
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-5.c
@@ -0,0 +1,56 @@
+/* { 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);
+ char *block3 = (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], 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. */
+#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 (block1, SIZE);
+ assert (!acc_is_present (block1, SIZE));
+
+ acc_copyout (block2, SIZE);
+ assert (!acc_is_present (block2, SIZE));
+#else
+#pragma acc exit data copyout(block1[0:SIZE])
+ assert (!acc_is_present (block1, SIZE));
+
+#pragma acc exit data copyout(block2[0:SIZE])
+ assert (!acc_is_present (block2, SIZE));
+#endif
+
+ free (block1);
+ free (block2);
+ free (block3);
+
+ 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..4a87dd72525
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-6-lib.c
@@ -0,0 +1,3 @@
+/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
+/* { dg-additional-options "-DOPENACC_API" } */
+#include "static-dynamic-lifetimes-6.c"
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..3e5c4d7ea11
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-6.c
@@ -0,0 +1,42 @@
+/* { 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);
+#else
+#pragma acc exit data copyout(block1[0:SIZE], block2[0:SIZE])
+#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..8ccbb126933
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-7-lib.c
@@ -0,0 +1,3 @@
+/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
+/* { dg-additional-options "-DOPENACC_API" } */
+#include "static-dynamic-lifetimes-7.c"
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..2735d6fa0eb
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-7.c
@@ -0,0 +1,42 @@
+/* { 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);
+#else
+#pragma acc exit data copyout(block1[0:SIZE])
+#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..f3104cbd035
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-8-lib.c
@@ -0,0 +1,3 @@
+/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
+/* { dg-additional-options "-DOPENACC_API" } */
+#include "static-dynamic-lifetimes-8.c"
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..919ee02b725
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-8.c
@@ -0,0 +1,47 @@
+/* { 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);
+#else
+#pragma acc exit data copyout(block1[0:SIZE])
+#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.17.1
[-- Warning: decoded text below may be mangled, UTF-8 assumed --]
[-- Attachment #3: 0001-Test-cases-for-mixed-structured-dynamic-data-life.g9.patch --]
[-- Type: text/x-diff, Size: 28765 bytes --]
From 3c7a476c5ad3761cb5373f8c59a92e04525c5638 Mon Sep 17 00:00:00 2001
From: Julian Brown <julian@codesourcery.com>
Date: Fri, 17 Jan 2020 13:18:18 -0800
Subject: [PATCH] Test cases for mixed structured/dynamic data lifetimes with
OpenACC [PR92843]
libgomp/
PR libgomp/92843
* testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-1-lib.c:
New file.
* testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-1.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-2-lib.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-2.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-3-lib.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-3.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-4-lib.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-4.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-5-lib.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-5.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-6-lib.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-6.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-7-lib.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-7.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-8-lib.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-8.c:
Likewise.
(cherry picked from commit be9862dd96945772ae0692bc95b37ec6dbcabda0)
---
libgomp/ChangeLog | 37 ++++
.../static-dynamic-lifetimes-1-lib.c | 3 +
.../static-dynamic-lifetimes-1.c | 160 +++++++++++++++
.../static-dynamic-lifetimes-2-lib.c | 3 +
.../static-dynamic-lifetimes-2.c | 166 ++++++++++++++++
.../static-dynamic-lifetimes-3-lib.c | 3 +
.../static-dynamic-lifetimes-3.c | 183 ++++++++++++++++++
.../static-dynamic-lifetimes-4-lib.c | 3 +
.../static-dynamic-lifetimes-4.c | 64 ++++++
.../static-dynamic-lifetimes-5-lib.c | 3 +
.../static-dynamic-lifetimes-5.c | 56 ++++++
.../static-dynamic-lifetimes-6-lib.c | 3 +
.../static-dynamic-lifetimes-6.c | 42 ++++
.../static-dynamic-lifetimes-7-lib.c | 3 +
.../static-dynamic-lifetimes-7.c | 42 ++++
.../static-dynamic-lifetimes-8-lib.c | 3 +
.../static-dynamic-lifetimes-8.c | 47 +++++
17 files changed, 821 insertions(+)
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-2-lib.c
create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-2.c
create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-3-lib.c
create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-3.c
create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-4-lib.c
create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-4.c
create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-5-lib.c
create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-5.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/ChangeLog b/libgomp/ChangeLog
index ee18a994590..39308ecbc03 100644
--- a/libgomp/ChangeLog
+++ b/libgomp/ChangeLog
@@ -1,3 +1,40 @@
+2020-04-10 Julian Brown <julian@codesourcery.com>
+ Thomas Schwinge <thomas@codesourcery.com>
+
+ PR libgomp/92843
+ * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-1-lib.c:
+ New file.
+ * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-1.c:
+ Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-2-lib.c:
+ Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-2.c:
+ Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-3-lib.c:
+ Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-3.c:
+ Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-4-lib.c:
+ Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-4.c:
+ Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-5-lib.c:
+ Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-5.c:
+ Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-6-lib.c:
+ Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-6.c:
+ Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-7-lib.c:
+ Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-7.c:
+ Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-8-lib.c:
+ Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-8.c:
+ Likewise.
+
2020-04-07 Jakub Jelinek <jakub@redhat.com>
Backported from mainline
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-2-lib.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-2-lib.c
new file mode 100644
index 00000000000..84f41a49dfd
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-2-lib.c
@@ -0,0 +1,3 @@
+/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
+/* { dg-additional-options "-DOPENACC_API" } */
+#include "static-dynamic-lifetimes-2.c"
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-2.c
new file mode 100644
index 00000000000..d3c6f5192d8
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-2.c
@@ -0,0 +1,166 @@
+/* Test nested dynamic/static data mappings. */
+
+/* { 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);
+
+#pragma acc data copy(block1[0:SIZE])
+ {
+#ifdef OPENACC_API
+ acc_copyin (block1, SIZE);
+ acc_copyout (block1, SIZE);
+#else
+#pragma acc enter data copyin(block1[0:SIZE])
+#pragma acc exit data copyout(block1[0:SIZE])
+#endif
+ }
+
+ assert (!acc_is_present (block1, SIZE));
+
+ 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
+
+ assert (!acc_is_present (block1, SIZE));
+
+ free (block1);
+}
+
+void
+f3 (void)
+{
+ char *block1 = (char *) malloc (SIZE);
+
+#pragma acc data copy(block1[0:SIZE])
+ {
+#ifdef OPENACC_API
+ acc_copyin (block1, SIZE);
+ acc_copyin (block1, SIZE);
+ acc_copyout (block1, SIZE);
+ acc_copyout (block1, SIZE);
+#else
+#pragma acc enter data copyin(block1[0:SIZE])
+#pragma acc enter data copyin(block1[0:SIZE])
+#pragma acc exit data copyout(block1[0:SIZE])
+#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);
+
+#pragma acc data copy(block1[0: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_copyin (block1, SIZE);
+ acc_copyout (block1, SIZE);
+#else
+#pragma acc enter data copyin(block1[0:SIZE])
+#pragma acc exit data copyout(block1[0:SIZE])
+#endif
+ }
+
+#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
+f5 (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_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
+ }
+#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);
+}
+
+int
+main (int argc, char *argv[])
+{
+ f1 ();
+ f2 ();
+ f3 ();
+ f4 ();
+ f5 ();
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-3-lib.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-3-lib.c
new file mode 100644
index 00000000000..d9e76c600f0
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-3-lib.c
@@ -0,0 +1,3 @@
+/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
+/* { dg-additional-options "-DOPENACC_API" } */
+#include "static-dynamic-lifetimes-3.c"
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-3.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-3.c
new file mode 100644
index 00000000000..59501864398
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-3.c
@@ -0,0 +1,183 @@
+/* Test nested dynamic/static data mappings (multiple blocks on data
+ regions). */
+
+/* { 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);
+ char *block2 = (char *) malloc (SIZE);
+
+#pragma acc data copy(block1[0:SIZE], block2[0:SIZE])
+ {
+#ifdef OPENACC_API
+ acc_copyin (block1, SIZE);
+ acc_copyout (block1, SIZE);
+#else
+#pragma acc enter data copyin(block1[0:SIZE])
+#pragma acc exit data copyout(block1[0:SIZE])
+#endif
+ }
+
+ assert (!acc_is_present (block1, SIZE));
+ assert (!acc_is_present (block2, SIZE));
+
+ free (block1);
+ free (block2);
+}
+
+void
+f2 (void)
+{
+ 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);
+#else
+#pragma acc exit data copyout(block1[0:SIZE])
+#endif
+
+ assert (!acc_is_present (block1, SIZE));
+ assert (!acc_is_present (block2, SIZE));
+
+ free (block1);
+ free (block2);
+}
+
+void
+f3 (void)
+{
+ char *block1 = (char *) malloc (SIZE);
+ char *block2 = (char *) malloc (SIZE);
+
+#pragma acc data copy(block1[0:SIZE], block2[0:SIZE])
+ {
+#ifdef OPENACC_API
+ acc_copyin (block1, SIZE);
+ acc_copyin (block2, SIZE);
+ acc_copyout (block2, SIZE);
+ acc_copyout (block1, SIZE);
+#else
+#pragma acc enter data copyin(block1[0:SIZE])
+#pragma acc enter data copyin(block2[0:SIZE])
+#pragma acc exit data copyout(block2[0:SIZE])
+#pragma acc exit data copyout(block1[0:SIZE])
+#endif
+ }
+
+ assert (!acc_is_present (block1, SIZE));
+ assert (!acc_is_present (block2, SIZE));
+
+ free (block1);
+ free (block2);
+}
+
+void
+f4 (void)
+{
+ char *block1 = (char *) malloc (SIZE);
+ char *block2 = (char *) malloc (SIZE);
+
+#pragma acc data copy(block1[0:SIZE], block2[0: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_copyin (block2, SIZE);
+ acc_copyout (block2, SIZE);
+#else
+#pragma acc enter data copyin(block2[0:SIZE])
+#pragma acc exit data copyout(block2[0:SIZE])
+#endif
+ }
+#ifdef OPENACC_API
+ acc_copyout (block1, SIZE);
+#else
+#pragma acc exit data copyout(block1[0:SIZE])
+#endif
+ }
+
+ assert (!acc_is_present (block1, SIZE));
+ assert (!acc_is_present (block2, SIZE));
+
+ free (block1);
+ free (block2);
+}
+
+void
+f5 (void)
+{
+ 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_copyin (block2, SIZE);
+#else
+#pragma acc enter data copyin(block2[0:SIZE])
+#endif
+
+#pragma acc data copy(block1[0:SIZE], block2[0:SIZE])
+ {
+ }
+#ifdef OPENACC_API
+ acc_copyout (block2, SIZE);
+#else
+#pragma acc exit data copyout(block2[0:SIZE])
+#endif
+ }
+
+#ifdef OPENACC_API
+ acc_copyout (block1, SIZE);
+#else
+#pragma acc exit data copyout(block1[0:SIZE])
+#endif
+
+ assert (!acc_is_present (block1, SIZE));
+ assert (!acc_is_present (block2, SIZE));
+
+ free (block1);
+ free (block2);
+}
+
+int
+main (int argc, char *argv[])
+{
+ f1 ();
+ f2 ();
+ f3 ();
+ f4 ();
+ f5 ();
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-4-lib.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-4-lib.c
new file mode 100644
index 00000000000..e3c1bfb473d
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-4-lib.c
@@ -0,0 +1,3 @@
+/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
+/* { dg-additional-options "-DOPENACC_API" } */
+#include "static-dynamic-lifetimes-4.c"
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-4.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-4.c
new file mode 100644
index 00000000000..e9a6510ace8
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-4.c
@@ -0,0 +1,64 @@
+/* { 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);
+ char *block3 = (char *) malloc (SIZE);
+
+ /* Doing this twice ensures that we have a non-zero virtual refcount. Make
+ sure that works too. */
+#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], 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. */
+#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 (block1, SIZE);
+ assert (acc_is_present (block1, SIZE));
+ acc_copyout (block1, SIZE);
+ assert (!acc_is_present (block1, SIZE));
+
+ acc_copyout (block2, SIZE);
+ assert (!acc_is_present (block2, 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(block2[0:SIZE])
+ assert (!acc_is_present (block2, SIZE));
+#endif
+
+ free (block1);
+ free (block2);
+ free (block3);
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-5-lib.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-5-lib.c
new file mode 100644
index 00000000000..77703122ad6
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-5-lib.c
@@ -0,0 +1,3 @@
+/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
+/* { dg-additional-options "-DOPENACC_API" } */
+#include "static-dynamic-lifetimes-5.c"
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-5.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-5.c
new file mode 100644
index 00000000000..9807076d3f4
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-5.c
@@ -0,0 +1,56 @@
+/* { 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);
+ char *block3 = (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], 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. */
+#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 (block1, SIZE);
+ assert (!acc_is_present (block1, SIZE));
+
+ acc_copyout (block2, SIZE);
+ assert (!acc_is_present (block2, SIZE));
+#else
+#pragma acc exit data copyout(block1[0:SIZE])
+ assert (!acc_is_present (block1, SIZE));
+
+#pragma acc exit data copyout(block2[0:SIZE])
+ assert (!acc_is_present (block2, SIZE));
+#endif
+
+ free (block1);
+ free (block2);
+ free (block3);
+
+ 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..4a87dd72525
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-6-lib.c
@@ -0,0 +1,3 @@
+/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
+/* { dg-additional-options "-DOPENACC_API" } */
+#include "static-dynamic-lifetimes-6.c"
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..3e5c4d7ea11
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-6.c
@@ -0,0 +1,42 @@
+/* { 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);
+#else
+#pragma acc exit data copyout(block1[0:SIZE], block2[0:SIZE])
+#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..8ccbb126933
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-7-lib.c
@@ -0,0 +1,3 @@
+/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
+/* { dg-additional-options "-DOPENACC_API" } */
+#include "static-dynamic-lifetimes-7.c"
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..2735d6fa0eb
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-7.c
@@ -0,0 +1,42 @@
+/* { 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);
+#else
+#pragma acc exit data copyout(block1[0:SIZE])
+#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..f3104cbd035
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-8-lib.c
@@ -0,0 +1,3 @@
+/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
+/* { dg-additional-options "-DOPENACC_API" } */
+#include "static-dynamic-lifetimes-8.c"
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..919ee02b725
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-8.c
@@ -0,0 +1,47 @@
+/* { 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);
+#else
+#pragma acc exit data copyout(block1[0:SIZE])
+#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.17.1
^ permalink raw reply [flat|nested] 8+ messages in thread
* Re: Test cases for mixed structured/dynamic data lifetimes with OpenACC [PR92843]
2020-04-10 14:44 ` Test cases for mixed structured/dynamic data lifetimes with OpenACC [PR92843] (was: [PATCH 0/3] Mixed static/dynamic data lifetimes with OpenACC (PR92843)) Thomas Schwinge
@ 2020-04-13 7:23 ` Thomas Schwinge
0 siblings, 0 replies; 8+ messages in thread
From: Thomas Schwinge @ 2020-04-13 7:23 UTC (permalink / raw)
To: gcc-patches; +Cc: Julian Brown, Jakub Jelinek
[-- Attachment #1: Type: text/plain, Size: 1156 bytes --]
Hi!
On 2020-04-10T16:44:42+0200, I wrote:
> On 2020-01-17T12:18:18-0800, Julian Brown <julian@codesourcery.com> wrote:
>> This patch series provides fixes for some cases of mixing static and
>
> (It's "structured", not "static".) ;-)
Of course, that should be reflected...
> libgomp/
> PR libgomp/92843
> * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-1-lib.c:
> New file.
> [...]
... in the test cases, too.
See attached for what I pushed to master branch in commit
af4c92573dc462a17a6c345756889d28054ed591 "Rename
'libgomp.oacc-c-c++-common/static-dynamic-lifetimes-*' to
'libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-*' [PR92843]",
and to releases/gcc-9 branch in commit
a99a8431e670a6c0ac861d738249ff4d94d6552e "Rename
'libgomp.oacc-c-c++-common/static-dynamic-lifetimes-*' to
'libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-*' [PR92843]".
Grüße
Thomas
-----------------
Mentor Graphics (Deutschland) GmbH, Arnulfstraße 201, 80634 München / Germany
Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Alexander Walter
[-- Warning: decoded text below may be mangled, UTF-8 assumed --]
[-- Attachment #2: 0001-Rename-libgomp.oacc-c-c-common-static-dynamic-lifeti.patch --]
[-- Type: text/x-diff, Size: 23070 bytes --]
From af4c92573dc462a17a6c345756889d28054ed591 Mon Sep 17 00:00:00 2001
From: Thomas Schwinge <thomas@codesourcery.com>
Date: Mon, 13 Apr 2020 08:56:03 +0200
Subject: [PATCH] Rename 'libgomp.oacc-c-c++-common/static-dynamic-lifetimes-*'
to 'libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-*' [PR92843]
Fix-up for commit be9862dd96945772ae0692bc95b37ec6dbcabda0 "Test cases for
mixed structured/dynamic data lifetimes with OpenACC [PR92843]": it's
"structured", not "static" data lifetimes/reference counters.
libgomp/
PR libgomp/92843
* testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-1-lib.c:
Rename to...
* testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-1-lib.c:
... this.
* testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-1.c:
Rename to...
* testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-1.c:
... this.
* testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-2-lib.c:
Rename to...
* testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-2-lib.c:
... this.
* testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-2.c:
Rename to...
* testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-2.c:
... this.
* testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-3-lib.c:
Rename to...
* testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-3-lib.c:
... this.
* testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-3.c:
Rename to...
* testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-3.c:
... this.
* testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-4-lib.c:
Rename to...
* testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-4-lib.c:
... this.
* testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-4.c:
Rename to...
* testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-4.c:
... this.
* testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-5-lib.c:
Rename to...
* testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-5-lib.c:
... this.
* testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-5.c:
Rename to...
* testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-5.c:
... this.
* testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-6-lib.c:
Rename to...
* testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-6-lib.c:
... this.
* testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-6.c:
Rename to...
* testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-6.c:
... this.
* testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-7-lib.c:
Rename to...
* testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-7-lib.c:
... this.
* testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-7.c:
Rename to...
* testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-7.c:
... this.
* testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-8-lib.c:
Rename to...
* testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-8-lib.c:
... this.
* testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-8.c::
Rename to...
* testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-8.c:
... this.
---
libgomp/ChangeLog | 68 +++++++++++++++++++
.../static-dynamic-lifetimes-5-lib.c | 3 -
.../static-dynamic-lifetimes-6-lib.c | 3 -
.../static-dynamic-lifetimes-7-lib.c | 3 -
.../static-dynamic-lifetimes-8-lib.c | 3 -
...c => structured-dynamic-lifetimes-1-lib.c} | 2 +-
...s-1.c => structured-dynamic-lifetimes-1.c} | 5 +-
...c => structured-dynamic-lifetimes-2-lib.c} | 2 +-
...s-2.c => structured-dynamic-lifetimes-2.c} | 2 +-
...c => structured-dynamic-lifetimes-3-lib.c} | 2 +-
...s-3.c => structured-dynamic-lifetimes-3.c} | 2 +-
...c => structured-dynamic-lifetimes-4-lib.c} | 2 +-
...s-4.c => structured-dynamic-lifetimes-4.c} | 0
.../structured-dynamic-lifetimes-5-lib.c | 3 +
...s-5.c => structured-dynamic-lifetimes-5.c} | 0
.../structured-dynamic-lifetimes-6-lib.c | 3 +
...s-6.c => structured-dynamic-lifetimes-6.c} | 3 +-
.../structured-dynamic-lifetimes-7-lib.c | 3 +
...s-7.c => structured-dynamic-lifetimes-7.c} | 6 +-
.../structured-dynamic-lifetimes-8-lib.c | 3 +
...s-8.c => structured-dynamic-lifetimes-8.c} | 0
21 files changed, 95 insertions(+), 23 deletions(-)
delete mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-5-lib.c
delete mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-6-lib.c
delete mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-7-lib.c
delete mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-8-lib.c
rename libgomp/testsuite/libgomp.oacc-c-c++-common/{static-dynamic-lifetimes-1-lib.c => structured-dynamic-lifetimes-1-lib.c} (70%)
rename libgomp/testsuite/libgomp.oacc-c-c++-common/{static-dynamic-lifetimes-1.c => structured-dynamic-lifetimes-1.c} (95%)
rename libgomp/testsuite/libgomp.oacc-c-c++-common/{static-dynamic-lifetimes-3-lib.c => structured-dynamic-lifetimes-2-lib.c} (70%)
rename libgomp/testsuite/libgomp.oacc-c-c++-common/{static-dynamic-lifetimes-2.c => structured-dynamic-lifetimes-2.c} (98%)
rename libgomp/testsuite/libgomp.oacc-c-c++-common/{static-dynamic-lifetimes-2-lib.c => structured-dynamic-lifetimes-3-lib.c} (70%)
rename libgomp/testsuite/libgomp.oacc-c-c++-common/{static-dynamic-lifetimes-3.c => structured-dynamic-lifetimes-3.c} (97%)
rename libgomp/testsuite/libgomp.oacc-c-c++-common/{static-dynamic-lifetimes-4-lib.c => structured-dynamic-lifetimes-4-lib.c} (70%)
rename libgomp/testsuite/libgomp.oacc-c-c++-common/{static-dynamic-lifetimes-4.c => structured-dynamic-lifetimes-4.c} (100%)
create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-5-lib.c
rename libgomp/testsuite/libgomp.oacc-c-c++-common/{static-dynamic-lifetimes-5.c => structured-dynamic-lifetimes-5.c} (100%)
create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-6-lib.c
rename libgomp/testsuite/libgomp.oacc-c-c++-common/{static-dynamic-lifetimes-6.c => structured-dynamic-lifetimes-6.c} (90%)
create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-7-lib.c
rename libgomp/testsuite/libgomp.oacc-c-c++-common/{static-dynamic-lifetimes-7.c => structured-dynamic-lifetimes-7.c} (82%)
create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-8-lib.c
rename libgomp/testsuite/libgomp.oacc-c-c++-common/{static-dynamic-lifetimes-8.c => structured-dynamic-lifetimes-8.c} (100%)
diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog
index b0f19845ad2..8f351f61f33 100644
--- a/libgomp/ChangeLog
+++ b/libgomp/ChangeLog
@@ -1,3 +1,71 @@
+2020-04-13 Thomas Schwinge <thomas@codesourcery.com>
+
+ PR libgomp/92843
+ * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-1-lib.c:
+ Rename to...
+ * testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-1-lib.c:
+ ... this.
+ * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-1.c:
+ Rename to...
+ * testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-1.c:
+ ... this.
+ * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-2-lib.c:
+ Rename to...
+ * testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-2-lib.c:
+ ... this.
+ * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-2.c:
+ Rename to...
+ * testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-2.c:
+ ... this.
+ * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-3-lib.c:
+ Rename to...
+ * testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-3-lib.c:
+ ... this.
+ * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-3.c:
+ Rename to...
+ * testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-3.c:
+ ... this.
+ * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-4-lib.c:
+ Rename to...
+ * testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-4-lib.c:
+ ... this.
+ * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-4.c:
+ Rename to...
+ * testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-4.c:
+ ... this.
+ * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-5-lib.c:
+ Rename to...
+ * testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-5-lib.c:
+ ... this.
+ * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-5.c:
+ Rename to...
+ * testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-5.c:
+ ... this.
+ * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-6-lib.c:
+ Rename to...
+ * testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-6-lib.c:
+ ... this.
+ * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-6.c:
+ Rename to...
+ * testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-6.c:
+ ... this.
+ * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-7-lib.c:
+ Rename to...
+ * testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-7-lib.c:
+ ... this.
+ * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-7.c:
+ Rename to...
+ * testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-7.c:
+ ... this.
+ * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-8-lib.c:
+ Rename to...
+ * testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-8-lib.c:
+ ... this.
+ * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-8.c::
+ Rename to...
+ * testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-8.c:
+ ... this.
+
2020-04-10 Julian Brown <julian@codesourcery.com>
Thomas Schwinge <thomas@codesourcery.com>
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-5-lib.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-5-lib.c
deleted file mode 100644
index 77703122ad6..00000000000
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-5-lib.c
+++ /dev/null
@@ -1,3 +0,0 @@
-/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
-/* { dg-additional-options "-DOPENACC_API" } */
-#include "static-dynamic-lifetimes-5.c"
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
deleted file mode 100644
index 4a87dd72525..00000000000
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-6-lib.c
+++ /dev/null
@@ -1,3 +0,0 @@
-/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
-/* { dg-additional-options "-DOPENACC_API" } */
-#include "static-dynamic-lifetimes-6.c"
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
deleted file mode 100644
index 8ccbb126933..00000000000
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-7-lib.c
+++ /dev/null
@@ -1,3 +0,0 @@
-/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
-/* { dg-additional-options "-DOPENACC_API" } */
-#include "static-dynamic-lifetimes-7.c"
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
deleted file mode 100644
index f3104cbd035..00000000000
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-8-lib.c
+++ /dev/null
@@ -1,3 +0,0 @@
-/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
-/* { dg-additional-options "-DOPENACC_API" } */
-#include "static-dynamic-lifetimes-8.c"
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-1-lib.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-1-lib.c
similarity index 70%
rename from libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-1-lib.c
rename to libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-1-lib.c
index 23c20d4fab7..8fa87777f1a 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-1-lib.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-1-lib.c
@@ -1,3 +1,3 @@
/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
/* { dg-additional-options "-DOPENACC_API" } */
-#include "static-dynamic-lifetimes-1.c"
+#include "structured-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/structured-dynamic-lifetimes-1.c
similarity index 95%
rename from libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-1.c
rename to libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-1.c
index a743660f53e..0d6b4159ad0 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-1.c
@@ -1,4 +1,4 @@
-/* Test transitioning of data lifetimes between static and dynamic. */
+/* Test transitioning of data lifetimes between structured and dynamic. */
/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
@@ -69,7 +69,8 @@ f2 (void)
#else
#pragma acc exit data copyout(block1[0:SIZE])
#endif
- /* This should stay present until the end of the static data lifetime. */
+ /* This should stay present until the end of the structured data
+ lifetime. */
assert (acc_is_present (block1, SIZE));
}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-3-lib.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-2-lib.c
similarity index 70%
rename from libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-3-lib.c
rename to libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-2-lib.c
index d9e76c600f0..365df8d7f7d 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-3-lib.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-2-lib.c
@@ -1,3 +1,3 @@
/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
/* { dg-additional-options "-DOPENACC_API" } */
-#include "static-dynamic-lifetimes-3.c"
+#include "structured-dynamic-lifetimes-2.c"
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-2.c
similarity index 98%
rename from libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-2.c
rename to libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-2.c
index d3c6f5192d8..726942c76fd 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-2.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-2.c
@@ -1,4 +1,4 @@
-/* Test nested dynamic/static data mappings. */
+/* Test nested dynamic/structured data mappings. */
/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-2-lib.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-3-lib.c
similarity index 70%
rename from libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-2-lib.c
rename to libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-3-lib.c
index 84f41a49dfd..469b35b76f1 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-2-lib.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-3-lib.c
@@ -1,3 +1,3 @@
/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
/* { dg-additional-options "-DOPENACC_API" } */
-#include "static-dynamic-lifetimes-2.c"
+#include "structured-dynamic-lifetimes-3.c"
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-3.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-3.c
similarity index 97%
rename from libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-3.c
rename to libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-3.c
index 59501864398..c13f3c56584 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-3.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-3.c
@@ -1,4 +1,4 @@
-/* Test nested dynamic/static data mappings (multiple blocks on data
+/* Test nested dynamic/structured data mappings (multiple blocks on data
regions). */
/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-4-lib.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-4-lib.c
similarity index 70%
rename from libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-4-lib.c
rename to libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-4-lib.c
index e3c1bfb473d..8e88b97b2ce 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-4-lib.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-4-lib.c
@@ -1,3 +1,3 @@
/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
/* { dg-additional-options "-DOPENACC_API" } */
-#include "static-dynamic-lifetimes-4.c"
+#include "structured-dynamic-lifetimes-4.c"
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-4.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-4.c
similarity index 100%
rename from libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-4.c
rename to libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-4.c
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-5-lib.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-5-lib.c
new file mode 100644
index 00000000000..59ef5626bd0
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-5-lib.c
@@ -0,0 +1,3 @@
+/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
+/* { dg-additional-options "-DOPENACC_API" } */
+#include "structured-dynamic-lifetimes-5.c"
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-5.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-5.c
similarity index 100%
rename from libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-5.c
rename to libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-5.c
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-6-lib.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-6-lib.c
new file mode 100644
index 00000000000..0401f73966f
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-6-lib.c
@@ -0,0 +1,3 @@
+/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
+/* { dg-additional-options "-DOPENACC_API" } */
+#include "structured-dynamic-lifetimes-6.c"
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-6.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-6.c
similarity index 90%
rename from libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-6.c
rename to libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-6.c
index 3e5c4d7ea11..9250b4af3ed 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-6.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-6.c
@@ -27,7 +27,8 @@ main (int argc, char *argv[])
#else
#pragma acc exit data copyout(block1[0:SIZE], block2[0:SIZE])
#endif
- /* These should stay present until the end of the static data lifetime. */
+ /* These should stay present until the end of the structured data
+ lifetime. */
assert (acc_is_present (block1, SIZE));
assert (acc_is_present (block2, SIZE));
}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-7-lib.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-7-lib.c
new file mode 100644
index 00000000000..07caefbd082
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-7-lib.c
@@ -0,0 +1,3 @@
+/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
+/* { dg-additional-options "-DOPENACC_API" } */
+#include "structured-dynamic-lifetimes-7.c"
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-7.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-7.c
similarity index 82%
rename from libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-7.c
rename to libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-7.c
index 2735d6fa0eb..52e8d4c9959 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-7.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-7.c
@@ -21,13 +21,15 @@ main (int argc, char *argv[])
#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. */
+ enclosing structured data region here, because that region maps block2
+ also. */
#ifdef OPENACC_API
acc_copyout (block1, SIZE);
#else
#pragma acc exit data copyout(block1[0:SIZE])
#endif
- /* These should stay present until the end of the static data lifetime. */
+ /* These should stay present until the end of the structured data
+ lifetime. */
assert (acc_is_present (block1, SIZE));
assert (acc_is_present (block2, SIZE));
}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-8-lib.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-8-lib.c
new file mode 100644
index 00000000000..1c2479ad96f
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-8-lib.c
@@ -0,0 +1,3 @@
+/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
+/* { dg-additional-options "-DOPENACC_API" } */
+#include "structured-dynamic-lifetimes-8.c"
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-8.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-8.c
similarity index 100%
rename from libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-8.c
rename to libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-8.c
--
2.17.1
[-- Warning: decoded text below may be mangled, UTF-8 assumed --]
[-- Attachment #3: 0001-Rename-libgomp.oacc-c-c-common-static-dynamic-lif.g9.patch --]
[-- Type: text/x-diff, Size: 23140 bytes --]
From a99a8431e670a6c0ac861d738249ff4d94d6552e Mon Sep 17 00:00:00 2001
From: Thomas Schwinge <thomas@codesourcery.com>
Date: Mon, 13 Apr 2020 08:56:03 +0200
Subject: [PATCH] Rename 'libgomp.oacc-c-c++-common/static-dynamic-lifetimes-*'
to 'libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-*' [PR92843]
Fix-up for commit be9862dd96945772ae0692bc95b37ec6dbcabda0 "Test cases for
mixed structured/dynamic data lifetimes with OpenACC [PR92843]": it's
"structured", not "static" data lifetimes/reference counters.
libgomp/
PR libgomp/92843
* testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-1-lib.c:
Rename to...
* testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-1-lib.c:
... this.
* testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-1.c:
Rename to...
* testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-1.c:
... this.
* testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-2-lib.c:
Rename to...
* testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-2-lib.c:
... this.
* testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-2.c:
Rename to...
* testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-2.c:
... this.
* testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-3-lib.c:
Rename to...
* testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-3-lib.c:
... this.
* testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-3.c:
Rename to...
* testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-3.c:
... this.
* testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-4-lib.c:
Rename to...
* testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-4-lib.c:
... this.
* testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-4.c:
Rename to...
* testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-4.c:
... this.
* testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-5-lib.c:
Rename to...
* testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-5-lib.c:
... this.
* testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-5.c:
Rename to...
* testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-5.c:
... this.
* testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-6-lib.c:
Rename to...
* testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-6-lib.c:
... this.
* testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-6.c:
Rename to...
* testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-6.c:
... this.
* testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-7-lib.c:
Rename to...
* testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-7-lib.c:
... this.
* testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-7.c:
Rename to...
* testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-7.c:
... this.
* testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-8-lib.c:
Rename to...
* testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-8-lib.c:
... this.
* testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-8.c::
Rename to...
* testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-8.c:
... this.
(cherry picked from commit af4c92573dc462a17a6c345756889d28054ed591)
---
libgomp/ChangeLog | 68 +++++++++++++++++++
.../static-dynamic-lifetimes-5-lib.c | 3 -
.../static-dynamic-lifetimes-6-lib.c | 3 -
.../static-dynamic-lifetimes-7-lib.c | 3 -
.../static-dynamic-lifetimes-8-lib.c | 3 -
...c => structured-dynamic-lifetimes-1-lib.c} | 2 +-
...s-1.c => structured-dynamic-lifetimes-1.c} | 5 +-
...c => structured-dynamic-lifetimes-2-lib.c} | 2 +-
...s-2.c => structured-dynamic-lifetimes-2.c} | 2 +-
...c => structured-dynamic-lifetimes-3-lib.c} | 2 +-
...s-3.c => structured-dynamic-lifetimes-3.c} | 2 +-
...c => structured-dynamic-lifetimes-4-lib.c} | 2 +-
...s-4.c => structured-dynamic-lifetimes-4.c} | 0
.../structured-dynamic-lifetimes-5-lib.c | 3 +
...s-5.c => structured-dynamic-lifetimes-5.c} | 0
.../structured-dynamic-lifetimes-6-lib.c | 3 +
...s-6.c => structured-dynamic-lifetimes-6.c} | 3 +-
.../structured-dynamic-lifetimes-7-lib.c | 3 +
...s-7.c => structured-dynamic-lifetimes-7.c} | 6 +-
.../structured-dynamic-lifetimes-8-lib.c | 3 +
...s-8.c => structured-dynamic-lifetimes-8.c} | 0
21 files changed, 95 insertions(+), 23 deletions(-)
delete mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-5-lib.c
delete mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-6-lib.c
delete mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-7-lib.c
delete mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-8-lib.c
rename libgomp/testsuite/libgomp.oacc-c-c++-common/{static-dynamic-lifetimes-1-lib.c => structured-dynamic-lifetimes-1-lib.c} (70%)
rename libgomp/testsuite/libgomp.oacc-c-c++-common/{static-dynamic-lifetimes-1.c => structured-dynamic-lifetimes-1.c} (95%)
rename libgomp/testsuite/libgomp.oacc-c-c++-common/{static-dynamic-lifetimes-3-lib.c => structured-dynamic-lifetimes-2-lib.c} (70%)
rename libgomp/testsuite/libgomp.oacc-c-c++-common/{static-dynamic-lifetimes-2.c => structured-dynamic-lifetimes-2.c} (98%)
rename libgomp/testsuite/libgomp.oacc-c-c++-common/{static-dynamic-lifetimes-2-lib.c => structured-dynamic-lifetimes-3-lib.c} (70%)
rename libgomp/testsuite/libgomp.oacc-c-c++-common/{static-dynamic-lifetimes-3.c => structured-dynamic-lifetimes-3.c} (97%)
rename libgomp/testsuite/libgomp.oacc-c-c++-common/{static-dynamic-lifetimes-4-lib.c => structured-dynamic-lifetimes-4-lib.c} (70%)
rename libgomp/testsuite/libgomp.oacc-c-c++-common/{static-dynamic-lifetimes-4.c => structured-dynamic-lifetimes-4.c} (100%)
create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-5-lib.c
rename libgomp/testsuite/libgomp.oacc-c-c++-common/{static-dynamic-lifetimes-5.c => structured-dynamic-lifetimes-5.c} (100%)
create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-6-lib.c
rename libgomp/testsuite/libgomp.oacc-c-c++-common/{static-dynamic-lifetimes-6.c => structured-dynamic-lifetimes-6.c} (90%)
create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-7-lib.c
rename libgomp/testsuite/libgomp.oacc-c-c++-common/{static-dynamic-lifetimes-7.c => structured-dynamic-lifetimes-7.c} (82%)
create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-8-lib.c
rename libgomp/testsuite/libgomp.oacc-c-c++-common/{static-dynamic-lifetimes-8.c => structured-dynamic-lifetimes-8.c} (100%)
diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog
index 39308ecbc03..2926096bb7c 100644
--- a/libgomp/ChangeLog
+++ b/libgomp/ChangeLog
@@ -1,3 +1,71 @@
+2020-04-13 Thomas Schwinge <thomas@codesourcery.com>
+
+ PR libgomp/92843
+ * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-1-lib.c:
+ Rename to...
+ * testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-1-lib.c:
+ ... this.
+ * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-1.c:
+ Rename to...
+ * testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-1.c:
+ ... this.
+ * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-2-lib.c:
+ Rename to...
+ * testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-2-lib.c:
+ ... this.
+ * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-2.c:
+ Rename to...
+ * testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-2.c:
+ ... this.
+ * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-3-lib.c:
+ Rename to...
+ * testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-3-lib.c:
+ ... this.
+ * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-3.c:
+ Rename to...
+ * testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-3.c:
+ ... this.
+ * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-4-lib.c:
+ Rename to...
+ * testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-4-lib.c:
+ ... this.
+ * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-4.c:
+ Rename to...
+ * testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-4.c:
+ ... this.
+ * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-5-lib.c:
+ Rename to...
+ * testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-5-lib.c:
+ ... this.
+ * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-5.c:
+ Rename to...
+ * testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-5.c:
+ ... this.
+ * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-6-lib.c:
+ Rename to...
+ * testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-6-lib.c:
+ ... this.
+ * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-6.c:
+ Rename to...
+ * testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-6.c:
+ ... this.
+ * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-7-lib.c:
+ Rename to...
+ * testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-7-lib.c:
+ ... this.
+ * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-7.c:
+ Rename to...
+ * testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-7.c:
+ ... this.
+ * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-8-lib.c:
+ Rename to...
+ * testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-8-lib.c:
+ ... this.
+ * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-8.c::
+ Rename to...
+ * testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-8.c:
+ ... this.
+
2020-04-10 Julian Brown <julian@codesourcery.com>
Thomas Schwinge <thomas@codesourcery.com>
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-5-lib.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-5-lib.c
deleted file mode 100644
index 77703122ad6..00000000000
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-5-lib.c
+++ /dev/null
@@ -1,3 +0,0 @@
-/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
-/* { dg-additional-options "-DOPENACC_API" } */
-#include "static-dynamic-lifetimes-5.c"
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
deleted file mode 100644
index 4a87dd72525..00000000000
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-6-lib.c
+++ /dev/null
@@ -1,3 +0,0 @@
-/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
-/* { dg-additional-options "-DOPENACC_API" } */
-#include "static-dynamic-lifetimes-6.c"
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
deleted file mode 100644
index 8ccbb126933..00000000000
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-7-lib.c
+++ /dev/null
@@ -1,3 +0,0 @@
-/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
-/* { dg-additional-options "-DOPENACC_API" } */
-#include "static-dynamic-lifetimes-7.c"
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
deleted file mode 100644
index f3104cbd035..00000000000
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-8-lib.c
+++ /dev/null
@@ -1,3 +0,0 @@
-/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
-/* { dg-additional-options "-DOPENACC_API" } */
-#include "static-dynamic-lifetimes-8.c"
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-1-lib.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-1-lib.c
similarity index 70%
rename from libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-1-lib.c
rename to libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-1-lib.c
index 23c20d4fab7..8fa87777f1a 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-1-lib.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-1-lib.c
@@ -1,3 +1,3 @@
/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
/* { dg-additional-options "-DOPENACC_API" } */
-#include "static-dynamic-lifetimes-1.c"
+#include "structured-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/structured-dynamic-lifetimes-1.c
similarity index 95%
rename from libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-1.c
rename to libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-1.c
index a743660f53e..0d6b4159ad0 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-1.c
@@ -1,4 +1,4 @@
-/* Test transitioning of data lifetimes between static and dynamic. */
+/* Test transitioning of data lifetimes between structured and dynamic. */
/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
@@ -69,7 +69,8 @@ f2 (void)
#else
#pragma acc exit data copyout(block1[0:SIZE])
#endif
- /* This should stay present until the end of the static data lifetime. */
+ /* This should stay present until the end of the structured data
+ lifetime. */
assert (acc_is_present (block1, SIZE));
}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-3-lib.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-2-lib.c
similarity index 70%
rename from libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-3-lib.c
rename to libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-2-lib.c
index d9e76c600f0..365df8d7f7d 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-3-lib.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-2-lib.c
@@ -1,3 +1,3 @@
/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
/* { dg-additional-options "-DOPENACC_API" } */
-#include "static-dynamic-lifetimes-3.c"
+#include "structured-dynamic-lifetimes-2.c"
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-2.c
similarity index 98%
rename from libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-2.c
rename to libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-2.c
index d3c6f5192d8..726942c76fd 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-2.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-2.c
@@ -1,4 +1,4 @@
-/* Test nested dynamic/static data mappings. */
+/* Test nested dynamic/structured data mappings. */
/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-2-lib.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-3-lib.c
similarity index 70%
rename from libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-2-lib.c
rename to libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-3-lib.c
index 84f41a49dfd..469b35b76f1 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-2-lib.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-3-lib.c
@@ -1,3 +1,3 @@
/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
/* { dg-additional-options "-DOPENACC_API" } */
-#include "static-dynamic-lifetimes-2.c"
+#include "structured-dynamic-lifetimes-3.c"
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-3.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-3.c
similarity index 97%
rename from libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-3.c
rename to libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-3.c
index 59501864398..c13f3c56584 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-3.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-3.c
@@ -1,4 +1,4 @@
-/* Test nested dynamic/static data mappings (multiple blocks on data
+/* Test nested dynamic/structured data mappings (multiple blocks on data
regions). */
/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-4-lib.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-4-lib.c
similarity index 70%
rename from libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-4-lib.c
rename to libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-4-lib.c
index e3c1bfb473d..8e88b97b2ce 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-4-lib.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-4-lib.c
@@ -1,3 +1,3 @@
/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
/* { dg-additional-options "-DOPENACC_API" } */
-#include "static-dynamic-lifetimes-4.c"
+#include "structured-dynamic-lifetimes-4.c"
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-4.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-4.c
similarity index 100%
rename from libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-4.c
rename to libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-4.c
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-5-lib.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-5-lib.c
new file mode 100644
index 00000000000..59ef5626bd0
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-5-lib.c
@@ -0,0 +1,3 @@
+/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
+/* { dg-additional-options "-DOPENACC_API" } */
+#include "structured-dynamic-lifetimes-5.c"
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-5.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-5.c
similarity index 100%
rename from libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-5.c
rename to libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-5.c
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-6-lib.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-6-lib.c
new file mode 100644
index 00000000000..0401f73966f
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-6-lib.c
@@ -0,0 +1,3 @@
+/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
+/* { dg-additional-options "-DOPENACC_API" } */
+#include "structured-dynamic-lifetimes-6.c"
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-6.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-6.c
similarity index 90%
rename from libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-6.c
rename to libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-6.c
index 3e5c4d7ea11..9250b4af3ed 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-6.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-6.c
@@ -27,7 +27,8 @@ main (int argc, char *argv[])
#else
#pragma acc exit data copyout(block1[0:SIZE], block2[0:SIZE])
#endif
- /* These should stay present until the end of the static data lifetime. */
+ /* These should stay present until the end of the structured data
+ lifetime. */
assert (acc_is_present (block1, SIZE));
assert (acc_is_present (block2, SIZE));
}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-7-lib.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-7-lib.c
new file mode 100644
index 00000000000..07caefbd082
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-7-lib.c
@@ -0,0 +1,3 @@
+/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
+/* { dg-additional-options "-DOPENACC_API" } */
+#include "structured-dynamic-lifetimes-7.c"
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-7.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-7.c
similarity index 82%
rename from libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-7.c
rename to libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-7.c
index 2735d6fa0eb..52e8d4c9959 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-7.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-7.c
@@ -21,13 +21,15 @@ main (int argc, char *argv[])
#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. */
+ enclosing structured data region here, because that region maps block2
+ also. */
#ifdef OPENACC_API
acc_copyout (block1, SIZE);
#else
#pragma acc exit data copyout(block1[0:SIZE])
#endif
- /* These should stay present until the end of the static data lifetime. */
+ /* These should stay present until the end of the structured data
+ lifetime. */
assert (acc_is_present (block1, SIZE));
assert (acc_is_present (block2, SIZE));
}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-8-lib.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-8-lib.c
new file mode 100644
index 00000000000..1c2479ad96f
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-8-lib.c
@@ -0,0 +1,3 @@
+/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
+/* { dg-additional-options "-DOPENACC_API" } */
+#include "structured-dynamic-lifetimes-8.c"
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-8.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-8.c
similarity index 100%
rename from libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-8.c
rename to libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-8.c
--
2.17.1
^ permalink raw reply [flat|nested] 8+ messages in thread
* Re: [PATCH 3/3] OpenACC dynamic data lifetimes ending within structured blocks
2020-01-17 21:31 ` [PATCH 3/3] OpenACC dynamic data lifetimes ending within structured blocks Julian Brown
@ 2020-05-11 15:05 ` Thomas Schwinge
0 siblings, 0 replies; 8+ messages in thread
From: Thomas Schwinge @ 2020-05-11 15:05 UTC (permalink / raw)
To: Julian Brown; +Cc: Jakub Jelinek, gcc-patches
Hi Julian!
On 2020-01-17T13:18:21-0800, Julian Brown <julian@codesourcery.com> wrote:
> 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.
As discussed before, all these test cases were already PASSing before any
of this thread's suggested patches (also for GCC 9), so "from a user's
point of view", all we get here are testsuite regressions:
- 'libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-6-lib.c'
- 'libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-6.c'
- 'libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-7-lib.c'
- 'libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-7.c'
- 'libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-8-lib.c'
- 'libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-8.c'
(Adjusted for the version of the test cases already committed; but
already XFAILed in your original patch submission, see below.)
And: the code changes proposed here are breaking compatibility with GCC
9, such that OpenACC/Fortran code compiled with GCC 9, but running with
recent runtime libraries (common case for users, distributions) would
then terminate with: 'libgomp: cannot handle 'exit data' within data
region'. For example, half of all 'libgomp.oacc-fortran' test cases
using OpenACC 'exit data':
- 'libgomp.oacc-fortran/data-2.f90'
- 'libgomp.oacc-fortran/data-3.f90'
- 'libgomp.oacc-fortran/data-4-2.f90'
- 'libgomp.oacc-fortran/data-4.f90'
- 'libgomp.oacc-fortran/if-1.f90'
Even though that "code generation problem" doesn't exist with GCC 10 and
newer, we still have to maintain ABI compatibility with existing binaries
compiled compiled with GCC 9. (Or, as a less preferred solution, arrange
so that they use host-fallback execution insted of offloading.)
Grüße
Thomas
> 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;
> +}
-----------------
Mentor Graphics (Deutschland) GmbH, Arnulfstraße 201, 80634 München / Germany
Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Alexander Walter
^ permalink raw reply [flat|nested] 8+ messages in thread
* Re: [PATCH 2/3] Don't copy back vars mapped with acc_map_data
2020-01-17 21:20 ` [PATCH 2/3] Don't copy back vars mapped with acc_map_data Julian Brown
@ 2020-07-03 15:53 ` Thomas Schwinge
0 siblings, 0 replies; 8+ messages in thread
From: Thomas Schwinge @ 2020-07-03 15:53 UTC (permalink / raw)
To: Julian Brown; +Cc: gcc-patches, Jakub Jelinek, catherine_moore
Hi!
On 2020-01-17T13:18:20-0800, Julian Brown <julian@codesourcery.com> wrote:
> This patch prevents "exit data" directives from copying back data that
> was mapped with an acc_map_data API call. This matches the behaviour
> expected by the pr92843-1.c test
> --- a/libgomp/oacc-mem.c
> +++ b/libgomp/oacc-mem.c
> @@ -1235,6 +1235,7 @@ goacc_exit_data_internal (struct gomp_device_descr *acc_dev, size_t mapnum,
> n->refcount--;
>
> if (copyfrom
> + && n->refcount != REFCOUNT_INFINITY
> && (kind != GOMP_MAP_FROM || n->refcount == 0))
> gomp_copy_dev2host (acc_dev, aq, (void *) cur_node.host_start,
> (void *) (n->tgt->tgt_start + n->tgt_offset
Such a change -- re-posted as
<http://mid.mail-archive.com/7b2f54faa0a25a7a445ad86bf4726202a1190a4f.1590182783.git.julian@codesourcery.com>,
and
<http://mid.mail-archive.com/0585b4fda1ba5c76dce2ac5053a55e2ceef06041.1592343756.git.julian@codesourcery.com>
-- does fix the 'libgomp.oacc-c-c++-common/pr92843-1.c' test case, but I
don't agree that the change is correct. Instead, I have pushed
"[OpenACC] Revert always-copyfrom behavior for 'GOMP_MAP_FORCE_FROM' in
'libgomp/oacc-mem.c:goacc_exit_data_internal'",
<http://mid.mail-archive.com/87wo3ky5vn.fsf@euler.schwinge.homeip.net>.
(This moves us past the point where it used to 'abort' before, but
doesn't resolve the XFAIL, yet, which needs changes from "[OpenACC]
Adjust dynamic reference count semantics",
<http://mid.mail-archive.com/5e9472b80dc475214a4a082ef54ee919d7f9dcff.1592343756.git.julian@codesourcery.com>.)
Grüße
Thomas
-----------------
Mentor Graphics (Deutschland) GmbH, Arnulfstraße 201, 80634 München / Germany
Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Alexander Walter
^ permalink raw reply [flat|nested] 8+ messages in thread
end of thread, other threads:[~2020-07-03 15:54 UTC | newest]
Thread overview: 8+ messages (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2020-01-17 21:18 [PATCH 0/3] Mixed static/dynamic data lifetimes with OpenACC (PR92843) Julian Brown
2020-01-17 21:20 ` [PATCH 2/3] Don't copy back vars mapped with acc_map_data Julian Brown
2020-07-03 15:53 ` Thomas Schwinge
2020-01-17 21:31 ` [PATCH 3/3] OpenACC dynamic data lifetimes ending within structured blocks Julian Brown
2020-05-11 15:05 ` Thomas Schwinge
2020-01-17 21:31 ` [PATCH 1/3] Introduce dynamic data mapping sentinel for OpenACC Julian Brown
2020-04-10 14:44 ` Test cases for mixed structured/dynamic data lifetimes with OpenACC [PR92843] (was: [PATCH 0/3] Mixed static/dynamic data lifetimes with OpenACC (PR92843)) Thomas Schwinge
2020-04-13 7:23 ` Test cases for mixed structured/dynamic data lifetimes with OpenACC [PR92843] Thomas Schwinge
This is a public inbox, see mirroring instructions
for how to clone and mirror all data and code used for this inbox;
as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).