diff options
author | Nathan Sidwell <nathan@codesourcery.com> | 2015-11-11 14:24:09 +0000 |
---|---|---|
committer | Nathan Sidwell <nathan@gcc.gnu.org> | 2015-11-11 14:24:09 +0000 |
commit | 182190f2b3a236eb30bd717981e1f2c9b51cea06 (patch) | |
tree | 295e38d6632d00a37b02d200a9f7a761f46af3f8 /libgomp | |
parent | 7700cd858ff97a257d6f48b0d079780f445821da (diff) | |
download | gcc-182190f2b3a236eb30bd717981e1f2c9b51cea06.zip gcc-182190f2b3a236eb30bd717981e1f2c9b51cea06.tar.gz gcc-182190f2b3a236eb30bd717981e1f2c9b51cea06.tar.bz2 |
gimplify.c (enum omp_region_type): Add ORT_ACC, ORT_ACC_DATA, ORT_ACC_PARALLEL, ORT_ACC_KERNELS.
gcc/
* gcc/gimplify.c (enum omp_region_type): Add ORT_ACC,
ORT_ACC_DATA, ORT_ACC_PARALLEL, ORT_ACC_KERNELS. Adjust ORT_NONE.
(gimple_add_tmp_var): Add ORT_ACC checks.
(gimplify_var_or_parm_decl): Likewise.
(omp_firstprivatize_variable): Likewise. Use ORT_TARGET_DATA as a
mask.
(omp_add_variable): Look in outer contexts for openacc and allow
reductions with other sharing. Add ORT_ACC and ORT_TARGET_DATA
checks.
(omp_notice_variable, omp_is_private, omp_check_private): Add
ORT_ACC checks.
(gimplify_scan_omp_clauses: Treat ORT_ACC as ORT_WORKSHARE.
Permit private openacc reductions.
(gimplify_oacc_cache): Specify ORT_ACC.
(gimplify_omp_workshare): Adjust OpenACC region types.
(gimplify_omp_target_update): Likewise.
* gcc/omp-low.c (scan_sharing_clauses): Remove Openacc
firstprivate sorry.
(lower-rec_input_clauses): Don't handle openacc firstprivate
references here.
(lower_omp_target): Emit initializers for openacc firstprivate vars.
gcc/testsuite/
* gfortran.dg/goacc/private-3.f95: Remove xfail.
* gfortran.dg/goacc/combined_loop.f90: Remove xfail.
libgomp/
* testsuite/libgomp.oacc-c-c++-common/loop-red-v-2.c: Remove xfail.
* testsuite/libgomp.oacc-c-c++-common/loop-red-w-2.c: Remove xfail.
* testsuite/libgomp.oacc-c-c++-common/firstprivate-1.c: New.
* testsuite/libgomp.oacc-c-c++-common/firstprivate-2.c: New.
Co-Authored-By: Cesar Philippidis <cesar@codesourcery.com>
From-SVN: r230169
Diffstat (limited to 'libgomp')
5 files changed, 77 insertions, 4 deletions
diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog index ed86943..406c572 100644 --- a/libgomp/ChangeLog +++ b/libgomp/ChangeLog @@ -1,3 +1,8 @@ +2015-11-1 Nathan Sidwell <nathan@codesourcery.com> + + * testsuite/libgomp.oacc-c-c++-common/firstprivate-1.c: New. + * testsuite/libgomp.oacc-c-c++-common/firstprivate-2.c: New. + 2015-11-09 Nathan Sidwell <nathan@codesourcery.com> * testsuite/libgomp.oacc-c-c++-common/firstprivate-1.c: Remove diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/firstprivate-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/firstprivate-1.c new file mode 100644 index 0000000..7f5d3d3 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/firstprivate-1.c @@ -0,0 +1,41 @@ +/* { dg-do run } */ + +#include <openacc.h> + +int main () +{ + int ok = 1; + int val = 2; + int ary[32]; + int ondev = 0; + + for (int i = 0; i < 32; i++) + ary[i] = ~0; + +#pragma acc parallel num_gangs (32) copy (ok) firstprivate (val) copy(ary, ondev) + { + ondev = acc_on_device (acc_device_not_host); +#pragma acc loop gang(static:1) + for (unsigned i = 0; i < 32; i++) + { + if (val != 2) + ok = 0; + val += i; + ary[i] = val; + } + } + + if (ondev) + { + if (!ok) + return 1; + if (val != 2) + return 1; + + for (int i = 0; i < 32; i++) + if (ary[i] != 2 + i) + return 1; + } + + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/firstprivate-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/firstprivate-2.c new file mode 100644 index 0000000..9666542 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/firstprivate-2.c @@ -0,0 +1,31 @@ +/* { dg-do run } */ + +#include <openacc.h> + +int main () +{ + int ok = 1; + int val = 2; + +#pragma acc data copy(val) + { +#pragma acc parallel present (val) + { + val = 7; + } + +#pragma acc parallel firstprivate (val) copy(ok) + { + ok = val == 7; + val = 9; + } + + } + + if (!ok) + return 1; + if(val != 7) + return 1; + + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-v-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-v-2.c index fbed589..e66732d 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-v-2.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-v-2.c @@ -1,7 +1,5 @@ /* { dg-do run } */ /* { dg-additional-options "-O2" */ -/* <http://news.gmane.org/find-root.php?message_id=%3C563B78B5.5090506%40acm.org%3E> - { dg-xfail-if "TODO" { *-*-* } } */ #include <stdio.h> diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-w-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-w-2.c index 47f1da0..0059077 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-w-2.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-w-2.c @@ -1,7 +1,5 @@ /* { dg-do run } */ /* { dg-additional-options "-O2" */ -/* <http://news.gmane.org/find-root.php?message_id=%3C563B78B5.5090506%40acm.org%3E> - { dg-xfail-if "TODO" { *-*-* } } */ #include <stdio.h> |