+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>
+++ /dev/null
-/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
-/* { dg-additional-options "-DOPENACC_API" } */
-#include "static-dynamic-lifetimes-1.c"
+++ /dev/null
-/* 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;
-}
+++ /dev/null
-/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
-/* { dg-additional-options "-DOPENACC_API" } */
-#include "static-dynamic-lifetimes-2.c"
+++ /dev/null
-/* 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;
-}
+++ /dev/null
-/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
-/* { dg-additional-options "-DOPENACC_API" } */
-#include "static-dynamic-lifetimes-3.c"
+++ /dev/null
-/* 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;
-}
+++ /dev/null
-/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
-/* { dg-additional-options "-DOPENACC_API" } */
-#include "static-dynamic-lifetimes-4.c"
+++ /dev/null
-/* { 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;
-}
+++ /dev/null
-/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
-/* { dg-additional-options "-DOPENACC_API" } */
-#include "static-dynamic-lifetimes-5.c"
+++ /dev/null
-/* { 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;
-}
+++ /dev/null
-/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
-/* { dg-additional-options "-DOPENACC_API" } */
-#include "static-dynamic-lifetimes-6.c"
+++ /dev/null
-/* { 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;
-}
+++ /dev/null
-/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
-/* { dg-additional-options "-DOPENACC_API" } */
-#include "static-dynamic-lifetimes-7.c"
+++ /dev/null
-/* { 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;
-}
+++ /dev/null
-/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
-/* { dg-additional-options "-DOPENACC_API" } */
-#include "static-dynamic-lifetimes-8.c"
+++ /dev/null
-/* { 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;
-}
--- /dev/null
+/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
+/* { dg-additional-options "-DOPENACC_API" } */
+#include "structured-dynamic-lifetimes-1.c"
--- /dev/null
+/* Test transitioning of data lifetimes between structured 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 structured 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;
+}
--- /dev/null
+/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
+/* { dg-additional-options "-DOPENACC_API" } */
+#include "structured-dynamic-lifetimes-2.c"
--- /dev/null
+/* Test nested dynamic/structured 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;
+}
--- /dev/null
+/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
+/* { dg-additional-options "-DOPENACC_API" } */
+#include "structured-dynamic-lifetimes-3.c"
--- /dev/null
+/* Test nested dynamic/structured 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;
+}
--- /dev/null
+/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
+/* { dg-additional-options "-DOPENACC_API" } */
+#include "structured-dynamic-lifetimes-4.c"
--- /dev/null
+/* { 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;
+}
--- /dev/null
+/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
+/* { dg-additional-options "-DOPENACC_API" } */
+#include "structured-dynamic-lifetimes-5.c"
--- /dev/null
+/* { 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;
+}
--- /dev/null
+/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
+/* { dg-additional-options "-DOPENACC_API" } */
+#include "structured-dynamic-lifetimes-6.c"
--- /dev/null
+/* { 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 structured 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;
+}
--- /dev/null
+/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
+/* { dg-additional-options "-DOPENACC_API" } */
+#include "structured-dynamic-lifetimes-7.c"
--- /dev/null
+/* { 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 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 structured 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;
+}
--- /dev/null
+/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
+/* { dg-additional-options "-DOPENACC_API" } */
+#include "structured-dynamic-lifetimes-8.c"
--- /dev/null
+/* { 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;
+}
projects / gcc.git / commitdiff