diff mbox series

[OpenACC,2.7] struct/array reductions for Fortran

Message ID 9209bd62-7ca1-4480-8497-d402b2889a72@baylibre.com
State New
Headers show
Series [OpenACC,2.7] struct/array reductions for Fortran | expand

Commit Message

Chung-Lin Tang Feb. 8, 2024, 2:47 p.m. UTC
Hi Tobias, Thomas,
this patch adds support for Fortran to use arrays and struct(record) types in OpenACC reductions.

There is still some shortcomings in the current state, mainly that only explicit-shaped arrays can be used (like its C counterpart). Anything else is currently a bit more complicated in the middle-end, since the existing reduction code creates an "init-op" (literal of initial values) which can't be done when say TYPE_MAX_VALUE (TYPE_DOMAIN (array_type)) is not a tree constant. I think we'll be on the hook to solve this later, but I think the current state is okay to submit.

Tested without regressions on mainline (on top of first struct/array reduction patch[1])

Thanks,
Chung-Lin

[1] https://gcc.gnu.org/pipermail/gcc-patches/2024-January/641669.html

2024-02-08  Chung-Lin Tang  <cltang@baylibre.com>

gcc/fortran/ChangeLog:
	* openmp.cc (oacc_reduction_defined_type_p): New function.
	(resolve_omp_clauses): Adjust OpenACC array reduction error case. Use
	oacc_reduction_defined_type_p for OpenACC.
	* trans-openmp.cc (gfc_trans_omp_array_reduction_or_udr):
	Add 'bool openacc' parameter, adjust part of function to be !openacc
	only.
	(gfc_trans_omp_reduction_list): Add 'bool openacc' parameter, pass to
	calls to gfc_trans_omp_array_reduction_or_udr.
	(gfc_trans_omp_clauses): Add 'openacc' argument to calls to
	gfc_trans_omp_reduction_list.
	(gfc_trans_omp_do): Pass 'op == EXEC_OACC_LOOP' as 'bool openacc'
	parameter in call to gfc_trans_omp_clauses.

gcc/ChangeLog:
	* omp-low.cc (omp_reduction_init_op): Add checking if reduced array
	has constant bounds.
	(lower_oacc_reductions): Add handling of error_mark_node.

gcc/testsuite/ChangeLog:
	* gfortran.dg/goacc/array-reduction.f90: Adjust testcase.
	* gfortran.dg/goacc/reduction.f95: Likewise.

libgomp/ChangeLog:
	* libgomp/testsuite/libgomp.oacc-fortran/reduction-9.f90: New testcase.
	* libgomp/testsuite/libgomp.oacc-fortran/reduction-10.f90: Likewise.
	* libgomp/testsuite/libgomp.oacc-fortran/reduction-11.f90: Likewise.
	* libgomp/testsuite/libgomp.oacc-fortran/reduction-12.f90: Likewise.
	* libgomp/testsuite/libgomp.oacc-fortran/reduction-13.f90: Likewise.

Comments

Tobias Burnus March 13, 2024, 6:59 p.m. UTC | #1
Hi Chung-Lin, hi Thomas, hello world,

some thoughts glancing at the patch.

Chung-Lin Tang wrote:
> There is still some shortcomings in the current state, mainly that only explicit-shaped arrays can be used (like its C counterpart). Anything else is currently a bit more complicated in the middle-end, since the existing reduction code creates an "init-op" (literal of initial values) which can't be done when say TYPE_MAX_VALUE (TYPE_DOMAIN (array_type)) is not a tree constant. I think we'll be on the hook to solve this later, but I think the current state is okay to submit.

I think having some initial support is fine, but it needs an 
understandable and somewhat complete error diagnostic and testcases. 
More to this below.

> +      if (!TREE_CONSTANT (min_tree) || !TREE_CONSTANT (max_tree))
> +	{
> +	  error_at (loc, "array in reduction must be of constant size");
> +	  return error_mark_node;
> +	}
Shouldn't this use a sorry_at instead?

> +	  /* OpenACC current only supports array reductions on explicit-shape
> +	     arrays.  */
> +	  if ((n->sym->as && n->sym->as->type != AS_EXPLICIT)
> +	      || n->sym->attr.codimension)
>   	    gfc_error ("Array %qs is not permitted in reduction at %L",
>   		       n->sym->name, &n->where);
[Coarray excursion. I am in favor of allowing it for the reasons above, 
but it could be also rejected but I would prefer to have a proper error 
message in that case.]

While coarrays are unspecified, I do not see a reason why a corray 
shouldn't be permitted here – as long as it is not coindexed. At the 
end, it is just a normal array with some additional properties, which 
make it possible to remotely access it.

Note: For coarray scalars, we have 'sym->as', thus the check should be 
'(n->sym->as && n->sym->as->rank)' to permit scalar coarrays.

* * *

Coarray excursion: A coarray variables exists in multiple processes 
("images", e.g. MPI processes). If 'caf' and 'caf2' are coarrays, then 
'caf = 5' and 'i = caf2' refer to the local variable.

On the other hand, 'caf[n] = 5' or 'i = caf[3,m]' refers to the 'caf' 
variable on image 'n' or [3,m]', respectively, which implies in general 
some function call to read or set the remote data, unless the memory is 
directly accessible (→ e.g. some offset calculation) and the compiler 
already knows how to handle this.

While a coarrary might be allocated in some special memory, as long as 
one uses the local version (i.e. not coindexed / without the image index 
in brackets).

Assume for the example above, e.g., integer :: caf[*], caf2[3:6, 7:*].

* * *

Thus, in terms of OpenACC or OpenMP, there is no reason to fret a 
coarray as long as it is not coindexed and as long as OpenMP/OpenACC 
does not interfere with the memory allocation – either directly ('!$omp 
allocators') or indirectly by placing it into special memory (pinned, 
pseudo-unified-shared memory → OG13's -foffload-memory=pinned/unified).

In the meanwhile, OpenMP actually explicitly allows coarrays with few 
exceptions while OpenACC talks about unspecified behavior.

* * *

Back to generic comments:

If I look at the existing code, I see at gfc_match_omp_clause_reduction:

>  if (gfc_match_omp_variable_list (" :", &c->lists[list_idx], false, NULL,
>                                   &head, openacc, allow_derived) != 
> MATCH_YES)

If 'openacc' is true, array sections are permitted - but the code added 
(see quote above) does not handle n->expr at all and only n->sym.

I think there needs to be at least a "gfc_error ("Sorry, subarrays/array 
sections not yet handled" [subarray is the OpenACC wording, 'array 
section' is the Fortran one, which might be clearer.

But you could consider to handle at least array elements, i.e. 
n->expr->rank == 0.

Additionally, I think the current error message is completely unhelpful 
given that some arrays are supported but most are not.

I think there should be also some testcases for the not-yet-supported 
case. I think the following will trigger the omp-low.cc 'sorry_at' (or 
currently 'error' - but I think it should be a sorry):

subroutine foo(n)

integer :: n, A(n)

... reduction(+:A)

And most others will trigger in openmp.cc; for those, you should have an 
allocatable/pointer and assumed-shape arrays for the diagnostic testcase 
as well.

* * *

I have not really experimented with the code, but does it handle 
multi-dimensional constant arrays like 'integer :: a(3:6,10,-1:1)' ? — I 
bet it does, at least after handling my example [2] for the C patch [1].

Thanks,

Tobias

[1] https://gcc.gnu.org/pipermail/gcc-patches/2024-January/641669.html

[2] https://gcc.gnu.org/pipermail/gcc-patches/2024-March/647704.html
Thomas Schwinge March 18, 2024, 4:39 p.m. UTC | #2
Hi Chung-Lin!

Thanks for your work here, which I'm beginning to look into (prerequisite
"[PATCH, OpenACC 2.7] Implement reductions for arrays and structs",
first, of course); it'll take me some time.


In non-offloading testing, I noticed for x86_64-pc-linux-gnu '-m32':

    +PASS: libgomp.oacc-fortran/reduction-13.f90 -DACC_DEVICE_TYPE_host=1 -DACC_MEM_SHARED=1 -foffload=disable  -O0  (test for excess errors)
    +PASS: libgomp.oacc-fortran/reduction-13.f90 -DACC_DEVICE_TYPE_host=1 -DACC_MEM_SHARED=1 -foffload=disable  -O0  execution test
    +PASS: libgomp.oacc-fortran/reduction-13.f90 -DACC_DEVICE_TYPE_host=1 -DACC_MEM_SHARED=1 -foffload=disable  -O1  (test for excess errors)
    +FAIL: libgomp.oacc-fortran/reduction-13.f90 -DACC_DEVICE_TYPE_host=1 -DACC_MEM_SHARED=1 -foffload=disable  -O1  execution test
    +PASS: libgomp.oacc-fortran/reduction-13.f90 -DACC_DEVICE_TYPE_host=1 -DACC_MEM_SHARED=1 -foffload=disable  -O2  (test for excess errors)
    +FAIL: libgomp.oacc-fortran/reduction-13.f90 -DACC_DEVICE_TYPE_host=1 -DACC_MEM_SHARED=1 -foffload=disable  -O2  execution test
    +PASS: libgomp.oacc-fortran/reduction-13.f90 -DACC_DEVICE_TYPE_host=1 -DACC_MEM_SHARED=1 -foffload=disable  -O3 -fomit-frame-pointer -funroll-loops -fpeel-loops -ftracer -finline-functions  (test for excess errors)
    +FAIL: libgomp.oacc-fortran/reduction-13.f90 -DACC_DEVICE_TYPE_host=1 -DACC_MEM_SHARED=1 -foffload=disable  -O3 -fomit-frame-pointer -funroll-loops -fpeel-loops -ftracer -finline-functions  execution test
    +PASS: libgomp.oacc-fortran/reduction-13.f90 -DACC_DEVICE_TYPE_host=1 -DACC_MEM_SHARED=1 -foffload=disable  -O3 -g  (test for excess errors)
    +FAIL: libgomp.oacc-fortran/reduction-13.f90 -DACC_DEVICE_TYPE_host=1 -DACC_MEM_SHARED=1 -foffload=disable  -O3 -g  execution test
    +PASS: libgomp.oacc-fortran/reduction-13.f90 -DACC_DEVICE_TYPE_host=1 -DACC_MEM_SHARED=1 -foffload=disable  -Os  (test for excess errors)
    +FAIL: libgomp.oacc-fortran/reduction-13.f90 -DACC_DEVICE_TYPE_host=1 -DACC_MEM_SHARED=1 -foffload=disable  -Os  execution test

With optimizations enabled, it runs into 'STOP 4'.

Per '-Wextra':

    [...]/libgomp.oacc-fortran/reduction-13.f90:40:6: Warning: Inequality comparison for REAL(4) at (1) [-Wcompare-reals]
    [...]/libgomp.oacc-fortran/reduction-13.f90:63:6: Warning: Inequality comparison for REAL(4) at (1) [-Wcompare-reals]
    [...]/libgomp.oacc-fortran/reduction-13.f90:64:6: Warning: Inequality comparison for REAL(8) at (1) [-Wcompare-reals]

Do we need to allow for some epsilon (generally in such test cases), or
is there another problem?

For reference:

