diff options
author | Julian Brown <julian@codesourcery.com> | 2019-09-20 13:53:10 -0700 |
---|---|---|
committer | Sandra Loosemore <sloosemore@baylibre.com> | 2025-05-15 20:25:46 +0000 |
commit | 3f1fd7de5e4de061acfeffb07d37b2f9b5c78d16 (patch) | |
tree | cb8071584e99ca00d2fa039211b3cc217928a483 | |
parent | 318087911456a67455b53d86f4be882f6db0c5c7 (diff) | |
download | gcc-3f1fd7de5e4de061acfeffb07d37b2f9b5c78d16.zip gcc-3f1fd7de5e4de061acfeffb07d37b2f9b5c78d16.tar.gz gcc-3f1fd7de5e4de061acfeffb07d37b2f9b5c78d16.tar.bz2 |
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 <tobias@codesourcery.com>
Co-Authored-By: Sandra Loosemore <sandra@baylibre.com>
6 files changed, 34 insertions, 20 deletions
diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc index 2a351b0..bd45816 100644 --- a/gcc/gimplify.cc +++ b/gcc/gimplify.cc @@ -16243,6 +16243,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 ad5e11a..c3fc774 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 68d084d..d4d548a 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 11e1cef..5c70260 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 df69362..30a55bc 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 b027d14..1b3367d 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 |