2016-04-04  Cesar Philippidis  <cesar@codesourcery.com>

	gcc/
	* gimplify.c (gimplify_adjust_acc_parallel_reductions): New function.
	(gimplify_omp_workshare): Call it.  Add new data clauses for acc
	parallel reductions as needed.

	gcc/testsuite/
	* c-c++-common/goacc/reduction-5.c: New test.
	* c-c++-common/goacc/reduction-promotions.c: New test.
	* gfortran.dg/goacc/reduction-3.f95: New test.
	* gfortran.dg/goacc/reduction-promotions.f90: New test.

	libgomp/
	* testsuite/libgomp.oacc-c-c++-common/pr70289.c: New test.
	* testsuite/libgomp.oacc-c-c++-common/pr70373.c: New test.
	* testsuite/libgomp.oacc-fortran/pr70289.f90: New test.


diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index b9757db..4625881 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -9484,6 +9484,108 @@ optimize_target_teams (tree target, gimple_seq *pre_p)
   OMP_TARGET_CLAUSES (target) = c;
 }
 
+/* OpenACC parallel reductions need a present_or_copy clause to ensure
+   that the original variable used in the reduction gets updated on
+   the host.  This function scans CLAUSES for reductions and adds or
+   adjusts the data clauses as necessary.  Any incompatible data clause
+   will be reported as a warning and promoted to present_or_copy.  Any
+   private reduction will be treated as an error.  This function
+   returns a list of new present_or_copy date clauses.  */
+
+static tree
+gimplify_adjust_acc_parallel_reductions (tree *clauses)
+{
+  tree c, list = NULL_TREE;
+  hash_set<tree> *reduction_decls;
+  reduction_decls = new hash_set<tree>;
+
+  /* Scan 1: Construct a hash set with all of the reduction decls.  */
+  for (c = *clauses; c; c = OMP_CLAUSE_CHAIN (c))
+    {
+      if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_REDUCTION)
+	reduction_decls->add (OMP_CLAUSE_DECL (c));
+    }
+
+  if (reduction_decls->elements () == 0)
+    goto cleanup;
+
+  /* Scan 2: Adjust the data clause for each reduction.  */
+  for (c = *clauses; c; c = OMP_CLAUSE_CHAIN (c))
+    {
+      int kind = -1;
+      tree decl;
+
+      switch (OMP_CLAUSE_CODE (c))
+	{
+	case OMP_CLAUSE_MAP:
+	  kind = OMP_CLAUSE_MAP_KIND (c);
+	case OMP_CLAUSE_PRIVATE:
+	case OMP_CLAUSE_FIRSTPRIVATE:
+	  decl = OMP_CLAUSE_DECL (c);
+
+	  /* Reference variables always have a GOMP_MAP_ALLOC.  Ignore it.  */
+	  if (POINTER_TYPE_P (TREE_TYPE (decl))
+	      && kind == GOMP_MAP_ALLOC)
+	    break;
+
+	  if (!DECL_P (decl))
+	    decl = TREE_OPERAND (decl, 0);
+	  gcc_assert (DECL_P (decl));
+
+	  if (!reduction_decls->contains (decl))
+	    break;
+
+	  if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP)
+	    {
+	      if (!((kind & GOMP_MAP_TOFROM) == GOMP_MAP_TOFROM
+		    || kind == GOMP_MAP_FORCE_PRESENT))
+		{
+		  warning_at (OMP_CLAUSE_LOCATION (c), 0, "incompatible data "
+			      "clause with reduction on %qE; promoting to "
+			      "present_or_copy", DECL_NAME (decl));
+
+		  OMP_CLAUSE_CODE (c) = OMP_CLAUSE_MAP;
+		  OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_TOFROM);
+		}
+	      reduction_decls->remove (decl);
+	      break;
+	    }
+
+	  if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_PRIVATE
+	      || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FIRSTPRIVATE)
+	    {
+	      error_at (OMP_CLAUSE_LOCATION (c), "invalid private reduction "
+			  "on %qE", DECL_NAME (decl));
+	      reduction_decls->remove (decl);
+	    }
+	default:;
+	}
+    }
+
+  if (reduction_decls->elements () == 0)
+    goto cleanup;
+  
+  /* Scan 3: Add a present_or_copy clause for any reduction variable which
+     doens't have a data clause already.  */
+
+  for (hash_set<tree>::iterator iter = reduction_decls->begin ();
+       iter != reduction_decls->end (); ++iter)
+    {
+      tree decl = *iter;
+
+      tree nc = build_omp_clause (UNKNOWN_LOCATION, OMP_CLAUSE_MAP);
+      OMP_CLAUSE_SET_MAP_KIND (nc, GOMP_MAP_TOFROM);
+      OMP_CLAUSE_DECL (nc) = decl;
+      TREE_CHAIN (nc) = list;
+      list = nc;
+    }
+
+ cleanup:
+  delete reduction_decls;
+
+  return list;
+}
+
 /* Gimplify the gross structure of several OMP constructs.  */
 
 static void