On 2024-02-08T22:47:13+0800, Chung-Lin Tang <cltang@baylibre.com> wrote:
> --- /dev/null
> +++ b/libgomp/testsuite/libgomp.oacc-fortran/reduction-13.f90
> @@ -0,0 +1,66 @@
> +! { dg-do run }
> +
> +! record type reductions
> +
> +program reduction_13
> +  implicit none
> +
> +  type t1
> +     integer :: i
> +     real :: r
> +  end type t1
> +
> +  type t2
> +     real :: r
> +     integer :: i
> +     double precision :: d
> +  end type t2
> +
> +  integer, parameter :: n = 10, ng = 8, nw = 4, vl = 32
> +  integer :: i
> +  type(t1) :: v1, a1
> +  type (t2) :: v2, a2
> +
> +  v1%i = 0
> +  v1%r = 0
> +  !$acc parallel num_gangs(ng) num_workers(nw) vector_length(vl) copy(v1)
> +  !$acc loop reduction (+:v1)
> +  do i = 1, n
> +     v1%i = v1%i + 1
> +     v1%r = v1%r + 2
> +  end do
> +  !$acc end parallel
> +  a1%i = 0
> +  a1%r = 0
> +  do i = 1, n
> +     a1%i = a1%i + 1
> +     a1%r = a1%r + 2
> +  end do
> +  if (v1%i .ne. a1%i) STOP 1
> +  if (v1%r .ne. a1%r) STOP 2
> +
> +  v2%i = 1
> +  v2%r = 1
> +  v2%d = 1
> +  !$acc parallel num_gangs(ng) num_workers(nw) vector_length(vl) copy(v2)
> +  !$acc loop reduction (*:v2)
> +  do i = 1, n
> +     v2%i = v2%i * 2
> +     v2%r = v2%r * 1.1
> +     v2%d = v2%d * 1.3
> +  end do
> +  !$acc end parallel
> +  a2%i = 1
> +  a2%r = 1
> +  a2%d = 1
> +  do i = 1, n
> +     a2%i = a2%i * 2
> +     a2%r = a2%r * 1.1
> +     a2%d = a2%d * 1.3
> +  end do
> +
> +  if (v2%i .ne. a2%i) STOP 3
> +  if (v2%r .ne. a2%r) STOP 4
> +  if (v2%d .ne. a2%d) STOP 5
> +
> +end program reduction_13


Grüße
 Thomas
diff mbox series

Patch

diff --git a/gcc/fortran/openmp.cc b/gcc/fortran/openmp.cc
index 0af80d54fad..4bba9e666d6 100644
--- a/gcc/fortran/openmp.cc
+++ b/gcc/fortran/openmp.cc
@@ -7047,6 +7047,72 @@  oacc_is_loop (gfc_code *code)
 	 || code->op == EXEC_OACC_LOOP;
 }
 
+static bool
+oacc_reduction_defined_type_p (enum gfc_omp_reduction_op rop, gfc_typespec *ts)
+{
+  if (rop == OMP_REDUCTION_USER || rop == OMP_REDUCTION_NONE)
+    return false;
+
+  if (ts->type == BT_INTEGER)
+    switch (rop)
+      {
+      case OMP_REDUCTION_AND:
+      case OMP_REDUCTION_OR:
+      case OMP_REDUCTION_EQV:
+      case OMP_REDUCTION_NEQV:
+	return false;
+      default:
+	return true;
+      }
+
+  if (ts->type == BT_LOGICAL)
+    switch (rop)
+      {
+      case OMP_REDUCTION_AND:
+      case OMP_REDUCTION_OR:
+      case OMP_REDUCTION_EQV:
+      case OMP_REDUCTION_NEQV:
+	return true;
+      default:
+	return false;
+      }
+
+  if (ts->type == BT_REAL || ts->type == BT_COMPLEX)
+    switch (rop)
+      {
+      case OMP_REDUCTION_PLUS:
+      case OMP_REDUCTION_TIMES:
+      case OMP_REDUCTION_MINUS:
+	return true;
+
+      case OMP_REDUCTION_AND:
+      case OMP_REDUCTION_OR:
+      case OMP_REDUCTION_EQV:
+      case OMP_REDUCTION_NEQV:
+	return false;
+
+      case OMP_REDUCTION_MAX:
+      case OMP_REDUCTION_MIN:
+	return ts->type != BT_COMPLEX;
+      case OMP_REDUCTION_IAND:
+      case OMP_REDUCTION_IOR:
+      case OMP_REDUCTION_IEOR:
+	return false;
+      default:
+	gcc_unreachable ();
+      }
+
+  if (ts->type == BT_DERIVED)
+    {
+      for (gfc_component *p = ts->u.derived->components; p; p = p->next)
+	if (!oacc_reduction_defined_type_p (rop, &p->ts))
+	  return false;
+      return true;
+    }
+
+  return false;
+}
+
 static void
 resolve_scalar_int_expr (gfc_expr *expr, const char *clause)
 {
@@ -8137,13 +8203,15 @@  resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses,
 	  else
 	    n->sym->mark = 1;
 
-	  /* OpenACC does not support reductions on arrays.  */
-	  if (n->sym->as)
+	  /* OpenACC current only supports array reductions on explicit-shape
+	     arrays.  */
+	  if ((n->sym->as && n->sym->as->type != AS_EXPLICIT)
+	      || n->sym->attr.codimension)
 	    gfc_error ("Array %qs is not permitted in reduction at %L",
 		       n->sym->name, &n->where);
 	}
     }
-  
+
   for (n = omp_clauses->lists[OMP_LIST_TO]; n; n = n->next)
     n->sym->mark = 0;
   for (n = omp_clauses->lists[OMP_LIST_FROM]; n; n = n->next)
@@ -8797,39 +8865,46 @@  resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses,
 		  case OMP_LIST_IN_REDUCTION:
 		  case OMP_LIST_TASK_REDUCTION:
 		  case OMP_LIST_REDUCTION_INSCAN:
-		    switch (n->u.reduction_op)
+		    if (openacc)
 		      {
-		      case OMP_REDUCTION_PLUS:
-		      case OMP_REDUCTION_TIMES:
-		      case OMP_REDUCTION_MINUS:
-			if (!gfc_numeric_ts (&n->sym->ts))
+			if (!oacc_reduction_defined_type_p (n->u.reduction_op,
+							    &n->sym->ts))
 			  bad = true;
-			break;
-		      case OMP_REDUCTION_AND:
-		      case OMP_REDUCTION_OR:
-		      case OMP_REDUCTION_EQV:
-		      case OMP_REDUCTION_NEQV:
-			if (n->sym->ts.type != BT_LOGICAL)
-			  bad = true;
-			break;
-		      case OMP_REDUCTION_MAX:
-		      case OMP_REDUCTION_MIN:
-			if (n->sym->ts.type != BT_INTEGER
-			    && n->sym->ts.type != BT_REAL)
-			  bad = true;
-			break;
-		      case OMP_REDUCTION_IAND:
-		      case OMP_REDUCTION_IOR:
-		      case OMP_REDUCTION_IEOR:
-			if (n->sym->ts.type != BT_INTEGER)
-			  bad = true;
-			break;
-		      case OMP_REDUCTION_USER:
-			bad = true;
-			break;
-		      default:
-			break;
 		      }
+		    else
+		      switch (n->u.reduction_op)
+			{
+			case OMP_REDUCTION_PLUS:
+			case OMP_REDUCTION_TIMES:
+			case OMP_REDUCTION_MINUS:
+			  if (!gfc_numeric_ts (&n->sym->ts))
+			    bad = true;
+			  break;
+			case OMP_REDUCTION_AND:
+			case OMP_REDUCTION_OR:
+			case OMP_REDUCTION_EQV:
+			case OMP_REDUCTION_NEQV:
+			  if (n->sym->ts.type != BT_LOGICAL)
+			    bad = true;
+			  break;
+			case OMP_REDUCTION_MAX:
+			case OMP_REDUCTION_MIN:
+			  if (n->sym->ts.type != BT_INTEGER
+			      && n->sym->ts.type != BT_REAL)
+			    bad = true;
+			  break;
+			case OMP_REDUCTION_IAND:
+			case OMP_REDUCTION_IOR:
+			case OMP_REDUCTION_IEOR:
+			  if (n->sym->ts.type != BT_INTEGER)
+			    bad = true;
+			  break;
+			case OMP_REDUCTION_USER:
+			  bad = true;
+			  break;
+			default:
+			  break;
+			}
 		    if (!bad)
 		      n->u2.udr = NULL;
 		    else
diff --git a/gcc/fortran/trans-openmp.cc b/gcc/fortran/trans-openmp.cc
index 9599521b97c..29ad880a30c 100644
--- a/gcc/fortran/trans-openmp.cc
+++ b/gcc/fortran/trans-openmp.cc
@@ -1996,7 +1996,8 @@  omp_udr_find_orig (gfc_expr **e, int *walk_subtrees ATTRIBUTE_UNUSED,
 }
 
 static void
-gfc_trans_omp_array_reduction_or_udr (tree c, gfc_omp_namelist *n, locus where)
+gfc_trans_omp_array_reduction_or_udr (tree c, gfc_omp_namelist *n, locus where,
+				      bool openacc)
 {
   gfc_symbol *sym = n->sym;
   gfc_symtree *root1 = NULL, *root2 = NULL, *root3 = NULL, *root4 = NULL;
@@ -2251,21 +2252,24 @@  gfc_trans_omp_array_reduction_or_udr (tree c, gfc_omp_namelist *n, locus where)
     poplevel (0, 0);
   OMP_CLAUSE_REDUCTION_INIT (c) = stmt;
 
-  /* Create the merge statement list.  */
-  pushlevel ();
-  if (e4)
-    stmt = gfc_trans_assignment (e3, e4, false, true);
-  else
-    stmt = gfc_trans_call (n->u2.udr->combiner, false,
-			   NULL_TREE, NULL_TREE, false);
-  if (TREE_CODE (stmt) != BIND_EXPR)
-    stmt = build3_v (BIND_EXPR, NULL, stmt, poplevel (1, 0));
-  else
-    poplevel (0, 0);
-  OMP_CLAUSE_REDUCTION_MERGE (c) = stmt;
+  if (!openacc)
+    {
+      /* Create the merge statement list.  */
+      pushlevel ();
+      if (e4)
+	stmt = gfc_trans_assignment (e3, e4, false, true);
+      else
+	stmt = gfc_trans_call (n->u2.udr->combiner, false,
+			       NULL_TREE, NULL_TREE, false);
+      if (TREE_CODE (stmt) != BIND_EXPR)
+	stmt = build3_v (BIND_EXPR, NULL, stmt, poplevel (1, 0));
+      else
+	poplevel (0, 0);
+      OMP_CLAUSE_REDUCTION_MERGE (c) = stmt;
 
-  /* And stick the placeholder VAR_DECL into the clause as well.  */
-  OMP_CLAUSE_REDUCTION_PLACEHOLDER (c) = outer_decl;
+      /* And stick the placeholder VAR_DECL into the clause as well.  */
+      OMP_CLAUSE_REDUCTION_PLACEHOLDER (c) = outer_decl;
+    }
 
   gfc_current_locus = old_loc;
 
@@ -2296,7 +2300,7 @@  gfc_trans_omp_array_reduction_or_udr (tree c, gfc_omp_namelist *n, locus where)
 
 static tree
 gfc_trans_omp_reduction_list (int kind, gfc_omp_namelist *namelist, tree list,
-			      locus where, bool mark_addressable)
+			      locus where, bool mark_addressable, bool openacc)
 {
   omp_clause_code clause = OMP_CLAUSE_REDUCTION;
   switch (kind)
@@ -2376,7 +2380,8 @@  gfc_trans_omp_reduction_list (int kind, gfc_omp_namelist *namelist, tree list,
 	    if (namelist->sym->attr.dimension
 		|| namelist->u.reduction_op == OMP_REDUCTION_USER
 		|| namelist->sym->attr.allocatable)
-	      gfc_trans_omp_array_reduction_or_udr (node, namelist, where);
+	      gfc_trans_omp_array_reduction_or_udr (node, namelist, where,
+						    openacc);
 	    list = gfc_trans_add_clause (node, list);
 	  }
       }
