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 <clt...@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.
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
+