@@ -9491,6 +9593,7 @@ gimplify_omp_workshare (tree *expr_p, gimple_seq *pre_p)
 {
   tree expr = *expr_p;
   gimple *stmt;
+  tree acc_reductions = NULL_TREE;
   gimple_seq body = NULL;
   enum omp_region_type ort;
 
@@ -9508,6 +9611,8 @@ gimplify_omp_workshare (tree *expr_p, gimple_seq *pre_p)
       break;
     case OACC_PARALLEL:
       ort = ORT_ACC_PARALLEL;
+      acc_reductions
+	= gimplify_adjust_acc_parallel_reductions (&OMP_CLAUSES (expr));
       break;
     case OACC_DATA:
       ort = ORT_ACC_DATA;
@@ -9524,6 +9629,7 @@ gimplify_omp_workshare (tree *expr_p, gimple_seq *pre_p)
     default:
       gcc_unreachable ();
     }
+
   gimplify_scan_omp_clauses (&OMP_CLAUSES (expr), pre_p, ort,
 			     TREE_CODE (expr));
   if (TREE_CODE (expr) == OMP_TARGET)
@@ -9606,6 +9712,41 @@ gimplify_omp_workshare (tree *expr_p, gimple_seq *pre_p)
 
   gimplify_seq_add_stmt (pre_p, stmt);
   *expr_p = NULL_TREE;
