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

+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

libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-1-lib.c

+/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
+/* { dg-additional-options "-DOPENACC_API" } */
+#include "static-dynamic-lifetimes-1.c"

libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-1.c

+/* 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;
+}

libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-2-lib.c

+/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
+/* { dg-additional-options "-DOPENACC_API" } */
+#include "static-dynamic-lifetimes-2.c"

libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-2.c

+/* 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;
+}

libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-3-lib.c

+/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
+/* { dg-additional-options "-DOPENACC_API" } */
+#include "static-dynamic-lifetimes-3.c"

libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-3.c

+/* 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;
+}

libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-4-lib.c

+/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
+/* { dg-additional-options "-DOPENACC_API" } */
+#include "static-dynamic-lifetimes-4.c"

libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-4.c

+/* { 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;
+}

libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-5-lib.c

+/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
+/* { dg-additional-options "-DOPENACC_API" } */
+#include "static-dynamic-lifetimes-5.c"

libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-5.c

+/* { 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;
+}

libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-6-lib.c

+/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
+/* { dg-additional-options "-DOPENACC_API" } */
+#include "static-dynamic-lifetimes-6.c"

libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-6.c

+/* { 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;
+}

libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-7-lib.c

+/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
+/* { dg-additional-options "-DOPENACC_API" } */
+#include "static-dynamic-lifetimes-7.c"

libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-7.c

+/* { 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;
+}

libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-8-lib.c

+/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
+/* { dg-additional-options "-DOPENACC_API" } */
+#include "static-dynamic-lifetimes-8.c"

libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-8.c

+/* { 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;
+}