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.

diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog
index beff3d65b4417cee9fb79bec561fcd01a4b418af..b0f19845ad268393bc021a18df563430f29b4322 100644 (file)
--- 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 (file)
index 0000000..23c20d4
--- /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 (file)
index 0000000..a743660
--- /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 (file)
index 0000000..84f41a4
--- /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 (file)
index 0000000..d3c6f51
--- /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 (file)
index 0000000..d9e76c6
--- /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 (file)
index 0000000..5950186
--- /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 (file)
index 0000000..e3c1bfb
--- /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 (file)
index 0000000..e9a6510
--- /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 (file)
index 0000000..7770312
--- /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 (file)
index 0000000..9807076
--- /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 (file)
index 0000000..4a87dd7
--- /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 (file)
index 0000000..3e5c4d7
--- /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 (file)
index 0000000..8ccbb12
--- /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 (file)
index 0000000..2735d6f
--- /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 (file)
index 0000000..f3104cb
--- /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 (file)
index 0000000..919ee02
--- /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;
+}