aboutsummaryrefslogtreecommitdiff
diff options
context:
space:
mode:
authorJulian Brown <julian@codesourcery.com>2019-09-20 13:53:10 -0700
committerSandra Loosemore <sloosemore@baylibre.com>2025-05-15 20:25:46 +0000
commit3f1fd7de5e4de061acfeffb07d37b2f9b5c78d16 (patch)
treecb8071584e99ca00d2fa039211b3cc217928a483
parent318087911456a67455b53d86f4be882f6db0c5c7 (diff)
downloadgcc-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>
-rw-r--r--gcc/gimplify.cc16
-rw-r--r--gcc/testsuite/gfortran.dg/goacc/privatization-1-compute-loop.f906
-rw-r--r--gcc/testsuite/gfortran.dg/goacc/privatization-1-compute.f9012
-rw-r--r--libgomp/testsuite/libgomp.oacc-c++/privatized-ref-3.C8
-rw-r--r--libgomp/testsuite/libgomp.oacc-fortran/optional-private.f904
-rw-r--r--libgomp/testsuite/libgomp.oacc-fortran/privatized-ref-1.f958
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