+
+  /* Finalize any parallel acc reductions.  */
+  if (acc_reductions)
+    {
+      tree c, nc, t;
+      tree clauses = NULL_TREE;
+
+      c = nc = acc_reductions;
+
+      while (c)
+	{
+	  nc = OMP_CLAUSE_CHAIN (c);
+	  OMP_CLAUSE_CHAIN (c) = NULL_TREE;
+	  lang_hooks.decls.omp_finish_clause (c, pre_p);
+
+	  /* Find the last data clause introduced by omp_finish_decls.  */
+	  for (t = c; TREE_CHAIN (t); t = TREE_CHAIN (t))
+	    ;
+
+	  /* Update the chain of clauses.  */
+	  TREE_CHAIN (t) = clauses;
+	  clauses = c;
+
+	  c = nc;
+	}
+
+      /* Update the list of clauses in the gimple stmt.  */
+      for (t = gimple_omp_target_clauses (stmt); OMP_CLAUSE_CHAIN (t);
+	   t = OMP_CLAUSE_CHAIN (t))
+	;
+
+      OMP_CLAUSE_CHAIN (t) = clauses;
+    }
+
+  return;
 }
 
 /* Gimplify the gross structure of OpenACC enter/exit data, update, and OpenMP
diff --git a/gcc/testsuite/c-c++-common/goacc/reduction-5.c b/gcc/testsuite/c-c++-common/goacc/reduction-5.c
new file mode 100644
index 0000000..74daad3
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/reduction-5.c
@@ -0,0 +1,16 @@
+/* Integer reductions.  */
+
+#define n 1000
+
+int
+main(void)
+{
+  int v1;
+
+#pragma acc parallel reduction(+:v1) private(v1) /* { dg-error "appears more than once in data clauses" } */
+  ;
+#pragma acc parallel reduction(+:v1) firstprivate(v1) /* { dg-error "appears more than once in data clauses" } */
+  ;
+
+  return 0;
+}
diff --git a/gcc/testsuite/c-c++-common/goacc/reduction-promotions.c b/gcc/testsuite/c-c++-common/goacc/reduction-promotions.c
new file mode 100644
index 0000000..4cc09da
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/reduction-promotions.c
@@ -0,0 +1,32 @@
+/* Integer reductions.  */
+
+#define n 1000
+
+int
+main(void)
+{
+  int v1, v2;
+
+#pragma acc parallel reduction(+:v1,v2)
+  ;
+#pragma acc parallel reduction(+:v1,v2) copy(v1,v2)
+  ;
+#pragma acc parallel reduction(+:v1,v2) pcopy(v1,v2)
+  ;
+#pragma acc parallel reduction(+:v1,v2) present(v1,v2)
+  ;
+#pragma acc parallel reduction(+:v1,v2) copyin(v1,v2) /* { dg-warning "incompatible data clause with reduction" } */
+  ;
+#pragma acc parallel reduction(+:v1,v2) pcopyin(v1,v2) /* { dg-warning "incompatible data clause with reduction" } */
+  ;
+#pragma acc parallel reduction(+:v1,v2) copyout(v1,v2) /* { dg-warning "incompatible data clause with reduction" } */
+  ;
+#pragma acc parallel reduction(+:v1,v2) pcopyout(v1,v2) /* { dg-warning "incompatible data clause with reduction" } */
+  ;
+#pragma acc parallel reduction(+:v1,v2) create(v1,v2) /* { dg-warning "incompatible data clause with reduction" } */
+  ;
+#pragma acc parallel reduction(+:v1,v2) pcreate(v1,v2) /* { dg-warning "incompatible data clause with reduction" } */
+  ;
+
+  return 0;
+}
diff --git a/gcc/testsuite/gfortran.dg/goacc/reduction-3.f95 b/gcc/testsuite/gfortran.dg/goacc/reduction-3.f95
new file mode 100644
index 0000000..72f0eb9
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/reduction-3.f95
@@ -0,0 +1,10 @@
+! { dg-do compile }
+
+subroutine foo (ia1)
+integer :: i1, i2
+
+!$acc parallel reduction (+:i1) private(i1) ! { dg-error "invalid private reduction on .i1." }
+!$acc end parallel
+!$acc parallel reduction (+:i2) firstprivate(i2) ! { dg-error "invalid private reduction on .i2." }
+!$acc end parallel
+end subroutine foo
diff --git a/gcc/testsuite/gfortran.dg/goacc/reduction-promotions.f90 b/gcc/testsuite/gfortran.dg/goacc/reduction-promotions.f90
new file mode 100644
index 0000000..2ae1707
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/reduction-promotions.f90
@@ -0,0 +1,45 @@
+! Ensure that each parallel reduction variable as a copy or pcopy
+! data clause.
+
+! { dg-additional-options "-fdump-tree-gimple" }
+
+program test
+  implicit none
+  integer :: v1, v2
+
+  !$acc parallel reduction(+:v1,v2)
+  !$acc end parallel
+
+  !$acc parallel reduction(+:v1,v2) copy(v1,v2)
+  !$acc end parallel
+
+  !$acc parallel reduction(+:v1,v2) pcopy(v1,v2)
+  !$acc end parallel
+
+  !$acc parallel reduction(+:v1,v2) present(v1,v2)
+  !$acc end parallel
+
+  !$acc parallel reduction(+:v1,v2) copyin(v1,v2) ! { dg-warning "incompatible data clause" }
+  !$acc end parallel
+
+  !$acc parallel reduction(+:v1,v2) pcopyin(v1,v2) ! { dg-warning "incompatible data clause" }
+  !$acc end parallel
+
+  !$acc parallel reduction(+:v1,v2) copyout(v1,v2) ! { dg-warning "incompatible data clause" }
+  !$acc end parallel
+
+  !$acc parallel reduction(+:v1,v2) pcopyout(v1,v2) ! { dg-warning "incompatible data clause" }
+  !$acc end parallel
+
+  !$acc parallel reduction(+:v1,v2) create(v1,v2) ! { dg-warning "incompatible data clause" }
+  !$acc end parallel
+
+  !$acc parallel reduction(+:v1,v2) pcreate(v1,v2) ! { dg-warning "incompatible data clause" }
+  !$acc end parallel
+end program test
+! { dg-final { scan-tree-dump-times "map.tofrom:v1" 8 "gimple" } }
+! { dg-final { scan-tree-dump-times "map.tofrom:v2" 8 "gimple" } }
+! { dg-final { scan-tree-dump-times "map.force_tofrom:v1" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "map.force_tofrom:v2" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "map.force_present:v1" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "map.force_present:v2" 1 "gimple" } }
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr70289.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr70289.c
new file mode 100644
index 0000000..6d52222
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr70289.c
@@ -0,0 +1,13 @@
+int
+main ()
+{
+  int i;
+  static int temp;
+
+#pragma acc parallel reduction(+:temp)
+  {
+    temp++;
+  }
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr70373.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr70373.c
new file mode 100644
index 0000000..af629c8
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr70373.c
@@ -0,0 +1,20 @@
+#define N 32
+
+int
+foo (unsigned int sum)
+{
+#pragma acc parallel reduction (+:sum)
+  {
+    sum;
+  }
+
+  return sum;
+}
+
+int
+main (void)
+{
+  unsigned int sum = 0;
+  foo (sum);
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/pr70289.f90 b/libgomp/testsuite/libgomp.oacc-fortran/pr70289.f90
new file mode 100644
index 0000000..63bde44
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/pr70289.f90
@@ -0,0 +1,20 @@
+program foo
+  implicit none
+  integer :: i
+  integer :: temp = 0
+  integer :: temp2 = 0
+
+  !$acc parallel
+  !$acc loop gang private(temp)
+  do i=1, 10000
+     temp = 0
+  enddo
+  !$acc end parallel
+
+  !$acc parallel reduction(+:temp2)
+  !$acc loop gang reduction(+:temp2)
+  do i=1, 10000
+     temp2 = 0
+  enddo
+  !$acc end parallel
+end program foo
