[gcc/devel/omp/gcc-16] Handle references in OpenACC "private" clauses
Sandra Loosemore via Gcc-cvs <[email protected]>
| Newsgroups | gmane.comp.gcc.cvs |
|---|---|
| Message-ID | <[email protected]> |
https://gcc.gnu.org/g:ee9ecdb03b378d271b5a4c700740452ec93d1913 commit ee9ecdb03b378d271b5a4c700740452ec93d1913 Author: Julian Brown <[email protected]> Date: Fri Sep 20 13:53:10 2019 -0700 Handle references in OpenACC "private" clauses Combination of OG14 commits 141a592bf147c91c28de7864fa12259687e827e3 8d7562192cc814c6d0d48b424d7751762871a37b + new testsuite fixes to add xfails for tests that already failed on OG14. gcc/ChangeLog * gimplify.cc (localize_reductions): Rewrite references for OMP_CLAUSE_PRIVATE also. Do not create local variable for privatized arrays as the size is not directly known by the type. gcc/testsuite/ChangeLog * gfortran.dg/goacc/privatization-1-compute-loop.f90: Add xfails. * gfortran.dg/goacc/privatization-1-compute.f90: Likewise. libgomp/ChangeLog * testsuite/libgomp.oacc-c++/privatized-ref-3.C: Add xfails. * testsuite/libgomp.oacc-fortran/optional-private.f90: Likewise. * testsuite/libgomp.oacc-fortran/privatized-ref-1.f95: Likewise. Co-Authored-By: Tobias Burnus <[email protected]> Co-Authored-By: Sandra Loosemore <[email protected]> Diff: --- gcc/gimplify.cc | 16 ++++++++++++++++ .../gfortran.dg/goacc/privatization-1-compute-loop.f90 | 6 ++---- .../gfortran.dg/goacc/privatization-1-compute.f90 | 12 ++++++------ libgomp/testsuite/libgomp.oacc-c++/privatized-ref-3.C | 8 ++++---- .../testsuite/libgomp.oacc-fortran/optional-private.f90 | 4 ++-- .../testsuite/libgomp.oacc-fortran/privatized-ref-1.f95 | 8 ++++---- 6 files changed, 34 insertions(+), 20 deletions(-) diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc index 4ecded552408..d4120308bb75 100644 --- a/gcc/gimplify.cc +++ b/gcc/gimplify.cc @@ -17416,6 +17416,22 @@ localize_reductions (tree clauses, tree body) OMP_CLAUSE_REDUCTION_PRIVATE_DECL (c) = new_var; } + else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_PRIVATE) + { + var = OMP_CLAUSE_DECL (c); + + if (!lang_hooks.decls.omp_privatize_by_reference (var)) + continue; + type = TREE_TYPE (TREE_TYPE (var)); + if (TREE_CODE (type) == ARRAY_TYPE) + continue; + new_var = create_tmp_var (type, IDENTIFIER_POINTER (DECL_NAME (var))); + + pr.ref_var = var; + pr.local_var = new_var; + + walk_tree (&body, localize_reductions_r, &pr, NULL); + } } diff --git a/gcc/testsuite/gfortran.dg/goacc/privatization-1-compute-loop.f90 b/gcc/testsuite/gfortran.dg/goacc/privatization-1-compute-loop.f90 index ad5e11abf913..c3fc774c9835 100644 --- a/gcc/testsuite/gfortran.dg/goacc/privatization-1-compute-loop.f90 +++ b/gcc/testsuite/gfortran.dg/goacc/privatization-1-compute-loop.f90 @@ -40,7 +40,8 @@ contains ! (See C/C++ example.) a = g (i, j, a, c) - ! { dg-warning {'a' is used uninitialized} TODO { xfail *-*-* } .-1 } + ! { dg-warning {'a\.[0-9]+' is used uninitialized} "" { target *-*-* } .-1 } + ! { dg-note {'a\.[0-9]+' was declared here} "" { target *-*-* } l_loop$c_loop } x = a !$acc atomic write y = a @@ -51,9 +52,6 @@ contains ! { dg-note {variable 'j\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } ! { dg-note {variable 'i\.[0-9]+' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } ! { dg-note {variable 'j\.[0-9]+' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } - ! { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } - ! { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } - ! { dg-note {variable 'a' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } ! { dg-note {variable 'x' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } ! { dg-note {variable 'y' in 'private' clause is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_loop$c_loop } ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_loop$c_loop } diff --git a/gcc/testsuite/gfortran.dg/goacc/privatization-1-compute.f90 b/gcc/testsuite/gfortran.dg/goacc/privatization-1-compute.f90 index 68d084dd492b..d4d548afee2f 100644 --- a/gcc/testsuite/gfortran.dg/goacc/privatization-1-compute.f90 +++ b/gcc/testsuite/gfortran.dg/goacc/privatization-1-compute.f90 @@ -37,12 +37,12 @@ contains ! (See C/C++ example.) a = g (i, j, a, c) - ! { dg-warning {'i' is used uninitialized} {} { target *-*-* } .-1 } - ! { dg-note {'i' was declared here} {} { target *-*-* } l_function$c_function } - ! { dg-warning {'j' is used uninitialized} {} { target *-*-* } .-3 } - ! { dg-note {'j' was declared here} {} { target *-*-* } l_function$c_function } - ! { dg-warning {'a' is used uninitialized} {} { target *-*-* } .-5 } - ! { dg-note {'a' was declared here} {} { target *-*-* } l_function$c_function } + ! { dg-warning {'i\.[0-9]+' is used uninitialized} {} { target *-*-* } .-1 } + ! { dg-note {'i\.[0-9]+' was declared here} {} { target *-*-* } l_compute$c_compute } + ! { dg-warning {'j\.[0-9]+' is used uninitialized} {} { target *-*-* } .-3 } + ! { dg-note {'j\.[0-9]+' was declared here} {} { target *-*-* } l_compute$c_compute } + ! { dg-warning {'a\.[0-9]+' is used uninitialized} {} { target *-*-* } .-5 } + ! { dg-note {'a\.[0-9]+' was declared here} {} { target *-*-* } l_compute$c_compute } x = a !$acc atomic write ! ... to force 'TREE_ADDRESSABLE'. y = a diff --git a/libgomp/testsuite/libgomp.oacc-c++/privatized-ref-3.C b/libgomp/testsuite/libgomp.oacc-c++/privatized-ref-3.C index 11e1cefbc396..5c70260941f9 100644 --- a/libgomp/testsuite/libgomp.oacc-c++/privatized-ref-3.C +++ b/libgomp/testsuite/libgomp.oacc-c++/privatized-ref-3.C @@ -47,7 +47,7 @@ void gangs (void) int tmpvar; int &tmpref = tmpvar; #pragma acc loop collapse(2) gang private(tmpref) /* { dg-line l_loop[incr c_loop] } */ - /* { dg-note {variable 'tmpref' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */ + /* { dg-note {variable 'tmpref' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { xfail *-*-* } l_loop$c_loop } */ /* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */ /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */ for (i = 0; i < 256; i++) @@ -96,7 +96,7 @@ void workers (void) for (i = 0; i < 256; i++) { #pragma acc loop worker private(tmpref) /* { dg-line l_loop[incr c_loop] } */ - /* { dg-note {variable 'tmpref' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */ + /* { dg-note {variable 'tmpref' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { xfail *-*-* } l_loop$c_loop } */ /* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */ for (j = 0; j < 256; j++) { @@ -142,7 +142,7 @@ void vectors (void) for (i = 0; i < 256; i++) { #pragma acc loop vector private(tmpref) /* { dg-line l_loop[incr c_loop] } */ - /* { dg-note {variable 'tmpref' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */ + /* { dg-note {variable 'tmpref' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { xfail *-*-* } l_loop$c_loop } */ /* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */ for (j = 0; j < 256; j++) { @@ -184,7 +184,7 @@ void gangs_workers_vectors (void) int tmpvar; int &tmpref = tmpvar; #pragma acc loop collapse(2) gang worker vector private(tmpref) /* { dg-line l_loop[incr c_loop] } */ - /* { dg-note {variable 'tmpref' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */ + /* { dg-note {variable 'tmpref' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { xfail *-*-* } l_loop$c_loop } */ /* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */ /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */ for (i = 0; i < 256; i++) diff --git a/libgomp/testsuite/libgomp.oacc-fortran/optional-private.f90 b/libgomp/testsuite/libgomp.oacc-fortran/optional-private.f90 index df693628c962..30a55bc23e4e 100644 --- a/libgomp/testsuite/libgomp.oacc-fortran/optional-private.f90 +++ b/libgomp/testsuite/libgomp.oacc-fortran/optional-private.f90 @@ -44,7 +44,7 @@ contains ! { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .-2 } !$acc loop gang private(x) ! { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-1 } - ! { dg-note {variable 'x' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-2 } + ! { dg-note {variable 'x' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { xfail *-*-* } .-2 } do i = 1, 32 x = i * 2; arr(i) = arr(i) + x @@ -72,7 +72,7 @@ contains ! { dg-warning "region is worker partitioned but does not contain worker partitioned code" "" { target *-*-* } .-1 } !$acc loop gang private(pt) ! { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-1 } - ! { dg-note {variable 'pt' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-2 } + ! { dg-note {variable 'pt' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { xfail *-*-* } .-2 } do i = 0, 31 pt%x = i pt%y = i * 2 diff --git a/libgomp/testsuite/libgomp.oacc-fortran/privatized-ref-1.f95 b/libgomp/testsuite/libgomp.oacc-fortran/privatized-ref-1.f95 index b027d14e7f52..1b3367d1fe1e 100644 --- a/libgomp/testsuite/libgomp.oacc-fortran/privatized-ref-1.f95 +++ b/libgomp/testsuite/libgomp.oacc-fortran/privatized-ref-1.f95 @@ -78,7 +78,7 @@ contains !$acc loop collapse(2) gang private(t1) ! { dg-line l_loop[incr c_loop] } ! { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } ! { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } - ! { dg-note {variable 't1' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } + ! { dg-note {variable 't1' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { xfail *-*-* } l_loop$c_loop } do i=0,255 do j=1,256 t1 = (i * 256 + j) * 97 @@ -103,7 +103,7 @@ contains do i=0,255 !$acc loop worker private(t1) ! { dg-line l_loop[incr c_loop] } ! { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } - ! { dg-note {variable 't1' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } + ! { dg-note {variable 't1' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { xfail *-*-* } l_loop$c_loop } do j=1,256 t1 = (i * 256 + j) * 99 res(i * 256 + j) = t1 @@ -127,7 +127,7 @@ contains do i=0,255 !$acc loop vector private(t1) ! { dg-line l_loop[incr c_loop] } ! { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } - ! { dg-note {variable 't1' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } + ! { dg-note {variable 't1' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { xfail *-*-* } l_loop$c_loop } do j=1,256 t1 = (i * 256 + j) * 101 res(i * 256 + j) = t1 @@ -149,7 +149,7 @@ contains !$acc loop collapse(2) gang worker vector private(t1) ! { dg-line l_loop[incr c_loop] } ! { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } ! { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } - ! { dg-note {variable 't1' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } + ! { dg-note {variable 't1' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { xfail *-*-* } l_loop$c_loop } do i=0,255 do j=1,256 t1 = (i * 256 + j) * 103