[gomp4] reductions on explicitly private variables

[Cesar Philippidis <cesar at codesourcery dot com>]

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.