diff options
author | Julian Brown <julian@codesourcery.com> | 2020-01-17 13:18:18 -0800 |
---|---|---|
committer | Thomas Schwinge <thomas@codesourcery.com> | 2020-04-10 15:53:04 +0200 |
commit | be9862dd96945772ae0692bc95b37ec6dbcabda0 (patch) | |
tree | f38e93522543be0066d2319d5d22c87c39d0ea14 /libgomp | |
parent | 6b816a5f0ed078cb2d380e10e68a95fb7e3d6778 (diff) | |
download | gcc-be9862dd96945772ae0692bc95b37ec6dbcabda0.zip gcc-be9862dd96945772ae0692bc95b37ec6dbcabda0.tar.gz gcc-be9862dd96945772ae0692bc95b37ec6dbcabda0.tar.bz2 |
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.
Diffstat (limited to 'libgomp')
17 files changed, 821 insertions, 0 deletions
diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog index beff3d6..b0f1984 100644 --- a/libgomp/ChangeLog +++ b/libgomp/ChangeLog @@ -1,3 +1,40 @@ +2020-04-10 Julian Brown <julian@codesourcery.com> + Thomas Schwinge <thomas@codesourcery.com> + + PR libgomp/92843 + * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-1-lib.c: + New file. + * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-1.c: + Likewise. + * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-2-lib.c: + Likewise. + * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-2.c: + Likewise. + * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-3-lib.c: + Likewise. + * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-3.c: + Likewise. + * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-4-lib.c: + Likewise. + * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-4.c: + Likewise. + * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-5-lib.c: + Likewise. + * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-5.c: + Likewise. + * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-6-lib.c: + Likewise. + * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-6.c: + Likewise. + * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-7-lib.c: + Likewise. + * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-7.c: + Likewise. + * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-8-lib.c: + Likewise. + * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-8.c: + Likewise. + 2020-04-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 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 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 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 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 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 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 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 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 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 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 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 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 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 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 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 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; +} |