[patch] Fortran fix for PR70289

[patch] Fortran fix for PR70289

[Cesar Philippidis]

[-- Attachment #1: Type: text/plain, Size: 2119 bytes --]

The bug in PR70289 is an assertion failure triggered by a static
variable used inside an offloaded acc region which doesn't have a data
clause associated with it. Basically, that static variable ends up in a
different lto partition, which was not streamed to the offloaded
compiler. I'm not sure if we should try to replicate the static storage
in the offloaded regions, but it probably doesn't make sense in a
parallel environment anyway.

My solution to this problem was to teach the fortran front end to create
a data clause for each reduction variable it encounters. Furthermore,
I've decided to update the semantics of the acc parallel reduction
clause such that gfortran will emit an error when a reduction variable
is private or firstprivate (note that an acc loop reduction still works
with private and firstprivate reductions, just not acc parallel
reductions). The second change is to emit a warning when an incompatible
data clause is used with a reduction, and promote that data clause to a
present_or_copy. My rationale behind the promotion is, you cannot have a
copyin reduction variable because the original variable on the host will
not be updated. Similarly, you cannot have a copyout reduction variable
because the reduction operator is supposed to combine the results of the
reduction with the original reduction variable, but in copyout the
original variable is not initialized on the accelerator. But perhaps the
copyout rule is too strict?

Tom and I are still working with the OpenACC technical committee to get
some clarification on how the reduction value should behave with respect
to data movement. In the meantime, I wanted to see if this type of
solution would be appropriate for trunk. I was trying to get this to
work in the gimplifier so that one patch could resolve the problem for
all of the front ends, but that was happening too late. Especially for
reference types.

By the way, we will also benefit from this patch too
<https://gcc.gnu.org/ml/gcc-patches/2016-01/msg00482.html>. If not for
these reduction, but for global acc variables which haven't been
properly declared.

Cesar

[-- Attachment #2: pr70289-20160331.diff --]
[-- Type: text/x-patch, Size: 6707 bytes --]

2016-03-31  Cesar Philippidis  <cesar@codesourcery.com>

	gcc/fortran/
	* openmp.c (resolve_oacc_private_reductions): New function.
	(resolve_omp_clauses): Ensure that acc parallel reductions
	have a copy, pcopy or present clause associated with it,
	otherwise emit a warning or error as appropriate.

	gcc/testsuite/
	* gfortran.dg/goacc/reduction-promotions.f90: New test.
	* gfortran.dg/goacc/reduction.f95:

	libgomp/
	* testsuite/libgomp.oacc-fortran/pr70289.f90: New test.


diff --git a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c
index a6c39cd..e59997f 100644
--- a/gcc/fortran/openmp.c
+++ b/gcc/fortran/openmp.c
@@ -3146,6 +3146,46 @@ resolve_omp_udr_clause (gfc_omp_namelist *n, gfc_namespace *ns,
 /* OpenMP directive resolving routines.  */
 
 static void
+resolve_oacc_private_reductions (gfc_omp_clauses *omp_clauses, int list)
+{
+  gfc_omp_namelist *n;
+
+  /* Check for bogus private reductions.  */
+  for (n = omp_clauses->lists[list]; n; n = n->next)
+    n->sym->mark = 1;
+
+  for (n = omp_clauses->lists[OMP_LIST_REDUCTION]; n; n = n->next)
+    if (n->sym->mark)
+      {
+	gfc_omp_namelist *prev = NULL, *tmp = NULL;
+
+	/* Remove it from the list of private clauses.  */
+	tmp = omp_clauses->lists[list];
+	while (tmp)
+	  {
+	    if (tmp->sym == n->sym)
+	      {
+		gfc_error ("Reduction symbol %qs cannot be private "
+			   "at %L", tmp->sym->name, &tmp->where);
+
+		if (omp_clauses->lists[list] == tmp)
+		  omp_clauses->lists[list] = tmp->next;
+		else
+		  prev->next = tmp->next;
+		break;
+	      }
+	    else
+	      {
+		prev = tmp;
+		tmp = tmp->next;
+	      }
+	  }
+
+	n->sym->mark = 0;
+      }
+}
+
+static void
 resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses,
 		     gfc_namespace *ns, bool openacc = false)
 {
@@ -3320,6 +3360,50 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses,
 	    gfc_error ("Array %qs is not permitted in reduction at %L",
 		       n->sym->name, &n->where);
 	}
+
+      /* Parallel reductions have their own set of rules.  */
+      if (code->op == EXEC_OACC_PARALLEL)
+	{
+	  for (n = omp_clauses->lists[OMP_LIST_REDUCTION]; n; n = n->next)
+	    n->sym->mark = 0;
+
+	  resolve_oacc_private_reductions (omp_clauses, OMP_LIST_PRIVATE);
+	  resolve_oacc_private_reductions (omp_clauses, OMP_LIST_FIRSTPRIVATE);
+
+	  /* Check for data maps which aren't copy or present_or_copy.  */
+	  for (n = omp_clauses->lists[OMP_LIST_MAP]; n; n = n->next)
+	    {
+	      if (n->sym->mark == 0
+		  && !(n->u.map_op == OMP_MAP_TOFROM
+		       || n->u.map_op == OMP_MAP_FORCE_TOFROM
+		       || n->u.map_op == OMP_MAP_FORCE_PRESENT))
+		{
+		  gfc_warning (0, "incompatible data clause associated with "
+			       "symbol %qs; promoting this clause to "
+			       "present_or_copy at %L", n->sym->name,
+			       &n->where);
+		  n->u.map_op = OMP_MAP_TOFROM;
+		}
+	      n->sym->mark = 1;
+	    }
+
+	  /* Add a present_or_copy data clause for any reduction which
+	     doesnot already have one.  */
+	  for (n = omp_clauses->lists[OMP_LIST_REDUCTION]; n; n = n->next)
+	    {
+	      if (n->sym->mark == 0)
+		{
+		  gfc_omp_namelist *p = gfc_get_omp_namelist ();
+		  p->next = omp_clauses->lists[OMP_LIST_MAP];
+		  omp_clauses->lists[OMP_LIST_MAP] = p;
+		  p->sym = n->sym;
+		  p->expr = NULL;
+		  p->where = n->where;
+		  p->u.map_op = OMP_MAP_TOFROM;
+		}
+	      n->sym->mark = 1;
+	    }
+	}
     }
   
   for (n = omp_clauses->lists[OMP_LIST_TO]; n; n = n->next)
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..28d0951
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/reduction-promotions.f90
@@ -0,0 +1,43 @@
+! Ensure that each parallel reduction variable as a copy or pcopy
+! data clause.
+
+! { dg-additional-options "-fdump-tree-gimple" }
+
+program test
+  implicit none
+  integer :: var
+
+  !$acc parallel reduction(+:var)
+  !$acc end parallel
+  
+  !$acc parallel reduction(+:var) copy(var)
+  !$acc end parallel
+
+  !$acc parallel reduction(+:var) pcopy(var)
+  !$acc end parallel
+
+  !$acc parallel reduction(+:var) present(var)
+  !$acc end parallel
+  
+  !$acc parallel reduction(+:var) copyin(var) ! { dg-warning "incompatible data clause" }
+  !$acc end parallel
+
+  !$acc parallel reduction(+:var) pcopyin(var) ! { dg-warning "incompatible data clause" }
+  !$acc end parallel
+
+  !$acc parallel reduction(+:var) copyout(var) ! { dg-warning "incompatible data clause" }
+  !$acc end parallel
+
+  !$acc parallel reduction(+:var) pcopyout(var) ! { dg-warning "incompatible data clause" }
+  !$acc end parallel
+
+  !$acc parallel reduction(+:var) create(var) ! { dg-warning "incompatible data clause" }
+  !$acc end parallel
+
+  !$acc parallel reduction(+:var) pcreate(var) ! { dg-warning "incompatible data clause" }
+  !$acc end parallel
+end program test
+
+! { dg-final { scan-tree-dump-times "target oacc_parallel map.tofrom:var" 8 "gimple" } }
+! { dg-final { scan-tree-dump-times "target oacc_parallel map.force_tofrom:var" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "target oacc_parallel map.force_present:var" 1 "gimple" } }
diff --git a/gcc/testsuite/gfortran.dg/goacc/reduction.f95 b/gcc/testsuite/gfortran.dg/goacc/reduction.f95
index a13574b..9bef09a 100644
--- a/gcc/testsuite/gfortran.dg/goacc/reduction.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/reduction.f95
@@ -134,8 +134,11 @@ common /blk/ i1
 !$acc end parallel
 !$acc parallel reduction (iand:ta1)	! { dg-error "OMP DECLARE REDUCTION iand not found for type TYPE" }
 !$acc end parallel
-
-end subroutine
+!$acc parallel reduction (+:i1) private(i1) ! { dg-error "Reduction symbol .i1. cannot be private" }
+!$acc end parallel
+!$acc parallel reduction (+:i2) firstprivate(i2) ! { dg-error "Reduction symbol .i2. cannot be private" }
+!$acc end parallel
+end subroutine foo
 
 ! { dg-error "Array 'ia2' is not permitted in reduction" "" { target "*-*-*" } 27 }
 ! { dg-error "Array 'ra1' is not permitted in reduction" "" { target "*-*-*" } 29 }
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

Re: [patch] Fortran fix for PR70289

[Jakub Jelinek]

On Fri, Apr 01, 2016 at 07:49:16AM -0700, Cesar Philippidis wrote:
> The bug in PR70289 is an assertion failure triggered by a static
> variable used inside an offloaded acc region which doesn't have a data
> clause associated with it. Basically, that static variable ends up in a
> different lto partition, which was not streamed to the offloaded
> compiler. I'm not sure if we should try to replicate the static storage
> in the offloaded regions, but it probably doesn't make sense in a
> parallel environment anyway.

Is this really Fortran specific?  I'd expect the diagnostics to be in
gimplify.c and handle it for all 3 FEs...

	Jakub

Re: [patch] Fortran fix for PR70289

[Cesar Philippidis]

On 04/01/2016 07:56 AM, Jakub Jelinek wrote:
> On Fri, Apr 01, 2016 at 07:49:16AM -0700, Cesar Philippidis wrote:
>> The bug in PR70289 is an assertion failure triggered by a static
>> variable used inside an offloaded acc region which doesn't have a data
>> clause associated with it. Basically, that static variable ends up in a
>> different lto partition, which was not streamed to the offloaded
>> compiler. I'm not sure if we should try to replicate the static storage
>> in the offloaded regions, but it probably doesn't make sense in a
>> parallel environment anyway.
> 
> Is this really Fortran specific?  I'd expect the diagnostics to be in
> gimplify.c and handle it for all 3 FEs...

By the time the variable reaches the gimplifier, the reduction variable
may no longer match the ones inside the data clause. E.g. consider this
directive inside a fortran subroutine:

  !$acc parallel copyout(temp) reduction(+:temp)

The gimplifier would see something like:

  map(force_from:*temp.2 [len: 4]) map(alloc:temp [pointer assign, bias:
0]) reduction(+:temp)

At this point, unless I'm mistaken, it would be difficult to tell if
temp.2 is a pointer to the same temp in the reduction. Maybe I'm missing
something?

Cesar

Re: [patch] Fortran fix for PR70289

[Jakub Jelinek]

On Fri, Apr 01, 2016 at 08:07:24AM -0700, Cesar Philippidis wrote:
> On 04/01/2016 07:56 AM, Jakub Jelinek wrote:
> > On Fri, Apr 01, 2016 at 07:49:16AM -0700, Cesar Philippidis wrote:
> >> The bug in PR70289 is an assertion failure triggered by a static
> >> variable used inside an offloaded acc region which doesn't have a data
> >> clause associated with it. Basically, that static variable ends up in a
> >> different lto partition, which was not streamed to the offloaded
> >> compiler. I'm not sure if we should try to replicate the static storage
> >> in the offloaded regions, but it probably doesn't make sense in a
> >> parallel environment anyway.
> > 
> > Is this really Fortran specific?  I'd expect the diagnostics to be in
> > gimplify.c and handle it for all 3 FEs...
> 
> By the time the variable reaches the gimplifier, the reduction variable
> may no longer match the ones inside the data clause. E.g. consider this
> directive inside a fortran subroutine:
> 
>   !$acc parallel copyout(temp) reduction(+:temp)
> 
> The gimplifier would see something like:
> 
>   map(force_from:*temp.2 [len: 4]) map(alloc:temp [pointer assign, bias:
> 0]) reduction(+:temp)
> 
> At this point, unless I'm mistaken, it would be difficult to tell if
> temp.2 is a pointer to the same temp in the reduction. Maybe I'm missing
> something?

All the info is still there, and this wouldn't be the only case where
we rely on exact clause ordering.  I think that is still much better than
doing it in all the FEs.

	Jakub

Re: [patch] Fortran fix for PR70289

[Cesar Philippidis]

[-- Attachment #1: Type: text/plain, Size: 1690 bytes --]

On 04/01/2016 08:17 AM, Jakub Jelinek wrote:
> On Fri, Apr 01, 2016 at 08:07:24AM -0700, Cesar Philippidis wrote:
>> On 04/01/2016 07:56 AM, Jakub Jelinek wrote:
>>> On Fri, Apr 01, 2016 at 07:49:16AM -0700, Cesar Philippidis wrote:
>>>> The bug in PR70289 is an assertion failure triggered by a static
>>>> variable used inside an offloaded acc region which doesn't have a data
>>>> clause associated with it. Basically, that static variable ends up in a
>>>> different lto partition, which was not streamed to the offloaded
>>>> compiler. I'm not sure if we should try to replicate the static storage
>>>> in the offloaded regions, but it probably doesn't make sense in a
>>>> parallel environment anyway.
>>>
>>> Is this really Fortran specific?  I'd expect the diagnostics to be in
>>> gimplify.c and handle it for all 3 FEs...
>>
>> By the time the variable reaches the gimplifier, the reduction variable
>> may no longer match the ones inside the data clause. E.g. consider this
>> directive inside a fortran subroutine:
>>
>>   !$acc parallel copyout(temp) reduction(+:temp)
>>
>> The gimplifier would see something like:
>>
>>   map(force_from:*temp.2 [len: 4]) map(alloc:temp [pointer assign, bias:
>> 0]) reduction(+:temp)
>>
>> At this point, unless I'm mistaken, it would be difficult to tell if
>> temp.2 is a pointer to the same temp in the reduction. Maybe I'm missing
>> something?
> 
> All the info is still there, and this wouldn't be the only case where
> we rely on exact clause ordering.  I think that is still much better than

The gimplify approach didn't turn out to be that bad after all. Is this
patch ok for trunk? It fixes the problem for all fo the FEs.

Cesar

[-- Attachment #2: pr70289-20160404.diff --]
[-- Type: text/x-patch, Size: 11651 bytes --]

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