@@ -2715,7 +2720,8 @@  gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
 	  /* An OpenACC async clause indicates the need to set reduction
 	     arguments addressable, to allow asynchronous copy-out.  */
 	  omp_clauses = gfc_trans_omp_reduction_list (list, n, omp_clauses,
-						      where, clauses->async);
+						      where, clauses->async,
+						      openacc);
 	  break;
 	case OMP_LIST_PRIVATE:
 	  clause_code = OMP_CLAUSE_PRIVATE;
@@ -5757,7 +5763,8 @@  gfc_trans_omp_do (gfc_code *code, gfc_exec_op op, stmtblock_t *pblock,
      on the simd construct and DO's clauses are translated elsewhere.  */
   do_clauses->sched_simd = false;
 
-  omp_clauses = gfc_trans_omp_clauses (pblock, do_clauses, code->loc);
+  omp_clauses = gfc_trans_omp_clauses (pblock, do_clauses, code->loc, false,
+				       op == EXEC_OACC_LOOP);
 
   for (i = 0; i < collapse; i++)
     {
diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc
index f3a056df8f2..4bbf30627c3 100644
--- a/gcc/omp-low.cc
+++ b/gcc/omp-low.cc
@@ -4426,9 +4426,16 @@  omp_reduction_init_op (location_t loc, enum tree_code op, tree type)
 {
   if (TREE_CODE (type) == ARRAY_TYPE)
     {
+      tree min_tree = TYPE_MIN_VALUE (TYPE_DOMAIN (type));
+      tree max_tree = TYPE_MAX_VALUE (TYPE_DOMAIN (type));
+      if (!TREE_CONSTANT (min_tree) || !TREE_CONSTANT (max_tree))
+	{
+	  error_at (loc, "array in reduction must be of constant size");
+	  return error_mark_node;
+	}
       vec<constructor_elt, va_gc> *v = NULL;
-      HOST_WIDE_INT min = tree_to_shwi (TYPE_MIN_VALUE (TYPE_DOMAIN (type)));
-      HOST_WIDE_INT max = tree_to_shwi (TYPE_MAX_VALUE (TYPE_DOMAIN (type)));
+      HOST_WIDE_INT min = tree_to_shwi (min_tree);
+      HOST_WIDE_INT max = tree_to_shwi (max_tree);
       tree t = omp_reduction_init_op (loc, op, TREE_TYPE (type));
       for (HOST_WIDE_INT i = min; i <= max; i++)
 	CONSTRUCTOR_APPEND_ELT (v, size_int (i), t);
@@ -7559,6 +7566,9 @@  lower_oacc_reductions (location_t loc, tree clauses, tree level, bool inner,
 	  has_outer_reduction:;
 	  }
 
+	if (incoming == error_mark_node)
+	  continue;
+
 	if (!ref_to_res)
 	  ref_to_res = integer_zero_node;
 
diff --git a/gcc/testsuite/gfortran.dg/goacc/array-reduction.f90 b/gcc/testsuite/gfortran.dg/goacc/array-reduction.f90
index d71c400a5bf..f9a3b43e7f3 100644
--- a/gcc/testsuite/gfortran.dg/goacc/array-reduction.f90
+++ b/gcc/testsuite/gfortran.dg/goacc/array-reduction.f90
@@ -1,74 +1,80 @@ 
+! { dg-do compile }
+! { dg-additional-options "-fdump-tree-gimple" }
+
 program test
   implicit none
   integer a(10), i
 
   a(:) = 0
-  
+
   ! Array reductions.
-  
-  !$acc parallel reduction (+:a) ! { dg-error "Array 'a' is not permitted in reduction" }
+
+  !$acc parallel reduction (+:a)
   do i = 1, 10
      a = a + 1
   end do
   !$acc end parallel
 
   !$acc parallel
-  !$acc loop reduction (+:a) ! { dg-error "Array 'a' is not permitted in reduction" }
+  !$acc loop reduction (+:a)
   do i = 1, 10
      a = a + 1
   end do
   !$acc end parallel
 
   !$acc kernels
-  !$acc loop reduction (+:a) ! { dg-error "Array 'a' is not permitted in reduction" }
+  !$acc loop reduction (+:a)
   do i = 1, 10
      a = a + 1
   end do
   !$acc end kernels
 
   ! Subarray reductions.
-  
-  !$acc parallel reduction (+:a(1:5)) ! { dg-error "Array 'a' is not permitted in reduction" }
+
+  !$acc parallel reduction (+:a(1:5))
   do i = 1, 10
      a = a + 1
   end do
   !$acc end parallel
 
   !$acc parallel
-  !$acc loop reduction (+:a(1:5)) ! { dg-error "Array 'a' is not permitted in reduction" }
+  !$acc loop reduction (+:a(1:5))
   do i = 1, 10
      a = a + 1
   end do
   !$acc end parallel
 
   !$acc kernels
-  !$acc loop reduction (+:a(1:5)) ! { dg-error "Array 'a' is not permitted in reduction" }
+  !$acc loop reduction (+:a(1:5))
   do i = 1, 10
      a = a + 1
   end do
   !$acc end kernels
 
   ! Reductions on array elements.
-  
-  !$acc parallel reduction (+:a(1)) ! { dg-error "Array 'a' is not permitted in reduction" }
+
+  !$acc parallel reduction (+:a(1))
   do i = 1, 10
      a(1) = a(1) + 1
   end do
   !$acc end parallel
 
   !$acc parallel
-  !$acc loop reduction (+:a(1)) ! { dg-error "Array 'a' is not permitted in reduction" }
+  !$acc loop reduction (+:a(1))
   do i = 1, 10
      a(1) = a(1) + 1
   end do
   !$acc end parallel
 
   !$acc kernels
-  !$acc loop reduction (+:a(1)) ! { dg-error "Array 'a' is not permitted in reduction" }
+  !$acc loop reduction (+:a(1))
   do i = 1, 10
      a(1) = a(1) + 1
   end do
   !$acc end kernels
-  
+
   print *, a
 end program test
+
+! { dg-final { scan-tree-dump-times "(?n)#pragma acc loop private\\(i\\) reduction\\(\\+:a\\)" 6 "gimple" } }
+! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_parallel reduction\\(\\+:a\\) map\\(tofrom:a \\\[len: \[0-9\]+\\\]\\)" 3 "gimple" } }
diff --git a/gcc/testsuite/gfortran.dg/goacc/reduction.f95 b/gcc/testsuite/gfortran.dg/goacc/reduction.f95
index a13574b150c..c425f00d87f 100644
--- a/gcc/testsuite/gfortran.dg/goacc/reduction.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/reduction.f95
@@ -72,9 +72,9 @@  common /blk/ i1
 !$acc end parallel
 !$acc parallel reduction (-:a1)		! { dg-error "OMP DECLARE REDUCTION - not found for type CHARACTER" }
 !$acc end parallel
-!$acc parallel reduction (+:t1)		! { dg-error "OMP DECLARE REDUCTION \\+ not found for type TYPE" }
+!$acc parallel reduction (+:t1)
 !$acc end parallel
-!$acc parallel reduction (*:ta1)	! { dg-error "OMP DECLARE REDUCTION \\* not found for type TYPE" }
+!$acc parallel reduction (*:ta1)
 !$acc end parallel
 !$acc parallel reduction (.and.:i3)	! { dg-error "OMP DECLARE REDUCTION \\.and\\. not found for type INTEGER" }
 !$acc end parallel
@@ -108,9 +108,9 @@  common /blk/ i1
 !$acc end parallel
 !$acc parallel reduction (max:a1)	! { dg-error "OMP DECLARE REDUCTION max not found for type CHARACTER" }
 !$acc end parallel
-!$acc parallel reduction (min:t1)	! { dg-error "OMP DECLARE REDUCTION min not found for type TYPE" }
+!$acc parallel reduction (min:t1)
 !$acc end parallel
-!$acc parallel reduction (max:ta1)	! { dg-error "OMP DECLARE REDUCTION max not found for type TYPE" }
+!$acc parallel reduction (max:ta1)
 !$acc end parallel
 !$acc parallel reduction (iand:r1)	! { dg-error "OMP DECLARE REDUCTION iand not found for type REAL" }
 !$acc end parallel
@@ -130,32 +130,12 @@  common /blk/ i1
 !$acc end parallel
 !$acc parallel reduction (ior:a1)	! { dg-error "OMP DECLARE REDUCTION ior not found for type CHARACTER" }
 !$acc end parallel
-!$acc parallel reduction (ieor:t1)	! { dg-error "OMP DECLARE REDUCTION ieor not found for type TYPE" }
+!$acc parallel reduction (ieor:t1)
 !$acc end parallel
-!$acc parallel reduction (iand:ta1)	! { dg-error "OMP DECLARE REDUCTION iand not found for type TYPE" }
+!$acc parallel reduction (iand:ta1)
 !$acc end parallel
 
 end subroutine
 
-! { dg-error "Array 'ia2' is not permitted in reduction" "" { target "*-*-*" } 27 }
-! { dg-error "Array 'ra1' is not permitted in reduction" "" { target "*-*-*" } 29 }
-! { dg-error "Array 'ca1' is not permitted in reduction" "" { target "*-*-*" } 31 }
-! { dg-error "Array 'da1' is not permitted in reduction" "" { target "*-*-*" } 33 }
-! { dg-error "Array 'la1' is not permitted in reduction" "" { target "*-*-*" } 35 }
 ! { dg-error "Array 'aa1' is not permitted in reduction" "" { target "*-*-*" } 65 }
 ! { dg-error "Array 'ia1' is not permitted in reduction" "" { target "*-*-*" } 67 }
-! { dg-error "Array 'la1' is not permitted in reduction" "" { target "*-*-*" } 71 }
-! { dg-error "Array 'ta1' is not permitted in reduction" "" { target "*-*-*" } 77 }
-! { dg-error "Array 'ia2' is not permitted in reduction" "" { target "*-*-*" } 81 }
-! { dg-error "Array 'ra1' is not permitted in reduction" "" { target "*-*-*" } 85 }
-! { dg-error "Array 'da1' is not permitted in reduction" "" { target "*-*-*" } 89 }
-! { dg-error "Array 'ca1' is not permitted in reduction" "" { target "*-*-*" } 93 }
-! { dg-error "Array 'ta1' is not permitted in reduction" "" { target "*-*-*" } 99 }
-! { dg-error "Array 'ca1' is not permitted in reduction" "" { target "*-*-*" } 103 }
-! { dg-error "Array 'la1' is not permitted in reduction" "" { target "*-*-*" } 107 }
-! { dg-error "Array 'ta1' is not permitted in reduction" "" { target "*-*-*" } 113 }
-! { dg-error "Array 'ra1' is not permitted in reduction" "" { target "*-*-*" } 117 }
-! { dg-error "Array 'da1' is not permitted in reduction" "" { target "*-*-*" } 121 }
-! { dg-error "Array 'ca1' is not permitted in reduction" "" { target "*-*-*" } 125 }
-! { dg-error "Array 'la1' is not permitted in reduction" "" { target "*-*-*" } 129 }
-! { dg-error "Array 'ta1' is not permitted in reduction" "" { target "*-*-*" } 135 }
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/reduction-10.f90 b/libgomp/testsuite/libgomp.oacc-fortran/reduction-10.f90
new file mode 100644
index 00000000000..506dfaf29f6
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/reduction-10.f90
@@ -0,0 +1,483 @@ 
+! { dg-do run }
+
+! real array reductions
+
+program reduction_10
+  implicit none
+
+  integer, parameter     :: n = 10, ng = 8, nw = 4, vl = 32
+  integer                :: i, j
+  real, dimension (n) :: vresult, rg, rw, rv, rc
+  logical, dimension (n) :: lrg, lrw, lrv, lrc, lvresult
+  real, dimension (n) :: array
+
+  do i = 1, n
+     array(i) = i
+  end do
+
+  !
+  ! '+' reductions
+  !
+
+  rg = 0
+  rw = 0
+  rv = 0
+  rc = 0
+  vresult = 0
+
+  !$acc parallel num_gangs(ng) copy(rg)
+  !$acc loop reduction(+:rg) gang
+  do i = 1, n
+    do j = 1, n
+      rg(j) = rg(j) + array(i)
+    end do
+  end do
+  !$acc end parallel
+
+  !$acc parallel num_workers(nw) copy(rw)
+  !$acc loop reduction(+:rw) worker
+  do i = 1, n
+    do j = 1, n
+      rw(j) = rw(j) + array(i)
+    end do
+  end do
+  !$acc end parallel
+
+  !$acc parallel vector_length(vl) copy(rv)
+  !$acc loop reduction(+:rv) vector
+  do i = 1, n
+    do j = 1, n
+      rv(j) = rv(j) + array(i)
+    end do
+  end do
+  !$acc end parallel
+
+  !$acc parallel num_gangs(ng) num_workers(nw) vector_length(vl) copy(rc)
+  !$acc loop reduction(+:rc) gang worker vector
+  do i = 1, n
+    do j = 1, n
+      rc(j) = rc(j) + array(i)
+    end do
+  end do
+  !$acc end parallel
+
+  ! Verify the results
+  do i = 1, n
+    do j = 1, n
+      vresult(j) = vresult(j) + array(i)
+    end do
+  end do
+
+  if (count (rg .ne. vresult) .ne. 0) STOP 1
+  if (count (rw .ne. vresult) .ne. 0) STOP 2
+  if (count (rv .ne. vresult) .ne. 0) STOP 3
+  if (count (rc .ne. vresult) .ne. 0) STOP 4
+
+  !
+  ! '*' reductions
+  !
+
+  rg = 1
+  rw = 1
+  rv = 1
+  rc = 1
+  vresult = 1
+
+  !$acc parallel num_gangs(ng) copy(rg)
+  !$acc loop reduction(*:rg) gang
+  do i = 1, n
+    do j = 1, n
+      rg(j) = rg(j) * array(i)
+    end do
+  end do
+  !$acc end parallel
+
+  !$acc parallel num_workers(nw) copy(rw)
+  !$acc loop reduction(*:rw) worker
+  do i = 1, n
+    do j = 1, n
+      rw(j) = rw(j) * array(i)
+    end do
+  end do
+  !$acc end parallel
+
+  !$acc parallel vector_length(vl) copy(rv)
+  !$acc loop reduction(*:rv) vector
+  do i = 1, n
+    do j = 1, n
+      rv(j) = rv(j) * array(i)
+    end do
+  end do
+  !$acc end parallel
+
+  !$acc parallel num_gangs(ng) num_workers(nw) vector_length(vl) copy(rc)
+  !$acc loop reduction(*:rc) gang worker vector
+  do i = 1, n
+    do j = 1, n
+      rc(j) = rc(j) * array(i)
+    end do
+  end do
+  !$acc end parallel
+
+  ! Verify the results
+  do i = 1, n
+    do j = 1, n
+      vresult(j) = vresult(j) * array(i)
+    end do
+  end do
+
+  if (count (rg .ne. vresult) .ne. 0) STOP 5
+  if (count (rw .ne. vresult) .ne. 0) STOP 6
+  if (count (rv .ne. vresult) .ne. 0) STOP 7
+  if (count (rc .ne. vresult) .ne. 0) STOP 8
+
+  !
+  ! 'max' reductions
+  !
+
+  rg = 0
+  rw = 0
+  rv = 0
+  rc = 0
+  vresult = 0
+
+  !$acc parallel num_gangs(ng) copy(rg)
+  !$acc loop reduction(max:rg) gang
+  do i = 1, n
+    do j = 1, n
+      rg(j) = max (rg(j), array(i))
+    end do
+  end do
+  !$acc end parallel
+
+  !$acc parallel num_workers(nw) copy(rw)
+  !$acc loop reduction(max:rw) worker
+  do i = 1, n
+    do j = 1, n
+      rw(j) = max (rw(j), array(i))
+    end do
+  end do
+  !$acc end parallel
+
+  !$acc parallel vector_length(vl) copy(rv)
+  !$acc loop reduction(max:rv) vector
+  do i = 1, n
+    do j = 1, n
+      rv(j) = max (rv(j), array(i))
+    end do
+  end do
+  !$acc end parallel
+
+  !$acc parallel num_gangs(ng) num_workers(nw) vector_length(vl) copy(rc)
+  !$acc loop reduction(max:rc) gang worker vector
+  do i = 1, n
+    do j = 1, n
+      rc(j) = max (rc(j), array(i))
+    end do
+  end do
+  !$acc end parallel
+
+  ! Verify the results
+  do i = 1, n
+    do j = 1, n
+      vresult(j) = max (vresult(j), array(i))
+    end do
+  end do
+
+  if (count (rg .ne. vresult) .ne. 0) STOP 9
+  if (count (rw .ne. vresult) .ne. 0) STOP 10
+  if (count (rv .ne. vresult) .ne. 0) STOP 11
+  if (count (rc .ne. vresult) .ne. 0) STOP 12
+
+  !
+  ! 'min' reductions
+  !
+
+  rg = 0
+  rw = 0
+  rv = 0
+  rc = 0
+  vresult = 0
+
+  !$acc parallel num_gangs(ng) copy(rg)
+  !$acc loop reduction(min:rg) gang
+  do i = 1, n
+    do j = 1, n
+      rg(j) = min (rg(j), array(i))
+    end do
+  end do
+  !$acc end parallel
+
+  !$acc parallel num_workers(nw) copy(rw)
+  !$acc loop reduction(min:rw) worker
+  do i = 1, n
+    do j = 1, n
+      rw(j) = min (rw(j), array(i))
+    end do
+  end do
+  !$acc end parallel
+
+  !$acc parallel vector_length(vl) copy(rv)
+  !$acc loop reduction(min:rv) vector
+  do i = 1, n
+    do j = 1, n
+      rv(j) = min (rv(j), array(i))
+    end do
+  end do
+  !$acc end parallel
+
+  !$acc parallel num_gangs(ng) num_workers(nw) vector_length(vl) copy(rc)
+  !$acc loop reduction(min:rc) gang worker vector
+  do i = 1, n
+    do j = 1, n
+      rc(j) = min (rc(j), array(i))
+    end do
+  end do
+  !$acc end parallel
+
+  ! Verify the results
+  do i = 1, n
+    do j = 1, n
+      vresult(j) = min (vresult(j), array(i))
+    end do
+  end do
+
+  if (count (rg .ne. vresult) .ne. 0) STOP 13
+  if (count (rw .ne. vresult) .ne. 0) STOP 14
+  if (count (rv .ne. vresult) .ne. 0) STOP 15
+  if (count (rc .ne. vresult) .ne. 0) STOP 16
+
+  !
+  ! '.and.' reductions
+  !
+
+  lrg = .true.
+  lrw = .true.
+  lrv = .true.
+  lrc = .true.
+  lvresult = .true.
+
+  !$acc parallel num_gangs(ng) copy(lrg)
+  !$acc loop reduction(.and.:lrg) gang
+  do i = 1, n
+    do j = 1, n
+      lrg(j) = lrg(j) .and. (array(i) .ge. 5)
+    end do
+  end do
+  !$acc end parallel
+
+  !$acc parallel num_workers(nw) copy(lrw)
+  !$acc loop reduction(.and.:lrw) worker
+  do i = 1, n
+    do j = 1, n
+      lrw(j) = lrw(j) .and. (array(i) .ge. 5)
+    end do
+  end do
+  !$acc end parallel
+
+  !$acc parallel vector_length(vl) copy(lrv)
+  !$acc loop reduction(.and.:lrv) vector
+  do i = 1, n
+    do j = 1, n
+      lrv(j) = lrv(j) .and. (array(i) .ge. 5)
+    end do
+  end do
+  !$acc end parallel
+
+  !$acc parallel num_gangs(ng) num_workers(nw) vector_length(vl) copy(lrc)
+  !$acc loop reduction(.and.:lrc) gang worker vector
+  do i = 1, n
+    do j = 1, n
+      lrc(j) = lrc(j) .and. (array(i) .ge. 5)
+    end do
+  end do
+  !$acc end parallel
+
+  ! Verify the results
+  do i = 1, n
+    do j = 1, n
+      lvresult(j) = lvresult(j) .and. (array(i) .ge. 5)
+    end do
+  end do
+
+  if (count (lrg .neqv. lvresult) .ne. 0) STOP 17
+  if (count (lrw .neqv. lvresult) .ne. 0) STOP 18
+  if (count (lrv .neqv. lvresult) .ne. 0) STOP 19
+  if (count (lrc .neqv. lvresult) .ne. 0) STOP 20
+
+  !
+  ! '.or.' reductions
+  !
+
+  lrg = .true.
+  lrw = .true.
+  lrv = .true.
+  lrc = .true.
+  lvresult = .true.
+
+  !$acc parallel num_gangs(ng) copy(lrg)
+  !$acc loop reduction(.or.:lrg) gang
+  do i = 1, n
+    do j = 1, n
+      lrg(j) = lrg(j) .or. (array(i) .ge. 5)
+    end do
+  end do
+  !$acc end parallel
+
+  !$acc parallel num_workers(nw) copy(lrw)
+  !$acc loop reduction(.or.:lrw) worker
+  do i = 1, n
+    do j = 1, n
+      lrw(j) = lrw(j) .or. (array(i) .ge. 5)
+    end do
+  end do
+  !$acc end parallel
+
+  !$acc parallel vector_length(vl) copy(lrv)
+  !$acc loop reduction(.or.:lrv) vector
+  do i = 1, n
+    do j = 1, n
+      lrv(j) = lrv(j) .or. (array(i) .ge. 5)
+    end do
+  end do
+  !$acc end parallel
+
+  !$acc parallel num_gangs(ng) num_workers(nw) vector_length(vl) copy(lrc)
+  !$acc loop reduction(.or.:lrc) gang worker vector
+  do i = 1, n
+    do j = 1, n
+      lrc(j) = lrc(j) .or. (array(i) .ge. 5)
+    end do
+  end do
+  !$acc end parallel
+
+  ! Verify the results
+  do i = 1, n
+    do j = 1, n
+      lvresult(j) = lvresult(j) .or. (array(i) .ge. 5)
+    end do
+  end do
+
+  if (count (lrg .neqv. lvresult) .ne. 0) STOP 21
+  if (count (lrw .neqv. lvresult) .ne. 0) STOP 22
+  if (count (lrv .neqv. lvresult) .ne. 0) STOP 23
+  if (count (lrc .neqv. lvresult) .ne. 0) STOP 24
+
+  !
+  ! '.eqv.' reductions
+  !
+
+  lrg = .true.
+  lrw = .true.
+  lrv = .true.
+  lrc = .true.
+  lvresult = .true.
+
+  !$acc parallel num_gangs(ng) copy(lrg)
+  !$acc loop reduction(.eqv.:lrg) gang
+  do i = 1, n
+    do j = 1, n
+      lrg(j) = lrg(j) .eqv. (array(i) .ge. 5)
+    end do
+  end do
+  !$acc end parallel
+
+  !$acc parallel num_workers(nw) copy(lrw)
+  !$acc loop reduction(.eqv.:lrw) worker
+  do i = 1, n
+    do j = 1, n
+      lrw(j) = lrw(j) .eqv. (array(i) .ge. 5)
+    end do
+  end do
+  !$acc end parallel
+
+  !$acc parallel vector_length(vl) copy(lrv)
+  !$acc loop reduction(.eqv.:lrv) vector
+  do i = 1, n
+    do j = 1, n
+      lrv(j) = lrv(j) .eqv. (array(i) .ge. 5)
+    end do
+  end do
+  !$acc end parallel
+
+  !$acc parallel num_gangs(ng) num_workers(nw) vector_length(vl) copy(lrc)
+  !$acc loop reduction(.eqv.:lrc) gang worker vector
+  do i = 1, n
+    do j = 1, n
+      lrc(j) = lrc(j) .eqv. (array(i) .ge. 5)
+    end do
+  end do
+  !$acc end parallel
+
+  ! Verify the results
+  do i = 1, n
+    do j = 1, n
+      lvresult(j) = lvresult(j) .eqv. (array(i) .ge. 5)
+    end do
+  end do
+
+  if (count (lrg .neqv. lvresult) .ne. 0) STOP 25
+  if (count (lrw .neqv. lvresult) .ne. 0) STOP 26
+  if (count (lrv .neqv. lvresult) .ne. 0) STOP 27
+  if (count (lrc .neqv. lvresult) .ne. 0) STOP 28
+
+  !
+  ! '.neqv.' reductions
+  !
+
+  lrg = .true.
+  lrw = .true.
+  lrv = .true.
+  lrc = .true.
+  lvresult = .true.
+
+  !$acc parallel num_gangs(ng) copy(lrg)
+  !$acc loop reduction(.neqv.:lrg) gang
+  do i = 1, n
+    do j = 1, n
+      lrg(j) = lrg(j) .neqv. (array(i) .ge. 5)
+    end do
+  end do
+  !$acc end parallel
+
+  !$acc parallel num_workers(nw) copy(lrw)
+  !$acc loop reduction(.neqv.:lrw) worker
+  do i = 1, n
+    do j = 1, n
+      lrw(j) = lrw(j) .neqv. (array(i) .ge. 5)
+    end do
+  end do
+  !$acc end parallel
+
+  !$acc parallel vector_length(vl) copy(lrv)
+  !$acc loop reduction(.neqv.:lrv) vector
+  do i = 1, n
+    do j = 1, n
+      lrv(j) = lrv(j) .neqv. (array(i) .ge. 5)
+    end do
+  end do
+  !$acc end parallel
+
+  !$acc parallel num_gangs(ng) num_workers(nw) vector_length(vl) copy(lrc)
+  !$acc loop reduction(.neqv.:lrc) gang worker vector
+  do i = 1, n
+    do j = 1, n
+      lrc(j) = lrc(j) .neqv. (array(i) .ge. 5)
+    end do
+  end do
+  !$acc end parallel
+
+  ! Verify the results
+  do i = 1, n
+    do j = 1, n
+      lvresult(j) = lvresult(j) .neqv. (array(i) .ge. 5)
+    end do
+  end do
+
+  if (count (lrg .neqv. lvresult) .ne. 0) STOP 29
+  if (count (lrw .neqv. lvresult) .ne. 0) STOP 30
+  if (count (lrv .neqv. lvresult) .ne. 0) STOP 31
+  if (count (lrc .neqv. lvresult) .ne. 0) STOP 32
+
+end program reduction_10
+
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/reduction-11.f90 b/libgomp/testsuite/libgomp.oacc-fortran/reduction-11.f90
new file mode 100644
index 00000000000..4bec1c797cd
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/reduction-11.f90
@@ -0,0 +1,483 @@ 
+! { dg-do run }
+
+! double precision array reductions
+
+program reduction_11
+  implicit none
+
+  integer, parameter     :: n = 10, ng = 8, nw = 4, vl = 32
+  integer                :: i, j
+  double precision, dimension (n) :: vresult, rg, rw, rv, rc
+  logical, dimension (n) :: lrg, lrw, lrv, lrc, lvresult
+  double precision, dimension (n) :: array
+
+  do i = 1, n
+     array(i) = i
+  end do
+
+  !
+  ! '+' reductions
+  !
+
+  rg = 0
+  rw = 0
+  rv = 0
+  rc = 0
+  vresult = 0
+
+  !$acc parallel num_gangs(ng) copy(rg)
+  !$acc loop reduction(+:rg) gang
+  do i = 1, n
+    do j = 1, n
+      rg(j) = rg(j) + array(i)
+    end do
+  end do
+  !$acc end parallel
+
+  !$acc parallel num_workers(nw) copy(rw)
+  !$acc loop reduction(+:rw) worker
+  do i = 1, n
+    do j = 1, n
+      rw(j) = rw(j) + array(i)
+    end do
+  end do
+  !$acc end parallel
+
+  !$acc parallel vector_length(vl) copy(rv)
+  !$acc loop reduction(+:rv) vector
+  do i = 1, n
+    do j = 1, n
+      rv(j) = rv(j) + array(i)
+    end do
+  end do
+  !$acc end parallel
+
+  !$acc parallel num_gangs(ng) num_workers(nw) vector_length(vl) copy(rc)
+  !$acc loop reduction(+:rc) gang worker vector
+  do i = 1, n
+    do j = 1, n
+      rc(j) = rc(j) + array(i)
+    end do
+  end do
+  !$acc end parallel
+
+  ! Verify the results
+  do i = 1, n
+    do j = 1, n
+      vresult(j) = vresult(j) + array(i)
+    end do
+  end do
+
+  if (count (rg .ne. vresult) .ne. 0) STOP 1
+  if (count (rw .ne. vresult) .ne. 0) STOP 2
+  if (count (rv .ne. vresult) .ne. 0) STOP 3
+  if (count (rc .ne. vresult) .ne. 0) STOP 4
+
+  !
+  ! '*' reductions
+  !
+
+  rg = 1
+  rw = 1
+  rv = 1
+  rc = 1
+  vresult = 1
+
+  !$acc parallel num_gangs(ng) copy(rg)
+  !$acc loop reduction(*:rg) gang
+  do i = 1, n
+    do j = 1, n
+      rg(j) = rg(j) * array(i)
+    end do
+  end do
+  !$acc end parallel
+
+  !$acc parallel num_workers(nw) copy(rw)
+  !$acc loop reduction(*:rw) worker
+  do i = 1, n
+    do j = 1, n
+      rw(j) = rw(j) * array(i)
+    end do
+  end do
+  !$acc end parallel
+
+  !$acc parallel vector_length(vl) copy(rv)
+  !$acc loop reduction(*:rv) vector
+  do i = 1, n
+    do j = 1, n
+      rv(j) = rv(j) * array(i)
+    end do
+  end do
+  !$acc end parallel
+
+  !$acc parallel num_gangs(ng) num_workers(nw) vector_length(vl) copy(rc)
+  !$acc loop reduction(*:rc) gang worker vector
+  do i = 1, n
+    do j = 1, n
+      rc(j) = rc(j) * array(i)
+    end do
+  end do
+  !$acc end parallel
+
+  ! Verify the results
+  do i = 1, n
+    do j = 1, n
+      vresult(j) = vresult(j) * array(i)
+    end do
+  end do
+
+  if (count (rg .ne. vresult) .ne. 0) STOP 5
+  if (count (rw .ne. vresult) .ne. 0) STOP 6
+  if (count (rv .ne. vresult) .ne. 0) STOP 7
+  if (count (rc .ne. vresult) .ne. 0) STOP 8
+
+  !
+  ! 'max' reductions
+  !
+
+  rg = 0
+  rw = 0
+  rv = 0
+  rc = 0
+  vresult = 0
+
+  !$acc parallel num_gangs(ng) copy(rg)
+  !$acc loop reduction(max:rg) gang
+  do i = 1, n
+    do j = 1, n
+      rg(j) = max (rg(j), array(i))
+    end do
+  end do
+  !$acc end parallel
+
+  !$acc parallel num_workers(nw) copy(rw)
+  !$acc loop reduction(max:rw) worker
+  do i = 1, n
+    do j = 1, n
+      rw(j) = max (rw(j), array(i))
+    end do
+  end do
+  !$acc end parallel
+
+  !$acc parallel vector_length(vl) copy(rv)
+  !$acc loop reduction(max:rv) vector
+  do i = 1, n
+    do j = 1, n
+      rv(j) = max (rv(j), array(i))
+    end do
+  end do
+  !$acc end parallel
+
+  !$acc parallel num_gangs(ng) num_workers(nw) vector_length(vl) copy(rc)
+  !$acc loop reduction(max:rc) gang worker vector
+  do i = 1, n
+    do j = 1, n
+      rc(j) = max (rc(j), array(i))
+    end do
+  end do
+  !$acc end parallel
+
+  ! Verify the results
+  do i = 1, n
+    do j = 1, n
+      vresult(j) = max (vresult(j), array(i))
+    end do
+  end do
+
+  if (count (rg .ne. vresult) .ne. 0) STOP 9
+  if (count (rw .ne. vresult) .ne. 0) STOP 10
+  if (count (rv .ne. vresult) .ne. 0) STOP 11
+  if (count (rc .ne. vresult) .ne. 0) STOP 12
+
+  !
+  ! 'min' reductions
+  !
+
+  rg = 0
+  rw = 0
+  rv = 0
+  rc = 0
+  vresult = 0
+
+  !$acc parallel num_gangs(ng) copy(rg)
+  !$acc loop reduction(min:rg) gang
+  do i = 1, n
+    do j = 1, n
+      rg(j) = min (rg(j), array(i))
+    end do
+  end do
+  !$acc end parallel
+
+  !$acc parallel num_workers(nw) copy(rw)
+  !$acc loop reduction(min:rw) worker
+  do i = 1, n
+    do j = 1, n
+      rw(j) = min (rw(j), array(i))
+    end do
+  end do
+  !$acc end parallel
+
+  !$acc parallel vector_length(vl) copy(rv)
+  !$acc loop reduction(min:rv) vector
+  do i = 1, n
+    do j = 1, n
+      rv(j) = min (rv(j), array(i))
+    end do
+  end do
+  !$acc end parallel
+
+  !$acc parallel num_gangs(ng) num_workers(nw) vector_length(vl) copy(rc)
+  !$acc loop reduction(min:rc) gang worker vector
+  do i = 1, n
+    do j = 1, n
+      rc(j) = min (rc(j), array(i))
+    end do
+  end do
+  !$acc end parallel
+
+  ! Verify the results
+  do i = 1, n
+    do j = 1, n
+      vresult(j) = min (vresult(j), array(i))
+    end do
+  end do
+
+  if (count (rg .ne. vresult) .ne. 0) STOP 13
+  if (count (rw .ne. vresult) .ne. 0) STOP 14
+  if (count (rv .ne. vresult) .ne. 0) STOP 15
+  if (count (rc .ne. vresult) .ne. 0) STOP 16
+
+  !
+  ! '.and.' reductions
+  !
+
+  lrg = .true.
+  lrw = .true.
+  lrv = .true.
+  lrc = .true.
+  lvresult = .true.
+
+  !$acc parallel num_gangs(ng) copy(lrg)
+  !$acc loop reduction(.and.:lrg) gang
+  do i = 1, n
+    do j = 1, n
+      lrg(j) = lrg(j) .and. (array(i) .ge. 5)
+    end do
+  end do
+  !$acc end parallel
+
+  !$acc parallel num_workers(nw) copy(lrw)
+  !$acc loop reduction(.and.:lrw) worker
+  do i = 1, n
+    do j = 1, n
+      lrw(j) = lrw(j) .and. (array(i) .ge. 5)
+    end do
+  end do
+  !$acc end parallel
+
+  !$acc parallel vector_length(vl) copy(lrv)
+  !$acc loop reduction(.and.:lrv) vector
+  do i = 1, n
+    do j = 1, n
+      lrv(j) = lrv(j) .and. (array(i) .ge. 5)
+    end do
+  end do
+  !$acc end parallel
+
+  !$acc parallel num_gangs(ng) num_workers(nw) vector_length(vl) copy(lrc)
+  !$acc loop reduction(.and.:lrc) gang worker vector
+  do i = 1, n
+    do j = 1, n
+      lrc(j) = lrc(j) .and. (array(i) .ge. 5)
+    end do
+  end do
+  !$acc end parallel
+
+  ! Verify the results
+  do i = 1, n
+    do j = 1, n
+      lvresult(j) = lvresult(j) .and. (array(i) .ge. 5)
+    end do
+  end do
+
+  if (count (lrg .neqv. lvresult) .ne. 0) STOP 17
+  if (count (lrw .neqv. lvresult) .ne. 0) STOP 18
+  if (count (lrv .neqv. lvresult) .ne. 0) STOP 19
+  if (count (lrc .neqv. lvresult) .ne. 0) STOP 20
+
+  !
+  ! '.or.' reductions
+  !
+
+  lrg = .true.
+  lrw = .true.
+  lrv = .true.
+  lrc = .true.
+  lvresult = .true.
+
+  !$acc parallel num_gangs(ng) copy(lrg)
+  !$acc loop reduction(.or.:lrg) gang
+  do i = 1, n
+    do j = 1, n
+      lrg(j) = lrg(j) .or. (array(i) .ge. 5)
+    end do
+  end do
+  !$acc end parallel
+
+  !$acc parallel num_workers(nw) copy(lrw)
+  !$acc loop reduction(.or.:lrw) worker
+  do i = 1, n
+    do j = 1, n
+      lrw(j) = lrw(j) .or. (array(i) .ge. 5)
+    end do
+  end do
+  !$acc end parallel
+
+  !$acc parallel vector_length(vl) copy(lrv)
+  !$acc loop reduction(.or.:lrv) vector
+  do i = 1, n
+    do j = 1, n
+      lrv(j) = lrv(j) .or. (array(i) .ge. 5)
+    end do
+  end do
+  !$acc end parallel
+
+  !$acc parallel num_gangs(ng) num_workers(nw) vector_length(vl) copy(lrc)
+  !$acc loop reduction(.or.:lrc) gang worker vector
+  do i = 1, n
+    do j = 1, n
+      lrc(j) = lrc(j) .or. (array(i) .ge. 5)
+    end do
+  end do
+  !$acc end parallel
+
+  ! Verify the results
+  do i = 1, n
+    do j = 1, n
+      lvresult(j) = lvresult(j) .or. (array(i) .ge. 5)
+    end do
+  end do
+
+  if (count (lrg .neqv. lvresult) .ne. 0) STOP 21
+  if (count (lrw .neqv. lvresult) .ne. 0) STOP 22
+  if (count (lrv .neqv. lvresult) .ne. 0) STOP 23
+  if (count (lrc .neqv. lvresult) .ne. 0) STOP 24
+
+  !
+  ! '.eqv.' reductions
+  !
+
+  lrg = .true.
+  lrw = .true.
+  lrv = .true.
+  lrc = .true.
+  lvresult = .true.
+
+  !$acc parallel num_gangs(ng) copy(lrg)
+  !$acc loop reduction(.eqv.:lrg) gang
+  do i = 1, n
+    do j = 1, n
+      lrg(j) = lrg(j) .eqv. (array(i) .ge. 5)
+    end do
+  end do
+  !$acc end parallel
+
+  !$acc parallel num_workers(nw) copy(lrw)
+  !$acc loop reduction(.eqv.:lrw) worker
+  do i = 1, n
+    do j = 1, n
+      lrw(j) = lrw(j) .eqv. (array(i) .ge. 5)
+    end do
+  end do
+  !$acc end parallel
+
+  !$acc parallel vector_length(vl) copy(lrv)
+  !$acc loop reduction(.eqv.:lrv) vector
+  do i = 1, n
+    do j = 1, n
+      lrv(j) = lrv(j) .eqv. (array(i) .ge. 5)
+    end do
+  end do
+  !$acc end parallel
+
+  !$acc parallel num_gangs(ng) num_workers(nw) vector_length(vl) copy(lrc)
+  !$acc loop reduction(.eqv.:lrc) gang worker vector
+  do i = 1, n
+    do j = 1, n
+      lrc(j) = lrc(j) .eqv. (array(i) .ge. 5)
+    end do
+  end do
+  !$acc end parallel
+
+  ! Verify the results
+  do i = 1, n
+    do j = 1, n
+      lvresult(j) = lvresult(j) .eqv. (array(i) .ge. 5)
+    end do
+  end do
+
+  if (count (lrg .neqv. lvresult) .ne. 0) STOP 25
+  if (count (lrw .neqv. lvresult) .ne. 0) STOP 26
+  if (count (lrv .neqv. lvresult) .ne. 0) STOP 27
+  if (count (lrc .neqv. lvresult) .ne. 0) STOP 28
+
+  !
+  ! '.neqv.' reductions
+  !
+
+  lrg = .true.
+  lrw = .true.
+  lrv = .true.
+  lrc = .true.
+  lvresult = .true.
+
+  !$acc parallel num_gangs(ng) copy(lrg)
+  !$acc loop reduction(.neqv.:lrg) gang
+  do i = 1, n
+    do j = 1, n
+      lrg(j) = lrg(j) .neqv. (array(i) .ge. 5)
+    end do
+  end do
+  !$acc end parallel
+
+  !$acc parallel num_workers(nw) copy(lrw)
+  !$acc loop reduction(.neqv.:lrw) worker
+  do i = 1, n
+    do j = 1, n
+      lrw(j) = lrw(j) .neqv. (array(i) .ge. 5)
+    end do
+  end do
+  !$acc end parallel
+
+  !$acc parallel vector_length(vl) copy(lrv)
+  !$acc loop reduction(.neqv.:lrv) vector
+  do i = 1, n
+    do j = 1, n
+      lrv(j) = lrv(j) .neqv. (array(i) .ge. 5)
+    end do
+  end do
+  !$acc end parallel
+
+  !$acc parallel num_gangs(ng) num_workers(nw) vector_length(vl) copy(lrc)
+  !$acc loop reduction(.neqv.:lrc) gang worker vector
+  do i = 1, n
+    do j = 1, n
+      lrc(j) = lrc(j) .neqv. (array(i) .ge. 5)
+    end do
+  end do
+  !$acc end parallel
+
+  ! Verify the results
+  do i = 1, n
+    do j = 1, n
+      lvresult(j) = lvresult(j) .neqv. (array(i) .ge. 5)
+    end do
+  end do
+
+  if (count (lrg .neqv. lvresult) .ne. 0) STOP 29
+  if (count (lrw .neqv. lvresult) .ne. 0) STOP 30
+  if (count (lrv .neqv. lvresult) .ne. 0) STOP 31
+  if (count (lrc .neqv. lvresult) .ne. 0) STOP 32
+
+end program reduction_11
+
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/reduction-12.f90 b/libgomp/testsuite/libgomp.oacc-fortran/reduction-12.f90
new file mode 100644
index 00000000000..b609c7a294e
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/reduction-12.f90
@@ -0,0 +1,135 @@ 
+! { dg-do run }
+
+! complex array reductions
+
+program reduction_12
+  implicit none
+
+  integer, parameter     :: n = 10, ng = 8, nw = 4, vl = 32
+  integer                :: i, j
+  complex, dimension (n) :: vresult, rg, rw, rv, rc
+  logical, dimension (n) :: lrg, lrw, lrv, lrc, lvresult
+  complex, dimension (n) :: array
+
+  do i = 1, n
+     array(i) = i
+  end do
+
+  !
+  ! '+' reductions
+  !
+
+  rg = 0
+  rw = 0
+  rv = 0
+  rc = 0
+  vresult = 0
+
+  !$acc parallel num_gangs(ng) copy(rg)
+  !$acc loop reduction(+:rg) gang
+  do i = 1, n
+    do j = 1, n
+      rg(j) = rg(j) + array(i)
+    end do
+  end do
+  !$acc end parallel
+
+  !$acc parallel num_workers(nw) copy(rw)
+  !$acc loop reduction(+:rw) worker
+  do i = 1, n
+    do j = 1, n
+      rw(j) = rw(j) + array(i)
+    end do
+  end do
+  !$acc end parallel
+
+  !$acc parallel vector_length(vl) copy(rv)
+  !$acc loop reduction(+:rv) vector
+  do i = 1, n
+    do j = 1, n
+      rv(j) = rv(j) + array(i)
+    end do
+  end do
+  !$acc end parallel
+
+  !$acc parallel num_gangs(ng) num_workers(nw) vector_length(vl) copy(rc)
+  !$acc loop reduction(+:rc) gang worker vector
+  do i = 1, n
+    do j = 1, n
+      rc(j) = rc(j) + array(i)
+    end do
+  end do
+  !$acc end parallel
+
+  ! Verify the results
+  do i = 1, n
+    do j = 1, n
+      vresult(j) = vresult(j) + array(i)
+    end do
+  end do
+
+  if (count (rg .ne. vresult) .ne. 0) STOP 1
+  if (count (rw .ne. vresult) .ne. 0) STOP 2
+  if (count (rv .ne. vresult) .ne. 0) STOP 3
+  if (count (rc .ne. vresult) .ne. 0) STOP 4
+
+  !
+  ! '*' reductions
+  !
+
+  rg = 1
+  rw = 1
+  rv = 1
+  rc = 1
+  vresult = 1
+
+  !$acc parallel num_gangs(ng) copy(rg)
+  !$acc loop reduction(*:rg) gang
+  do i = 1, n
+    do j = 1, n
+      rg(j) = rg(j) * array(i)
+    end do
+  end do
+  !$acc end parallel
+
+  !$acc parallel num_workers(nw) copy(rw)
+  !$acc loop reduction(*:rw) worker
+  do i = 1, n
+    do j = 1, n
+      rw(j) = rw(j) * array(i)
+    end do
+  end do
+  !$acc end parallel
+
+  !$acc parallel vector_length(vl) copy(rv)
+  !$acc loop reduction(*:rv) vector
+  do i = 1, n
+    do j = 1, n
+      rv(j) = rv(j) * array(i)
+    end do
+  end do
+  !$acc end parallel
+
+  !$acc parallel num_gangs(ng) num_workers(nw) vector_length(vl) copy(rc)
+  !$acc loop reduction(*:rc) gang worker vector
+  do i = 1, n
+    do j = 1, n
+      rc(j) = rc(j) * array(i)
+    end do
+  end do
+  !$acc end parallel
+
+  ! Verify the results
+  do i = 1, n
+    do j = 1, n
+      vresult(j) = vresult(j) * array(i)
+    end do
+  end do
+
+  if (count (rg .ne. vresult) .ne. 0) STOP 5
+  if (count (rw .ne. vresult) .ne. 0) STOP 6
+  if (count (rv .ne. vresult) .ne. 0) STOP 7
+  if (count (rc .ne. vresult) .ne. 0) STOP 8
+
+end program reduction_12
+
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/reduction-13.f90 b/libgomp/testsuite/libgomp.oacc-fortran/reduction-13.f90
new file mode 100644
index 00000000000..088c5cd3b04
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/reduction-13.f90
@@ -0,0 +1,66 @@ 
+! { dg-do run }
+
+! record type reductions
+
+program reduction_13
+  implicit none
+
+  type t1
+     integer :: i
+     real :: r
+  end type t1
+
+  type t2
+     real :: r
+     integer :: i
+     double precision :: d
+  end type t2
+
+  integer, parameter :: n = 10, ng = 8, nw = 4, vl = 32
+  integer :: i
+  type(t1) :: v1, a1
+  type (t2) :: v2, a2
+
+  v1%i = 0
+  v1%r = 0
+  !$acc parallel num_gangs(ng) num_workers(nw) vector_length(vl) copy(v1)
+  !$acc loop reduction (+:v1)
+  do i = 1, n
+     v1%i = v1%i + 1
+     v1%r = v1%r + 2
+  end do
+  !$acc end parallel
+  a1%i = 0
+  a1%r = 0
+  do i = 1, n
+     a1%i = a1%i + 1
+     a1%r = a1%r + 2
+  end do
+  if (v1%i .ne. a1%i) STOP 1
+  if (v1%r .ne. a1%r) STOP 2
+
+  v2%i = 1
+  v2%r = 1
+  v2%d = 1
+  !$acc parallel num_gangs(ng) num_workers(nw) vector_length(vl) copy(v2)
+  !$acc loop reduction (*:v2)
+  do i = 1, n
+     v2%i = v2%i * 2
+     v2%r = v2%r * 1.1
+     v2%d = v2%d * 1.3
+  end do
+  !$acc end parallel
+  a2%i = 1
+  a2%r = 1
+  a2%d = 1
+  do i = 1, n
+     a2%i = a2%i * 2
+     a2%r = a2%r * 1.1
+     a2%d = a2%d * 1.3
+  end do
+
+  if (v2%i .ne. a2%i) STOP 3
+  if (v2%r .ne. a2%r) STOP 4
+  if (v2%d .ne. a2%d) STOP 5
+
+end program reduction_13
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/reduction-9.f90 b/libgomp/testsuite/libgomp.oacc-fortran/reduction-9.f90
new file mode 100644
index 00000000000..43ab155aa73
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/reduction-9.f90
@@ -0,0 +1,657 @@ 
+! { dg-do run }
+
+! integer array reductions
+
+program reduction_9
+  implicit none
+
+  integer, parameter     :: n = 10, ng = 8, nw = 4, vl = 32
+  integer                :: i, j
+  integer, dimension (n) :: vresult, rg, rw, rv, rc
+  logical, dimension (n) :: lrg, lrw, lrv, lrc, lvresult
+  integer, dimension (n) :: array
+
+  do i = 1, n
+     array(i) = i
+  end do
+
+  !
+  ! '+' reductions
+  !
+
+  rg = 0
+  rw = 0
+  rv = 0
+  rc = 0
+  vresult = 0
+
+  !$acc parallel num_gangs(ng) copy(rg)
+  !$acc loop reduction(+:rg) gang
+  do i = 1, n
+    do j = 1, n
+      rg(j) = rg(j) + array(i)
+    end do
+  end do
+  !$acc end parallel
+
+  !$acc parallel num_workers(nw) copy(rw)
+  !$acc loop reduction(+:rw) worker
+  do i = 1, n
+    do j = 1, n
+      rw(j) = rw(j) + array(i)
+    end do
+  end do
+  !$acc end parallel
+
+  !$acc parallel vector_length(vl) copy(rv)
+  !$acc loop reduction(+:rv) vector
+  do i = 1, n
+    do j = 1, n
+      rv(j) = rv(j) + array(i)
+    end do
+  end do
+  !$acc end parallel
+
+  !$acc parallel num_gangs(ng) num_workers(nw) vector_length(vl) copy(rc)
+  !$acc loop reduction(+:rc) gang worker vector
+  do i = 1, n
+    do j = 1, n
+      rc(j) = rc(j) + array(i)
+    end do
+  end do
+  !$acc end parallel
+
+  ! Verify the results
+  do i = 1, n
+    do j = 1, n
+      vresult(j) = vresult(j) + array(i)
+    end do
+  end do
+
+  if (count (rg .ne. vresult) .ne. 0) STOP 1
+  if (count (rw .ne. vresult) .ne. 0) STOP 2
+  if (count (rv .ne. vresult) .ne. 0) STOP 3
+  if (count (rc .ne. vresult) .ne. 0) STOP 4
+
+  !
+  ! '*' reductions
+  !
+
+  rg = 1
+  rw = 1
+  rv = 1
+  rc = 1
+  vresult = 1
+
+  !$acc parallel num_gangs(ng) copy(rg)
+  !$acc loop reduction(*:rg) gang
+  do i = 1, n
+    do j = 1, n
+      rg(j) = rg(j) * array(i)
+    end do
+  end do
+  !$acc end parallel
+
+  !$acc parallel num_workers(nw) copy(rw)
+  !$acc loop reduction(*:rw) worker
+  do i = 1, n
+    do j = 1, n
+      rw(j) = rw(j) * array(i)
+    end do
+  end do
+  !$acc end parallel
+
+  !$acc parallel vector_length(vl) copy(rv)
+  !$acc loop reduction(*:rv) vector
+  do i = 1, n
+    do j = 1, n
+      rv(j) = rv(j) * array(i)
+    end do
+  end do
+  !$acc end parallel
+
+  !$acc parallel num_gangs(ng) num_workers(nw) vector_length(vl) copy(rc)
+  !$acc loop reduction(*:rc) gang worker vector
+  do i = 1, n
+    do j = 1, n
+      rc(j) = rc(j) * array(i)
+    end do
+  end do
+  !$acc end parallel
+
+  ! Verify the results
+  do i = 1, n
+    do j = 1, n
+      vresult(j) = vresult(j) * array(i)
+    end do
+  end do
+
+  if (count (rg .ne. vresult) .ne. 0) STOP 5
+  if (count (rw .ne. vresult) .ne. 0) STOP 6
+  if (count (rv .ne. vresult) .ne. 0) STOP 7
+  if (count (rc .ne. vresult) .ne. 0) STOP 8
+
+  !
+  ! 'max' reductions
+  !
+
+  rg = 0
+  rw = 0
+  rv = 0
+  rc = 0
+  vresult = 0
+
+  !$acc parallel num_gangs(ng) copy(rg)
+  !$acc loop reduction(max:rg) gang
+  do i = 1, n
+    do j = 1, n
+      rg(j) = max (rg(j), array(i))
+    end do
+  end do
+  !$acc end parallel
+
+  !$acc parallel num_workers(nw) copy(rw)
+  !$acc loop reduction(max:rw) worker
+  do i = 1, n
+    do j = 1, n
+      rw(j) = max (rw(j), array(i))
+    end do
+  end do
+  !$acc end parallel
+
+  !$acc parallel vector_length(vl) copy(rv)
+  !$acc loop reduction(max:rv) vector
+  do i = 1, n
+    do j = 1, n
+      rv(j) = max (rv(j), array(i))
+    end do
+  end do
+  !$acc end parallel
+
+  !$acc parallel num_gangs(ng) num_workers(nw) vector_length(vl) copy(rc)
+  !$acc loop reduction(max:rc) gang worker vector
+  do i = 1, n
+    do j = 1, n
+      rc(j) = max (rc(j), array(i))
+    end do
+  end do
+  !$acc end parallel
+
+  ! Verify the results
+  do i = 1, n
+    do j = 1, n
+      vresult(j) = max (vresult(j), array(i))
+    end do
+  end do
+
+  if (count (rg .ne. vresult) .ne. 0) STOP 9
+  if (count (rw .ne. vresult) .ne. 0) STOP 10
+  if (count (rv .ne. vresult) .ne. 0) STOP 11
+  if (count (rc .ne. vresult) .ne. 0) STOP 12
+
+  !
+  ! 'min' reductions
+  !
+
+  rg = 0
+  rw = 0
+  rv = 0
+  rc = 0
+  vresult = 0
+
+  !$acc parallel num_gangs(ng) copy(rg)
+  !$acc loop reduction(min:rg) gang
+  do i = 1, n
+    do j = 1, n
+      rg(j) = min (rg(j), array(i))
+    end do
+  end do
+  !$acc end parallel
+
+  !$acc parallel num_workers(nw) copy(rw)
+  !$acc loop reduction(min:rw) worker
+  do i = 1, n
+    do j = 1, n
+      rw(j) = min (rw(j), array(i))
+    end do
+  end do
+  !$acc end parallel
+
+  !$acc parallel vector_length(vl) copy(rv)
+  !$acc loop reduction(min:rv) vector
+  do i = 1, n
+    do j = 1, n
+      rv(j) = min (rv(j), array(i))
+    end do
+  end do
+  !$acc end parallel
+
+  !$acc parallel num_gangs(ng) num_workers(nw) vector_length(vl) copy(rc)
+  !$acc loop reduction(min:rc) gang worker vector
+  do i = 1, n
+    do j = 1, n
+      rc(j) = min (rc(j), array(i))
+    end do
+  end do
+  !$acc end parallel
+
+  ! Verify the results
+  do i = 1, n
+    do j = 1, n
+      vresult(j) = min (vresult(j), array(i))
+    end do
+  end do
+
+  if (count (rg .ne. vresult) .ne. 0) STOP 13
+  if (count (rw .ne. vresult) .ne. 0) STOP 14
+  if (count (rv .ne. vresult) .ne. 0) STOP 15
+  if (count (rc .ne. vresult) .ne. 0) STOP 16
+
+  !
+  ! 'iand' reductions
+  !
+
+  rg = 1
+  rw = 1
+  rv = 1
+  rc = 1
+  vresult = 1
+
+  !$acc parallel num_gangs(ng) copy(rg)
+  !$acc loop reduction(iand:rg) gang
+  do i = 1, n
+    do j = 1, n
+      rg(j) = iand (rg(j), array(i))
+    end do
+  end do
+  !$acc end parallel
+
+  !$acc parallel num_workers(nw) copy(rw)
+  !$acc loop reduction(iand:rw) worker
+  do i = 1, n
+    do j = 1, n
+      rw(j) = iand (rw(j), array(i))
+    end do
+  end do
+  !$acc end parallel
+
+  !$acc parallel vector_length(vl) copy(rv)
+  !$acc loop reduction(iand:rv) vector
+  do i = 1, n
+    do j = 1, n
+      rv(j) = iand (rv(j), array(i))
+    end do
+  end do
+  !$acc end parallel
+
+  !$acc parallel num_gangs(ng) num_workers(nw) vector_length(vl) copy(rc)
+  !$acc loop reduction(iand:rc) gang worker vector
+  do i = 1, n
+    do j = 1, n
+      rc(j) = iand (rc(j), array(i))
+    end do
+  end do
+  !$acc end parallel
+
+  ! Verify the results
+  do i = 1, n
+    do j = 1, n
+      vresult(j) = iand (vresult(j), array(i))
+    end do
+  end do
+
+  if (count (rg .ne. vresult) .ne. 0) STOP 17
+  if (count (rw .ne. vresult) .ne. 0) STOP 18
+  if (count (rv .ne. vresult) .ne. 0) STOP 19
+  if (count (rc .ne. vresult) .ne. 0) STOP 20
+
+  !
+  ! 'ior' reductions
+  !
+
+  rg = 0
+  rw = 0
+  rv = 0
+  rc = 0
+  vresult = 0
+
+  !$acc parallel num_gangs(ng) copy(rg)
+  !$acc loop reduction(ior:rg) gang
+  do i = 1, n
+    do j = 1, n
+      rg(j) = ior (rg(j), array(i))
+    end do
+  end do
+  !$acc end parallel
+
+  !$acc parallel num_workers(nw) copy(rw)
+  !$acc loop reduction(ior:rw) worker
+  do i = 1, n
+    do j = 1, n
+      rw(j) = ior (rw(j), array(i))
+    end do
+  end do
+  !$acc end parallel
+
+  !$acc parallel vector_length(vl) copy(rv)
+  !$acc loop reduction(ior:rv) vector
+  do i = 1, n
+    do j = 1, n
+      rv(j) = ior (rv(j), array(i))
+    end do
+  end do
+  !$acc end parallel
+
+  !$acc parallel num_gangs(ng) num_workers(nw) vector_length(vl) copy(rc)
+  !$acc loop reduction(ior:rc) gang worker vector
+  do i = 1, n
+    do j = 1, n
+      rc(j) = ior (rc(j), array(i))
+    end do
+  end do
+  !$acc end parallel
+
+  ! Verify the results
+  do i = 1, n
+    do j = 1, n
+      vresult(j) = ior (vresult(j), array(i))
+    end do
+  end do
+
+  if (count (rg .ne. vresult) .ne. 0) STOP 21
+  if (count (rw .ne. vresult) .ne. 0) STOP 22
+  if (count (rv .ne. vresult) .ne. 0) STOP 23
+  if (count (rc .ne. vresult) .ne. 0) STOP 24
+
+  !
+  ! 'ieor' reductions
+  !
+
+  rg = 0
+  rw = 0
+  rv = 0
+  rc = 0
+  vresult = 0
+
+  !$acc parallel num_gangs(ng) copy(rg)
+  !$acc loop reduction(ieor:rg) gang
+  do i = 1, n
+    do j = 1, n
+      rg(j) = ieor (rg(j), array(i))
+    end do
+  end do
+  !$acc end parallel
+
+  !$acc parallel num_workers(nw) copy(rw)
+  !$acc loop reduction(ieor:rw) worker
+  do i = 1, n
+    do j = 1, n
+      rw(j) = ieor (rw(j), array(i))
+    end do
+  end do
+  !$acc end parallel
+
+  !$acc parallel vector_length(vl) copy(rv)
+  !$acc loop reduction(ieor:rv) vector
+  do i = 1, n
+    do j = 1, n
+      rv(j) = ieor (rv(j), array(i))
+    end do
+  end do
+  !$acc end parallel
+
+  !$acc parallel num_gangs(ng) num_workers(nw) vector_length(vl) copy(rc)
+  !$acc loop reduction(ieor:rc) gang worker vector
+  do i = 1, n
+    do j = 1, n
+      rc(j) = ieor (rc(j), array(i))
+    end do
+  end do
+  !$acc end parallel
+
+  ! Verify the results
+  do i = 1, n
+    do j = 1, n
+      vresult(j) = ieor (vresult(j), array(i))
+    end do
+  end do
+
+  if (count (rg .ne. vresult) .ne. 0) STOP 25
+  if (count (rw .ne. vresult) .ne. 0) STOP 26
+  if (count (rv .ne. vresult) .ne. 0) STOP 27
+  if (count (rc .ne. vresult) .ne. 0) STOP 28
+
+  !
+  ! '.and.' reductions
+  !
+
+  lrg = .true.
+  lrw = .true.
+  lrv = .true.
+  lrc = .true.
+  lvresult = .true.
+
+  !$acc parallel num_gangs(ng) copy(lrg)
+  !$acc loop reduction(.and.:lrg) gang
+  do i = 1, n
+    do j = 1, n
+      lrg(j) = lrg(j) .and. (array(i) .ge. 5)
+    end do
+  end do
+  !$acc end parallel
+
+  !$acc parallel num_workers(nw) copy(lrw)
+  !$acc loop reduction(.and.:lrw) worker
+  do i = 1, n
+    do j = 1, n
+      lrw(j) = lrw(j) .and. (array(i) .ge. 5)
+    end do
+  end do
+  !$acc end parallel
+
+  !$acc parallel vector_length(vl) copy(lrv)
+  !$acc loop reduction(.and.:lrv) vector
+  do i = 1, n
+    do j = 1, n
+      lrv(j) = lrv(j) .and. (array(i) .ge. 5)
+    end do
+  end do
+  !$acc end parallel
+
+  !$acc parallel num_gangs(ng) num_workers(nw) vector_length(vl) copy(lrc)
+  !$acc loop reduction(.and.:lrc) gang worker vector
+  do i = 1, n
+    do j = 1, n
+      lrc(j) = lrc(j) .and. (array(i) .ge. 5)
+    end do
+  end do
+  !$acc end parallel
+
+  ! Verify the results
+  do i = 1, n
+    do j = 1, n
+      lvresult(j) = lvresult(j) .and. (array(i) .ge. 5)
+    end do
+  end do
+
+  if (count (lrg .neqv. lvresult) .ne. 0) STOP 29
+  if (count (lrw .neqv. lvresult) .ne. 0) STOP 30
+  if (count (lrv .neqv. lvresult) .ne. 0) STOP 31
+  if (count (lrc .neqv. lvresult) .ne. 0) STOP 32
+
+  !
+  ! '.or.' reductions
+  !
+
+  lrg = .true.
+  lrw = .true.
+  lrv = .true.
+  lrc = .true.
+  lvresult = .true.
+
+  !$acc parallel num_gangs(ng) copy(lrg)
+  !$acc loop reduction(.or.:lrg) gang
+  do i = 1, n
+    do j = 1, n
+      lrg(j) = lrg(j) .or. (array(i) .ge. 5)
+    end do
+  end do
+  !$acc end parallel
+
+  !$acc parallel num_workers(nw) copy(lrw)
+  !$acc loop reduction(.or.:lrw) worker
+  do i = 1, n
+    do j = 1, n
+      lrw(j) = lrw(j) .or. (array(i) .ge. 5)
+    end do
+  end do
+  !$acc end parallel
+
+  !$acc parallel vector_length(vl) copy(lrv)
+  !$acc loop reduction(.or.:lrv) vector
+  do i = 1, n
+    do j = 1, n
+      lrv(j) = lrv(j) .or. (array(i) .ge. 5)
+    end do
+  end do
+  !$acc end parallel
+
+  !$acc parallel num_gangs(ng) num_workers(nw) vector_length(vl) copy(lrc)
+  !$acc loop reduction(.or.:lrc) gang worker vector
+  do i = 1, n
+    do j = 1, n
+      lrc(j) = lrc(j) .or. (array(i) .ge. 5)
+    end do
+  end do
+  !$acc end parallel
+
+  ! Verify the results
+  do i = 1, n
+    do j = 1, n
+      lvresult(j) = lvresult(j) .or. (array(i) .ge. 5)
+    end do
+  end do
+
+  if (count (lrg .neqv. lvresult) .ne. 0) STOP 33
+  if (count (lrw .neqv. lvresult) .ne. 0) STOP 34
+  if (count (lrv .neqv. lvresult) .ne. 0) STOP 35
+  if (count (lrc .neqv. lvresult) .ne. 0) STOP 36
+
+  !
+  ! '.eqv.' reductions
+  !
+
+  lrg = .true.
+  lrw = .true.
+  lrv = .true.
+  lrc = .true.
+  lvresult = .true.
+
+  !$acc parallel num_gangs(ng) copy(lrg)
+  !$acc loop reduction(.eqv.:lrg) gang
+  do i = 1, n
+    do j = 1, n
+      lrg(j) = lrg(j) .eqv. (array(i) .ge. 5)
+    end do
+  end do
+  !$acc end parallel
+
+  !$acc parallel num_workers(nw) copy(lrw)
+  !$acc loop reduction(.eqv.:lrw) worker
+  do i = 1, n
+    do j = 1, n
+      lrw(j) = lrw(j) .eqv. (array(i) .ge. 5)
+    end do
+  end do
+  !$acc end parallel
+
+  !$acc parallel vector_length(vl) copy(lrv)
+  !$acc loop reduction(.eqv.:lrv) vector
+  do i = 1, n
+    do j = 1, n
+      lrv(j) = lrv(j) .eqv. (array(i) .ge. 5)
+    end do
+  end do
+  !$acc end parallel
+
+  !$acc parallel num_gangs(ng) num_workers(nw) vector_length(vl) copy(lrc)
+  !$acc loop reduction(.eqv.:lrc) gang worker vector
+  do i = 1, n
+    do j = 1, n
+      lrc(j) = lrc(j) .eqv. (array(i) .ge. 5)
+    end do
+  end do
+  !$acc end parallel
+
+  ! Verify the results
+  do i = 1, n
+    do j = 1, n
+      lvresult(j) = lvresult(j) .eqv. (array(i) .ge. 5)
+    end do
+  end do
+
+  if (count (lrg .neqv. lvresult) .ne. 0) STOP 37
+  if (count (lrw .neqv. lvresult) .ne. 0) STOP 38
+  if (count (lrv .neqv. lvresult) .ne. 0) STOP 39
+  if (count (lrc .neqv. lvresult) .ne. 0) STOP 40
+
+  !
+  ! '.neqv.' reductions
+  !
+
+  lrg = .true.
+  lrw = .true.
+  lrv = .true.
+  lrc = .true.
+  lvresult = .true.
+
+  !$acc parallel num_gangs(ng) copy(lrg)
+  !$acc loop reduction(.neqv.:lrg) gang
+  do i = 1, n
+    do j = 1, n
+      lrg(j) = lrg(j) .neqv. (array(i) .ge. 5)
+    end do
+  end do
+  !$acc end parallel
+
+  !$acc parallel num_workers(nw) copy(lrw)
+  !$acc loop reduction(.neqv.:lrw) worker
+  do i = 1, n
+    do j = 1, n
+      lrw(j) = lrw(j) .neqv. (array(i) .ge. 5)
+    end do
+  end do
+  !$acc end parallel
+
+  !$acc parallel vector_length(vl) copy(lrv)
+  !$acc loop reduction(.neqv.:lrv) vector
+  do i = 1, n
+    do j = 1, n
+      lrv(j) = lrv(j) .neqv. (array(i) .ge. 5)
+    end do
+  end do
+  !$acc end parallel
+
+  !$acc parallel num_gangs(ng) num_workers(nw) vector_length(vl) copy(lrc)
+  !$acc loop reduction(.neqv.:lrc) gang worker vector
+  do i = 1, n
+    do j = 1, n
+      lrc(j) = lrc(j) .neqv. (array(i) .ge. 5)
+    end do
+  end do
+  !$acc end parallel
+
+  ! Verify the results
+  do i = 1, n
+    do j = 1, n
+      lvresult(j) = lvresult(j) .neqv. (array(i) .ge. 5)
+    end do
+  end do
+
+  if (count (lrg .neqv. lvresult) .ne. 0) STOP 41
+  if (count (lrw .neqv. lvresult) .ne. 0) STOP 42
+  if (count (lrv .neqv. lvresult) .ne. 0) STOP 43
+  if (count (lrc .neqv. lvresult) .ne. 0) STOP 44
+
+end program reduction_9
+