From: Julian Brown Date: Fri, 17 Jan 2020 21:18:18 +0000 (-0800) Subject: Test cases for mixed structured/dynamic data lifetimes with OpenACC [PR92843] X-Git-Url: https://git.libre-soc.org/?a=commitdiff_plain;h=be9862dd96945772ae0692bc95b37ec6dbcabda0;p=gcc.git Test cases for mixed structured/dynamic data lifetimes with OpenACC [PR92843] libgomp/ PR libgomp/92843 * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-1-lib.c: New file. * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-1.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-2-lib.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-2.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-3-lib.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-3.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-4-lib.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-4.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-5-lib.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-5.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-6-lib.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-6.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-7-lib.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-7.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-8-lib.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-8.c: Likewise. --- diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog index beff3d65b44..b0f19845ad2 100644 --- a/libgomp/ChangeLog +++ b/libgomp/ChangeLog @@ -1,3 +1,40 @@ +2020-04-10 Julian Brown + Thomas Schwinge + + 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 * 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 00000000000..23c20d4fab7 --- /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 00000000000..a743660f53e --- /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 +#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 new file mode 100644 index 00000000000..84f41a49dfd --- /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 00000000000..d3c6f5192d8 --- /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 +#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 new file mode 100644 index 00000000000..d9e76c600f0 --- /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 00000000000..59501864398 --- /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 +#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 new file mode 100644 index 00000000000..e3c1bfb473d --- /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 00000000000..e9a6510ace8 --- /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 +#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 new file mode 100644 index 00000000000..77703122ad6 --- /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 00000000000..9807076d3f4 --- /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 +#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 new file mode 100644 index 00000000000..4a87dd72525 --- /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 00000000000..3e5c4d7ea11 --- /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 +#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 new file mode 100644 index 00000000000..8ccbb126933 --- /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 00000000000..2735d6fa0eb --- /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 +#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 new file mode 100644 index 00000000000..f3104cbd035 --- /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 00000000000..919ee02b725 --- /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 +#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; +}