From: Thomas Schwinge Date: Mon, 13 Apr 2020 06:56:03 +0000 (+0200) Subject: Rename 'libgomp.oacc-c-c++-common/static-dynamic-lifetimes-*' to 'libgomp.oacc-c... X-Git-Url: https://git.libre-soc.org/?a=commitdiff_plain;h=af4c92573dc462a17a6c345756889d28054ed591;p=gcc.git Rename 'libgomp.oacc-c-c++-common/static-dynamic-lifetimes-*' to 'libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-*' [PR92843] Fix-up for commit be9862dd96945772ae0692bc95b37ec6dbcabda0 "Test cases for mixed structured/dynamic data lifetimes with OpenACC [PR92843]": it's "structured", not "static" data lifetimes/reference counters. libgomp/ PR libgomp/92843 * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-1-lib.c: Rename to... * testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-1-lib.c: ... this. * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-1.c: Rename to... * testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-1.c: ... this. * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-2-lib.c: Rename to... * testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-2-lib.c: ... this. * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-2.c: Rename to... * testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-2.c: ... this. * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-3-lib.c: Rename to... * testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-3-lib.c: ... this. * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-3.c: Rename to... * testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-3.c: ... this. * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-4-lib.c: Rename to... * testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-4-lib.c: ... this. * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-4.c: Rename to... * testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-4.c: ... this. * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-5-lib.c: Rename to... * testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-5-lib.c: ... this. * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-5.c: Rename to... * testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-5.c: ... this. * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-6-lib.c: Rename to... * testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-6-lib.c: ... this. * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-6.c: Rename to... * testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-6.c: ... this. * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-7-lib.c: Rename to... * testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-7-lib.c: ... this. * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-7.c: Rename to... * testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-7.c: ... this. * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-8-lib.c: Rename to... * testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-8-lib.c: ... this. * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-8.c:: Rename to... * testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-8.c: ... this. --- diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog index b0f19845ad2..8f351f61f33 100644 --- a/libgomp/ChangeLog +++ b/libgomp/ChangeLog @@ -1,3 +1,71 @@ +2020-04-13 Thomas Schwinge + + 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 Thomas Schwinge 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 deleted file mode 100644 index 23c20d4fab7..00000000000 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-1-lib.c +++ /dev/null @@ -1,3 +0,0 @@ -/* { 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 deleted file mode 100644 index a743660f53e..00000000000 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-1.c +++ /dev/null @@ -1,160 +0,0 @@ -/* Test transitioning of data lifetimes between static and dynamic. */ - -/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */ - -#include -#include -#include - -#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 deleted file mode 100644 index 84f41a49dfd..00000000000 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-2-lib.c +++ /dev/null @@ -1,3 +0,0 @@ -/* { 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 deleted file mode 100644 index d3c6f5192d8..00000000000 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-2.c +++ /dev/null @@ -1,166 +0,0 @@ -/* Test nested dynamic/static data mappings. */ - -/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */ - -#include -#include -#include - -#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 deleted file mode 100644 index d9e76c600f0..00000000000 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-3-lib.c +++ /dev/null @@ -1,3 +0,0 @@ -/* { 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 deleted file mode 100644 index 59501864398..00000000000 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-3.c +++ /dev/null @@ -1,183 +0,0 @@ -/* Test nested dynamic/static data mappings (multiple blocks on data - regions). */ - -/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */ - -#include -#include -#include - -#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 deleted file mode 100644 index e3c1bfb473d..00000000000 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-4-lib.c +++ /dev/null @@ -1,3 +0,0 @@ -/* { 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 deleted file mode 100644 index e9a6510ace8..00000000000 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-4.c +++ /dev/null @@ -1,64 +0,0 @@ -/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */ - -#include -#include -#include - -#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 deleted file mode 100644 index 77703122ad6..00000000000 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-5-lib.c +++ /dev/null @@ -1,3 +0,0 @@ -/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */ -/* { dg-additional-options "-DOPENACC_API" } */ -#include "static-dynamic-lifetimes-5.c" diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-5.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-5.c deleted file mode 100644 index 9807076d3f4..00000000000 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-5.c +++ /dev/null @@ -1,56 +0,0 @@ -/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */ - -#include -#include -#include - -#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 deleted file mode 100644 index 4a87dd72525..00000000000 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-6-lib.c +++ /dev/null @@ -1,3 +0,0 @@ -/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */ -/* { dg-additional-options "-DOPENACC_API" } */ -#include "static-dynamic-lifetimes-6.c" diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-6.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-6.c deleted file mode 100644 index 3e5c4d7ea11..00000000000 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-6.c +++ /dev/null @@ -1,42 +0,0 @@ -/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */ - -#include -#include -#include - -#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 deleted file mode 100644 index 8ccbb126933..00000000000 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-7-lib.c +++ /dev/null @@ -1,3 +0,0 @@ -/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */ -/* { dg-additional-options "-DOPENACC_API" } */ -#include "static-dynamic-lifetimes-7.c" diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-7.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-7.c deleted file mode 100644 index 2735d6fa0eb..00000000000 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-7.c +++ /dev/null @@ -1,42 +0,0 @@ -/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */ - -#include -#include -#include - -#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 deleted file mode 100644 index f3104cbd035..00000000000 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-8-lib.c +++ /dev/null @@ -1,3 +0,0 @@ -/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */ -/* { dg-additional-options "-DOPENACC_API" } */ -#include "static-dynamic-lifetimes-8.c" diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-8.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-8.c deleted file mode 100644 index 919ee02b725..00000000000 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-8.c +++ /dev/null @@ -1,47 +0,0 @@ -/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */ - -#include -#include -#include - -#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; -} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-1-lib.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-1-lib.c new file mode 100644 index 00000000000..8fa87777f1a --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-1-lib.c @@ -0,0 +1,3 @@ +/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */ +/* { dg-additional-options "-DOPENACC_API" } */ +#include "structured-dynamic-lifetimes-1.c" diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-1.c new file mode 100644 index 00000000000..0d6b4159ad0 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-1.c @@ -0,0 +1,161 @@ +/* Test transitioning of data lifetimes between structured and dynamic. */ + +/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */ + +#include +#include +#include + +#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; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-2-lib.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-2-lib.c new file mode 100644 index 00000000000..365df8d7f7d --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-2-lib.c @@ -0,0 +1,3 @@ +/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */ +/* { dg-additional-options "-DOPENACC_API" } */ +#include "structured-dynamic-lifetimes-2.c" diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-2.c new file mode 100644 index 00000000000..726942c76fd --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-2.c @@ -0,0 +1,166 @@ +/* Test nested dynamic/structured data mappings. */ + +/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */ + +#include +#include +#include + +#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/structured-dynamic-lifetimes-3-lib.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-3-lib.c new file mode 100644 index 00000000000..469b35b76f1 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-3-lib.c @@ -0,0 +1,3 @@ +/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */ +/* { dg-additional-options "-DOPENACC_API" } */ +#include "structured-dynamic-lifetimes-3.c" diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-3.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-3.c new file mode 100644 index 00000000000..c13f3c56584 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-3.c @@ -0,0 +1,183 @@ +/* Test nested dynamic/structured data mappings (multiple blocks on data + regions). */ + +/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */ + +#include +#include +#include + +#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/structured-dynamic-lifetimes-4-lib.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-4-lib.c new file mode 100644 index 00000000000..8e88b97b2ce --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-4-lib.c @@ -0,0 +1,3 @@ +/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */ +/* { dg-additional-options "-DOPENACC_API" } */ +#include "structured-dynamic-lifetimes-4.c" diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-4.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-4.c new file mode 100644 index 00000000000..e9a6510ace8 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-4.c @@ -0,0 +1,64 @@ +/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */ + +#include +#include +#include + +#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/structured-dynamic-lifetimes-5-lib.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-5-lib.c new file mode 100644 index 00000000000..59ef5626bd0 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-5-lib.c @@ -0,0 +1,3 @@ +/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */ +/* { dg-additional-options "-DOPENACC_API" } */ +#include "structured-dynamic-lifetimes-5.c" diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-5.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-5.c new file mode 100644 index 00000000000..9807076d3f4 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-5.c @@ -0,0 +1,56 @@ +/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */ + +#include +#include +#include + +#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/structured-dynamic-lifetimes-6-lib.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-6-lib.c new file mode 100644 index 00000000000..0401f73966f --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-6-lib.c @@ -0,0 +1,3 @@ +/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */ +/* { dg-additional-options "-DOPENACC_API" } */ +#include "structured-dynamic-lifetimes-6.c" diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-6.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-6.c new file mode 100644 index 00000000000..9250b4af3ed --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-6.c @@ -0,0 +1,43 @@ +/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */ + +#include +#include +#include + +#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; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-7-lib.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-7-lib.c new file mode 100644 index 00000000000..07caefbd082 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-7-lib.c @@ -0,0 +1,3 @@ +/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */ +/* { dg-additional-options "-DOPENACC_API" } */ +#include "structured-dynamic-lifetimes-7.c" diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-7.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-7.c new file mode 100644 index 00000000000..52e8d4c9959 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-7.c @@ -0,0 +1,44 @@ +/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */ + +#include +#include +#include + +#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; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-8-lib.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-8-lib.c new file mode 100644 index 00000000000..1c2479ad96f --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-8-lib.c @@ -0,0 +1,3 @@ +/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */ +/* { dg-additional-options "-DOPENACC_API" } */ +#include "structured-dynamic-lifetimes-8.c" diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-8.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-8.c new file mode 100644 index 00000000000..919ee02b725 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-8.c @@ -0,0 +1,47 @@ +/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */ + +#include +#include +#include + +#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; +}