Test cases for mixed structured/dynamic data lifetimes with OpenACC [PR92843] (was: [PATCH 0/3] Mixed static/dynamic data lifetimes with OpenACC (PR92843))

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