This is the mail archive of the
gcc-patches@gcc.gnu.org
mailing list for the GCC project.
[gomp4] reductions on explicitly private variables
- From: Cesar Philippidis <cesar at codesourcery dot com>
- To: "gcc-patches at gcc dot gnu dot org" <gcc-patches at gcc dot gnu dot org>
- Date: Thu, 7 Jan 2016 20:53:57 -0800
- Subject: [gomp4] reductions on explicitly private variables
- Authentication-results: sourceware.org; auth=none
I noticed that libgomp is lacking test coverage reductions involving
explicitly private variables. E.g.
#pragma acc parallel copyout(red)
{
#pragma acc loop gang private(v)
for (j = 0; j < 10; j++)
{
v = j;
#pragma acc loop vector reduction (+:v)
for (i = 0; i < 100; i++)
v++;
red[j] = v;
}
}
This used to case a problem in lower_oacc_reductions because the
reduction finalizer wasn't updating the correct variable. The attached
patch fixes it.
I've applied this patch to gomp-4_0-branch.
Cesar
2016-01-07 Cesar Philippidis <cesar@codesourcery.com>
gcc/
* omp-low.c (lower_oacc_reductions): Handle explicitly private
reduction varaibles.
libgomp/
* testsuite/libgomp.oacc-c-c++-common/reduction-6.c: New test.
* testsuite/libgomp.oacc-fortran/reduction-7.f90: Add a test with
an explicitly private reduction variable.
diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index bbfd642..0be25b8 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -5616,9 +5616,15 @@ lower_oacc_reductions (location_t loc, tree clauses, tree level, bool inner,
outer = probe;
for (; cls; cls = OMP_CLAUSE_CHAIN (cls))
- if (OMP_CLAUSE_CODE (cls) == OMP_CLAUSE_REDUCTION
- && orig == OMP_CLAUSE_DECL (cls))
- goto has_outer_reduction;
+ {
+ if (OMP_CLAUSE_CODE (cls) == OMP_CLAUSE_REDUCTION
+ && orig == OMP_CLAUSE_DECL (cls))
+ goto has_outer_reduction;
+ else if (OMP_CLAUSE_CODE (cls) == OMP_CLAUSE_PRIVATE
+ && orig == OMP_CLAUSE_DECL (cls)
+ && gimple_code (outer->stmt) != GIMPLE_OMP_TARGET)
+ goto has_outer_reduction;
+ }
}
do_lookup:
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/reduction-6.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/reduction-6.c
new file mode 100644
index 0000000..af30b31
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/reduction-6.c
@@ -0,0 +1,36 @@
+/* { dg-do run } */
+/* { dg-additional-options "-w" } */
+
+/* Test reductions on explicitly private variables. */
+
+#include <assert.h>
+
+int
+main ()
+{
+ int i, j, red[10];
+ int v;
+
+ for (i = 0; i < 10; i++)
+ red[i] = -1;
+
+#pragma acc parallel copyout(red)
+ {
+#pragma acc loop gang private(v)
+ for (j = 0; j < 10; j++)
+ {
+ v = j;
+
+#pragma acc loop vector reduction (+:v)
+ for (i = 0; i < 100; i++)
+ v++;
+
+ red[j] = v;
+ }
+ }
+
+ for (i = 0; i < 10; i++)
+ assert (red[i] == i + 100);
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/reduction-7.f90 b/libgomp/testsuite/libgomp.oacc-fortran/reduction-7.f90
index e80004d..8ec36ad 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/reduction-7.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/reduction-7.f90
@@ -1,12 +1,13 @@
! { dg-do run }
! { dg-additional-options "-w" }
-! subroutine reduction with firstprivate variables
+! subroutine reduction with private and firstprivate variables
program reduction
integer, parameter :: n = 100
integer :: i, j, vsum, cs, arr(n)
+ call redsub_private (cs, n, arr)
call redsub_bogus (cs, n)
call redsub_combined (cs, n, arr)
@@ -22,6 +23,33 @@ program reduction
end do
end program reduction
+! This subroutine tests a reduction with an explicit private variable.
+
+subroutine redsub_private(sum, n, arr)
+ integer :: sum, n, arr(n)
+ integer :: i, j, v
+
+ !$acc parallel copyout (arr)
+ !$acc loop gang private (v)
+ do j = 1, n
+ v = j
+
+ !$acc loop vector reduction (+:v)
+ do i = 1, 100
+ v = v + 1
+ end do
+
+ arr(j) = v
+ end do
+ !$acc end parallel
+
+ ! verify the results
+ do i = 1, 10
+ if (arr(i) .ne. 100+i) call abort ()
+ end do
+end subroutine redsub_private
+
+
! Bogus reduction on an impliclitly firstprivate variable. The results do
! survive the parallel region. The goal here is to ensure that gfortran
! doesn't ICE.