Hi,
OpenACC requires that, if a variable is used in reduction clauses on two nested
loops, then there
must be reduction clauses for that variable on all loops that are nested in
between the two loops
and all these reduction clauses must use the same operator; this has been first
clarified by
OpenACC 2.6. This commit introduces a check for that property which reports
errors if the property
is violated.
I have tested the patch by comparing "make check" results and I am not aware of
any regressions.
Gergö has implemented the check and it works, but I was wondering if the way in
which the patch
avoids issuing errors about operator switches more than once by modifying the
clauses (cf. the
corresponding comment in omp-low.c) could lead to problems - the processing
might still continue
after the error on the modified tree, right? I was also wondering about the
best place for such
checks. Should this be a part of "pass_lower_omp" (as in the patch) or should
it run earlier
like, for instance, "pass_diagnose_omp_blocks".
Can the patch be included in trunk?
Frederik
>From 99796969c1bf91048c6383dfb1b8576bdd9efd7d Mon Sep 17 00:00:00 2001
From: Frederik Harwath <frede...@codesourcery.com>
Date: Mon, 21 Oct 2019 08:27:58 +0200
Subject: [PATCH] Report errors on inconsistent OpenACC nested reduction
clauses
MIME-Version: 1.0
Content-Type: text/plain; charset=UTF-8
Content-Transfer-Encoding: 8bit
OpenACC (cf. OpenACC 2.7, section 2.9.11. "reduction clause";
this was first clarified by OpenACC 2.6) requires that, if a
variable is used in reduction clauses on two nested loops, then
there must be reduction clauses for that variable on all loops
that are nested in between the two loops and all these reduction
clauses must use the same operator.
This commit introduces a check for that property which reports
errors if it is violated.
In gcc/testsuite/c-c++-common/goacc/reduction-6.c, we remove the erroneous
reductions on variable b; adding a reduction clause to make it compile cleanly
would make it a duplicate of the test for variable c.
2010-10-21 Gergö Barany <ge...@codesourcery.com>
Frederik Harwath <frede...@codesourcery.com>
gcc/
* omp-low.c (struct omp_context): New fields
local_reduction_clauses, outer_reduction_clauses.
(new_omp_context): Initialize these.
(scan_sharing_clauses): Record reduction clauses on OpenACC
constructs.
(scan_omp_for): Check reduction clauses for incorrect nesting.
gcc/testsuite/
* c-c++-common/goacc/nested-reductions-fail.c: New test.
* c-c++-common/goacc/nested-reductions.c: New test.
* c-c++-common/goacc/reduction-6.c: Adjust.
libgomp/
* testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-1.c:
Add missing reduction clauses.
* testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-2.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-3.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-4.c:
Likewise.
---
gcc/omp-low.c | 107 +++-
.../goacc/nested-reductions-fail.c | 492 ++++++++++++++++++
.../c-c++-common/goacc/nested-reductions.c | 420 +++++++++++++++
.../c-c++-common/goacc/reduction-6.c | 11 -
.../par-loop-comb-reduction-1.c | 2 +-
.../par-loop-comb-reduction-2.c | 2 +-
.../par-loop-comb-reduction-3.c | 2 +-
.../par-loop-comb-reduction-4.c | 2 +-
8 files changed, 1022 insertions(+), 16 deletions(-)
create mode 100644 gcc/testsuite/c-c++-common/goacc/nested-reductions-fail.c
create mode 100644 gcc/testsuite/c-c++-common/goacc/nested-reductions.c
diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index 279b6ef893a..a2212274685 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -127,6 +127,12 @@ struct omp_context
corresponding tracking loop iteration variables. */
hash_map<tree, tree> *lastprivate_conditional_map;
+ /* A tree_list of the reduction clauses in this context. */
+ tree local_reduction_clauses;
+
+ /* A tree_list of the reduction clauses in outer contexts. */
+ tree outer_reduction_clauses;
+
/* Nesting depth of this context. Used to beautify error messages re
invalid gotos. The outermost ctx is depth 1, with depth 0 being
reserved for the main body of the function. */
@@ -902,6 +908,8 @@ new_omp_context (gimple *stmt, omp_context *outer_ctx)
ctx->cb = outer_ctx->cb;
ctx->cb.block = NULL;
ctx->depth = outer_ctx->depth + 1;
+ ctx->local_reduction_clauses = NULL;
+ ctx->outer_reduction_clauses = ctx->outer_reduction_clauses;
}
else
{
@@ -917,6 +925,8 @@ new_omp_context (gimple *stmt, omp_context *outer_ctx)
ctx->cb.adjust_array_error_bounds = true;
ctx->cb.dont_remap_vla_if_no_change = true;
ctx->depth = 1;
+ ctx->local_reduction_clauses = NULL;
+ ctx->outer_reduction_clauses = NULL;
}
ctx->cb.decl_map = new hash_map<tree, tree>;
@@ -1131,6 +1141,9 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
case OMP_CLAUSE_REDUCTION:
case OMP_CLAUSE_IN_REDUCTION:
+ if (is_oacc_parallel (ctx) || is_oacc_kernels (ctx))
+ ctx->local_reduction_clauses
+ = tree_cons (NULL, c, ctx->local_reduction_clauses);
decl = OMP_CLAUSE_DECL (c);
if (TREE_CODE (decl) == MEM_REF)
{
@@ -2410,7 +2423,99 @@ scan_omp_for (gomp_for *stmt, omp_context *outer_ctx)
gimple_omp_for_set_clauses (stmt, clauses);
check_oacc_kernel_gwv (stmt, ctx);
}
- }
+
+ /* Collect all variables named in reductions on this loop. Ensure
+ that, if this loop has a reduction on some variable v, and there is
+ a reduction on v somewhere in an outer context, then there is a
+ reduction on v on all intervening loops as well. */
+ tree local_reduction_clauses = NULL;
+ for (tree c = gimple_omp_for_clauses (stmt); c; c = OMP_CLAUSE_CHAIN (c))
+ {
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_REDUCTION)
+ local_reduction_clauses
+ = tree_cons (NULL, c, local_reduction_clauses);
+ }
+ if (ctx->outer_reduction_clauses == NULL && ctx->outer != NULL)
+ ctx->outer_reduction_clauses
+ = chainon (unshare_expr (ctx->outer->local_reduction_clauses),
+ ctx->outer->outer_reduction_clauses);
+ tree outer_reduction_clauses = ctx->outer_reduction_clauses;
+ tree local_iter = local_reduction_clauses;
+ for (; local_iter; local_iter = TREE_CHAIN (local_iter))
+ {
+ tree local_clause = TREE_VALUE (local_iter);
+ tree local_var = OMP_CLAUSE_DECL (local_clause);
+ tree_code local_op = OMP_CLAUSE_REDUCTION_CODE (local_clause);
+ bool have_outer_reduction = false;
+ tree ctx_iter = outer_reduction_clauses;
+ for (; ctx_iter; ctx_iter = TREE_CHAIN (ctx_iter))
+ {
+ tree outer_clause = TREE_VALUE (ctx_iter);
+ tree outer_var = OMP_CLAUSE_DECL (outer_clause);
+ tree_code outer_op = OMP_CLAUSE_REDUCTION_CODE (outer_clause);
+ if (outer_var == local_var && outer_op != local_op)
+ {
+ error_at (gimple_location (stmt),
+ "conflicting reduction operations for %qE",
+ local_var);
+ inform (OMP_CLAUSE_LOCATION (outer_clause),
+ "location of the previous reduction for %qE",
+ outer_var);
+ /* Change this operation to be equal to the outer one.
+ This is meant to suppress spurious errors; for example,
+ in nested +, -, + reductions, we would generate errors
+ for both the change from + to - and from - to +. */
+ OMP_CLAUSE_REDUCTION_CODE (local_clause) = outer_op;
+ /* Also change the location so that in nested +, -, -
+ reductions, the second error message also refers to the
+ outermost + reduction. */
+ OMP_CLAUSE_LOCATION (local_clause)
+ = OMP_CLAUSE_LOCATION (outer_clause);
+ }
+ if (outer_var == local_var)
+ {
+ have_outer_reduction = true;
+ break;
+ }
+ }
+ if (have_outer_reduction)
+ {
+ /* There is a reduction on outer_var both on this loop and on
+ some enclosing loop. Walk up the context tree until such a
+ loop with a reduction on outer_var is found, and complain
+ about all intervening loops that do not have such a
+ reduction. */
+ struct omp_context *curr_loop = ctx->outer;
+ bool found = false;
+ while (curr_loop != NULL)
+ {
+ tree curr_iter = curr_loop->local_reduction_clauses;
+ for (; curr_iter; curr_iter = TREE_CHAIN (curr_iter))
+ {
+ tree curr_clause = TREE_VALUE (curr_iter);
+ tree curr_var = OMP_CLAUSE_DECL (curr_clause);
+ if (curr_var == local_var)
+ {
+ found = true;
+ break;
+ }
+ }
+ if (!found)
+ error_at (gimple_location (curr_loop->stmt),
+ "nested loop in reduction needs "
+ "reduction clause for %qE",
+ local_var);
+ else
+ break;
+ curr_loop = curr_loop->outer;
+ }
+ }
+ }
+ ctx->local_reduction_clauses = local_reduction_clauses;
+ ctx->outer_reduction_clauses
+ = chainon (unshare_expr (ctx->local_reduction_clauses),
+ ctx->outer_reduction_clauses);
+ }
scan_sharing_clauses (clauses, ctx);
diff --git a/gcc/testsuite/c-c++-common/goacc/nested-reductions-fail.c b/gcc/testsuite/c-c++-common/goacc/nested-reductions-fail.c
new file mode 100644
index 00000000000..a642dd03870
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/nested-reductions-fail.c
@@ -0,0 +1,492 @@
+/* Test erroneous cases of nested reduction loops. */
+
+void acc_parallel (void)
+{
+ int i, j, k, l, sum, diff;
+
+ #pragma acc parallel
+ {
+ #pragma acc loop reduction(+:sum)
+ for (i = 0; i < 10; i++)
+ #pragma acc loop // { dg-error "nested loop in reduction needs reduction clause for .sum." }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(+:sum)
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ #pragma acc loop reduction(+:sum)
+ for (i = 0; i < 10; i++)
+ #pragma acc loop collapse(2) // { dg-error "nested loop in reduction needs reduction clause for .sum." }
+ for (j = 0; j < 10; j++)
+ for (k = 0; k < 10; k++)
+ #pragma acc loop reduction(+:sum)
+ for (l = 0; l < 10; l++)
+ sum = 1;
+
+ #pragma acc loop reduction(+:sum)
+ for (i = 0; i < 10; i++)
+ #pragma acc loop // { dg-error "nested loop in reduction needs reduction clause for .sum." }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop // { dg-error "nested loop in reduction needs reduction clause for .sum." }
+ for (k = 0; k < 10; k++)
+ #pragma acc loop reduction(+:sum)
+ for (l = 0; l < 10; l++)
+ sum = 1;
+
+ #pragma acc loop reduction(+:sum)
+ for (i = 0; i < 10; i++)
+ #pragma acc loop reduction(-:sum) // { dg-error "conflicting reduction operations for .sum." }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(+:sum)
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ #pragma acc loop reduction(+:sum)
+ for (i = 0; i < 10; i++)
+ #pragma acc loop reduction(-:sum) // { dg-error "conflicting reduction operations for .sum." }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(-:sum) // { dg-error "conflicting reduction operations for .sum." }
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ #pragma acc loop reduction(+:sum)
+ for (i = 0; i < 10; i++)
+ #pragma acc loop reduction(-:sum) // { dg-error "conflicting reduction operations for .sum." }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop // { dg-error "nested loop in reduction needs reduction clause for .sum." }
+ for (k = 0; k < 10; k++)
+ #pragma acc loop reduction(*:sum) // { dg-error "conflicting reduction operations for .sum." }
+ for (l = 0; l < 10; l++)
+ sum = 1;
+
+ #pragma acc loop reduction(+:sum)
+ for (i = 0; i < 10; i++)
+ #pragma acc loop reduction(-:sum) // { dg-error "conflicting reduction operations for .sum." }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(+:sum)
+ for (k = 0; k < 10; k++)
+ #pragma acc loop reduction(*:sum) // { dg-error "conflicting reduction operations for .sum." }
+ for (l = 0; l < 10; l++)
+ sum = 1;
+
+ #pragma acc loop reduction(+:sum) reduction(-:diff)
+ for (i = 0; i < 10; i++)
+ {
+ #pragma acc loop reduction(-:diff) // { dg-error "nested loop in reduction needs reduction clause for .sum." }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(+:sum)
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ #pragma acc loop reduction(+:sum) // { dg-error "nested loop in reduction needs reduction clause for .diff." }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(-:diff)
+ for (k = 0; k < 10; k++)
+ diff = 1;
+ }
+ }
+}
+
+/* The same tests as above, but using a combined parallel loop construct. */
+
+void acc_parallel_loop (void)
+{
+ int i, j, k, l, sum, diff;
+
+ #pragma acc parallel loop
+ for (int h = 0; h < 10; ++h)
+ {
+ #pragma acc loop reduction(+:sum)
+ for (i = 0; i < 10; i++)
+ #pragma acc loop // { dg-error "nested loop in reduction needs reduction clause for .sum." }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(+:sum)
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ #pragma acc loop reduction(+:sum)
+ for (i = 0; i < 10; i++)
+ #pragma acc loop collapse(2) // { dg-error "nested loop in reduction needs reduction clause for .sum." }
+ for (j = 0; j < 10; j++)
+ for (k = 0; k < 10; k++)
+ #pragma acc loop reduction(+:sum)
+ for (l = 0; l < 10; l++)
+ sum = 1;
+
+ #pragma acc loop reduction(+:sum)
+ for (i = 0; i < 10; i++)
+ #pragma acc loop // { dg-error "nested loop in reduction needs reduction clause for .sum." }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop // { dg-error "nested loop in reduction needs reduction clause for .sum." }
+ for (k = 0; k < 10; k++)
+ #pragma acc loop reduction(+:sum)
+ for (l = 0; l < 10; l++)
+ sum = 1;
+
+ #pragma acc loop reduction(+:sum)
+ for (i = 0; i < 10; i++)
+ #pragma acc loop reduction(-:sum) // { dg-error "conflicting reduction operations for .sum." }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(+:sum)
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ #pragma acc loop reduction(+:sum)
+ for (i = 0; i < 10; i++)
+ #pragma acc loop reduction(-:sum) // { dg-error "conflicting reduction operations for .sum." }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(-:sum) // { dg-error "conflicting reduction operations for .sum." }
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ #pragma acc loop reduction(+:sum)
+ for (i = 0; i < 10; i++)
+ #pragma acc loop reduction(-:sum) // { dg-error "conflicting reduction operations for .sum." }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop // { dg-error "nested loop in reduction needs reduction clause for .sum." }
+ for (k = 0; k < 10; k++)
+ #pragma acc loop reduction(*:sum) // { dg-error "conflicting reduction operations for .sum." }
+ for (l = 0; l < 10; l++)
+ sum = 1;
+
+ #pragma acc loop reduction(+:sum)
+ for (i = 0; i < 10; i++)
+ #pragma acc loop reduction(-:sum) // { dg-error "conflicting reduction operations for .sum." }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(+:sum)
+ for (k = 0; k < 10; k++)
+ #pragma acc loop reduction(*:sum) // { dg-error "conflicting reduction operations for .sum." }
+ for (l = 0; l < 10; l++)
+ sum = 1;
+
+ #pragma acc loop reduction(+:sum) reduction(-:diff)
+ for (i = 0; i < 10; i++)
+ {
+ #pragma acc loop reduction(-:diff) // { dg-error "nested loop in reduction needs reduction clause for .sum." }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(+:sum)
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ #pragma acc loop reduction(+:sum) // { dg-error "nested loop in reduction needs reduction clause for .diff." }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(-:diff)
+ for (k = 0; k < 10; k++)
+ diff = 1;
+ }
+ }
+}
+
+/* The same tests as above, but now the outermost reduction clause is on
+ the parallel region, not the outermost loop. */
+void acc_parallel_reduction (void)
+{
+ int i, j, k, l, sum, diff;
+
+ #pragma acc parallel reduction(+:sum)
+ {
+ #pragma acc loop // { dg-error "nested loop in reduction needs reduction clause for .sum." }
+ for (i = 0; i < 10; i++)
+ #pragma acc loop // { dg-error "nested loop in reduction needs reduction clause for .sum." }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(+:sum)
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ #pragma acc loop // { dg-error "nested loop in reduction needs reduction clause for .sum." }
+ for (i = 0; i < 10; i++)
+ #pragma acc loop collapse(2) // { dg-error "nested loop in reduction needs reduction clause for .sum." }
+ for (j = 0; j < 10; j++)
+ for (k = 0; k < 10; k++)
+ #pragma acc loop reduction(+:sum)
+ for (l = 0; l < 10; l++)
+ sum = 1;
+
+ #pragma acc loop // { dg-error "nested loop in reduction needs reduction clause for .sum." }
+ for (i = 0; i < 10; i++)
+ #pragma acc loop // { dg-error "nested loop in reduction needs reduction clause for .sum." }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop // { dg-error "nested loop in reduction needs reduction clause for .sum." }
+ for (k = 0; k < 10; k++)
+ #pragma acc loop reduction(+:sum)
+ for (l = 0; l < 10; l++)
+ sum = 1;
+
+ #pragma acc loop // { dg-error "nested loop in reduction needs reduction clause for .sum." }
+ for (i = 0; i < 10; i++)
+ #pragma acc loop reduction(-:sum) // { dg-error "conflicting reduction operations for .sum." }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(+:sum)
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ #pragma acc loop // { dg-error "nested loop in reduction needs reduction clause for .sum." }
+ for (i = 0; i < 10; i++)
+ #pragma acc loop reduction(-:sum) // { dg-error "conflicting reduction operations for .sum." }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(-:sum) // { dg-error "conflicting reduction operations for .sum." }
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ #pragma acc loop reduction(max:sum) // { dg-error "conflicting reduction operations for .sum." }
+ for (i = 0; i < 10; i++)
+ #pragma acc loop reduction(-:sum) // { dg-error "conflicting reduction operations for .sum." }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop // { dg-error "nested loop in reduction needs reduction clause for .sum." }
+ for (k = 0; k < 10; k++)
+ #pragma acc loop reduction(*:sum) // { dg-error "conflicting reduction operations for .sum." }
+ for (l = 0; l < 10; l++)
+ sum = 1;
+
+ #pragma acc loop reduction(max:sum) // { dg-error "conflicting reduction operations for .sum." }
+ for (i = 0; i < 10; i++)
+ #pragma acc loop reduction(-:sum) // { dg-error "conflicting reduction operations for .sum." }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(+:sum)
+ for (k = 0; k < 10; k++)
+ #pragma acc loop reduction(*:sum) // { dg-error "conflicting reduction operations for .sum." }
+ for (l = 0; l < 10; l++)
+ sum = 1;
+
+ #pragma acc loop reduction(-:diff) // { dg-error "nested loop in reduction needs reduction clause for .sum." }
+ for (i = 0; i < 10; i++)
+ {
+ #pragma acc loop reduction(-:diff) // { dg-error "nested loop in reduction needs reduction clause for .sum." }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(+:sum)
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ #pragma acc loop reduction(+:sum) // { dg-error "nested loop in reduction needs reduction clause for .diff." }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(-:diff)
+ for (k = 0; k < 10; k++)
+ diff = 1;
+ }
+ }
+}
+
+/* The same tests as above, but using a combined parallel loop construct, and
+ the outermost reduction clause is on that one, not the outermost loop. */
+void acc_parallel_loop_reduction (void)
+{
+ int i, j, k, l, sum, diff;
+
+ #pragma acc parallel loop reduction(+:sum)
+ for (int h = 0; h < 10; ++h)
+ {
+ #pragma acc loop // { dg-error "nested loop in reduction needs reduction clause for .sum." }
+ for (i = 0; i < 10; i++)
+ #pragma acc loop // { dg-error "nested loop in reduction needs reduction clause for .sum." }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(+:sum)
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ #pragma acc loop // { dg-error "nested loop in reduction needs reduction clause for .sum." }
+ for (i = 0; i < 10; i++)
+ #pragma acc loop collapse(2) // { dg-error "nested loop in reduction needs reduction clause for .sum." }
+ for (j = 0; j < 10; j++)
+ for (k = 0; k < 10; k++)
+ #pragma acc loop reduction(+:sum)
+ for (l = 0; l < 10; l++)
+ sum = 1;
+
+ #pragma acc loop // { dg-error "nested loop in reduction needs reduction clause for .sum." }
+ for (i = 0; i < 10; i++)
+ #pragma acc loop // { dg-error "nested loop in reduction needs reduction clause for .sum." }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop // { dg-error "nested loop in reduction needs reduction clause for .sum." }
+ for (k = 0; k < 10; k++)
+ #pragma acc loop reduction(+:sum)
+ for (l = 0; l < 10; l++)
+ sum = 1;
+
+ #pragma acc loop // { dg-error "nested loop in reduction needs reduction clause for .sum." }
+ for (i = 0; i < 10; i++)
+ #pragma acc loop reduction(-:sum) // { dg-error "conflicting reduction operations for .sum." }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(+:sum)
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ #pragma acc loop // { dg-error "nested loop in reduction needs reduction clause for .sum." }
+ for (i = 0; i < 10; i++)
+ #pragma acc loop reduction(-:sum) // { dg-error "conflicting reduction operations for .sum." }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(-:sum) // { dg-error "conflicting reduction operations for .sum." }
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ #pragma acc loop reduction(max:sum) // { dg-error "conflicting reduction operations for .sum." }
+ for (i = 0; i < 10; i++)
+ #pragma acc loop reduction(-:sum) // { dg-error "conflicting reduction operations for .sum." }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop // { dg-error "nested loop in reduction needs reduction clause for .sum." }
+ for (k = 0; k < 10; k++)
+ #pragma acc loop reduction(*:sum) // { dg-error "conflicting reduction operations for .sum." }
+ for (l = 0; l < 10; l++)
+ sum = 1;
+
+ #pragma acc loop reduction(max:sum) // { dg-error "conflicting reduction operations for .sum." }
+ for (i = 0; i < 10; i++)
+ #pragma acc loop reduction(-:sum) // { dg-error "conflicting reduction operations for .sum." }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(+:sum)
+ for (k = 0; k < 10; k++)
+ #pragma acc loop reduction(*:sum) // { dg-error "conflicting reduction operations for .sum." }
+ for (l = 0; l < 10; l++)
+ sum = 1;
+
+ #pragma acc loop reduction(-:diff) // { dg-error "nested loop in reduction needs reduction clause for .sum." }
+ for (i = 0; i < 10; i++)
+ {
+ #pragma acc loop reduction(-:diff) // { dg-error "nested loop in reduction needs reduction clause for .sum." }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(+:sum)
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ #pragma acc loop reduction(+:sum) // { dg-error "nested loop in reduction needs reduction clause for .diff." }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(-:diff)
+ for (k = 0; k < 10; k++)
+ diff = 1;
+ }
+ }
+}
+
+/* The same tests as above, but inside a routine construct. */
+#pragma acc routine gang
+void acc_routine (void)
+{
+ int i, j, k, l, sum, diff;
+
+ {
+ #pragma acc loop reduction(+:sum)
+ for (i = 0; i < 10; i++)
+ #pragma acc loop // { dg-error "nested loop in reduction needs reduction clause for .sum." }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(+:sum)
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ #pragma acc loop reduction(+:sum)
+ for (i = 0; i < 10; i++)
+ #pragma acc loop collapse(2) // { dg-error "nested loop in reduction needs reduction clause for .sum." }
+ for (j = 0; j < 10; j++)
+ for (k = 0; k < 10; k++)
+ #pragma acc loop reduction(+:sum)
+ for (l = 0; l < 10; l++)
+ sum = 1;
+
+ #pragma acc loop reduction(+:sum)
+ for (i = 0; i < 10; i++)
+ #pragma acc loop // { dg-error "nested loop in reduction needs reduction clause for .sum." }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop // { dg-error "nested loop in reduction needs reduction clause for .sum." }
+ for (k = 0; k < 10; k++)
+ #pragma acc loop reduction(+:sum)
+ for (l = 0; l < 10; l++)
+ sum = 1;
+
+ #pragma acc loop reduction(+:sum)
+ for (i = 0; i < 10; i++)
+ #pragma acc loop reduction(-:sum) // { dg-error "conflicting reduction operations for .sum." }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(+:sum)
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ #pragma acc loop reduction(+:sum)
+ for (i = 0; i < 10; i++)
+ #pragma acc loop reduction(-:sum) // { dg-error "conflicting reduction operations for .sum." }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(-:sum) // { dg-error "conflicting reduction operations for .sum." }
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ #pragma acc loop reduction(+:sum)
+ for (i = 0; i < 10; i++)
+ #pragma acc loop reduction(-:sum) // { dg-error "conflicting reduction operations for .sum." }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop // { dg-error "nested loop in reduction needs reduction clause for .sum." }
+ for (k = 0; k < 10; k++)
+ #pragma acc loop reduction(*:sum) // { dg-error "conflicting reduction operations for .sum." }
+ for (l = 0; l < 10; l++)
+ sum = 1;
+
+ #pragma acc loop reduction(+:sum)
+ for (i = 0; i < 10; i++)
+ #pragma acc loop reduction(-:sum) // { dg-error "conflicting reduction operations for .sum." }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(+:sum)
+ for (k = 0; k < 10; k++)
+ #pragma acc loop reduction(*:sum) // { dg-error "conflicting reduction operations for .sum." }
+ for (l = 0; l < 10; l++)
+ sum = 1;
+
+ #pragma acc loop reduction(+:sum) reduction(-:diff)
+ for (i = 0; i < 10; i++)
+ {
+ #pragma acc loop reduction(-:diff) // { dg-error "nested loop in reduction needs reduction clause for .sum." }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(+:sum)
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ #pragma acc loop reduction(+:sum) // { dg-error "nested loop in reduction needs reduction clause for .diff." }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(-:diff)
+ for (k = 0; k < 10; k++)
+ diff = 1;
+ }
+ }
+}
+
+void acc_kernels (void)
+{
+ int i, j, k, sum, diff;
+
+ /* FIXME: No diagnostics are produced for these loops because reductions
+ in kernels regions are not supported yet. */
+ #pragma acc kernels
+ {
+ #pragma acc loop reduction(+:sum)
+ for (i = 0; i < 10; i++)
+ for (j = 0; j < 10; j++)
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ #pragma acc loop reduction(+:sum)
+ for (i = 0; i < 10; i++)
+ #pragma acc loop
+ for (j = 0; j < 10; j++)
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ #pragma acc loop reduction(+:sum)
+ for (i = 0; i < 10; i++)
+ #pragma acc loop reduction(-:diff)
+ for (j = 0; j < 10; j++)
+ #pragma acc loop
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ #pragma acc loop reduction(+:sum)
+ for (i = 0; i < 10; i++)
+ #pragma acc loop
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(+:sum)
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ #pragma acc loop reduction(+:sum)
+ for (i = 0; i < 10; i++)
+ #pragma acc loop reduction(-:sum)
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(+:sum)
+ for (k = 0; k < 10; k++)
+ sum = 1;
+ }
+}
diff --git a/gcc/testsuite/c-c++-common/goacc/nested-reductions.c b/gcc/testsuite/c-c++-common/goacc/nested-reductions.c
new file mode 100644
index 00000000000..15385c4da99
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/nested-reductions.c
@@ -0,0 +1,420 @@
+/* Test cases of nested reduction loops that should compile cleanly. */
+
+void acc_parallel (void)
+{
+ int i, j, k, sum, diff;
+
+ #pragma acc parallel
+ {
+ #pragma acc loop reduction(+:sum)
+ for (i = 0; i < 10; i++)
+ for (j = 0; j < 10; j++)
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ #pragma acc loop collapse(2) reduction(+:sum)
+ for (i = 0; i < 10; i++)
+ for (j = 0; j < 10; j++)
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ #pragma acc loop reduction(+:sum)
+ for (i = 0; i < 10; i++)
+ #pragma acc loop reduction(+:sum)
+ for (j = 0; j < 10; j++)
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ #pragma acc loop reduction(+:sum)
+ for (i = 0; i < 10; i++)
+ #pragma acc loop collapse(2) reduction(+:sum)
+ for (j = 0; j < 10; j++)
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ #pragma acc loop reduction(+:sum)
+ for (i = 0; i < 10; i++)
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(+:sum)
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ #pragma acc loop reduction(+:sum)
+ for (i = 0; i < 10; i++)
+ #pragma acc loop reduction(+:sum)
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(+:sum)
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ #pragma acc loop reduction(+:sum) reduction(-:diff)
+ for (i = 0; i < 10; i++)
+ {
+ #pragma acc loop reduction(+:sum)
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(+:sum)
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ #pragma acc loop reduction(-:diff)
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(-:diff)
+ for (k = 0; k < 10; k++)
+ diff = 1;
+ }
+ }
+}
+
+/* The same tests as above, but using a combined parallel loop construct. */
+
+void acc_parallel_loop (void)
+{
+ int i, j, k, l, sum, diff;
+
+ #pragma acc parallel loop
+ for (int h = 0; h < 10; ++h)
+ {
+ #pragma acc loop reduction(+:sum)
+ for (i = 0; i < 10; i++)
+ for (j = 0; j < 10; j++)
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ #pragma acc loop collapse(2) reduction(+:sum)
+ for (i = 0; i < 10; i++)
+ for (j = 0; j < 10; j++)
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ #pragma acc loop reduction(+:sum)
+ for (i = 0; i < 10; i++)
+ #pragma acc loop reduction(+:sum)
+ for (j = 0; j < 10; j++)
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ #pragma acc loop reduction(+:sum)
+ for (i = 0; i < 10; i++)
+ #pragma acc loop collapse(2) reduction(+:sum)
+ for (j = 0; j < 10; j++)
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ #pragma acc loop reduction(+:sum)
+ for (i = 0; i < 10; i++)
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(+:sum)
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ #pragma acc loop reduction(+:sum)
+ for (i = 0; i < 10; i++)
+ #pragma acc loop reduction(+:sum) // { dg-warning "insufficient partitioning available to parallelize loop" }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(+:sum)
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ #pragma acc loop reduction(+:sum) reduction(-:diff)
+ for (i = 0; i < 10; i++)
+ {
+ #pragma acc loop reduction(+:sum) // { dg-warning "insufficient partitioning available to parallelize loop" }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(+:sum)
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ #pragma acc loop reduction(-:diff) // { dg-warning "insufficient partitioning available to parallelize loop" }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(-:diff)
+ for (k = 0; k < 10; k++)
+ diff = 1;
+ }
+ }
+}
+
+/* The same tests as above, but now the outermost reduction clause is on
+ the parallel region, not the outermost loop. */
+
+void acc_parallel_reduction (void)
+{
+ int i, j, k, sum, diff;
+
+ #pragma acc parallel reduction(+:sum)
+ {
+ for (i = 0; i < 10; i++)
+ for (j = 0; j < 10; j++)
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ for (i = 0; i < 10; i++)
+ #pragma acc loop
+ for (j = 0; j < 10; j++)
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ #pragma acc loop reduction(+:sum)
+ for (i = 0; i < 10; i++)
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(+:sum)
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ for (i = 0; i < 10; i++)
+ for (j = 0; j < 10; j++)
+ #pragma acc loop
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ #pragma acc loop reduction(+:sum)
+ for (i = 0; i < 10; i++)
+ #pragma acc loop reduction(+:sum)
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(+:sum)
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ #pragma acc loop reduction(+:sum) reduction(-:diff)
+ for (i = 0; i < 10; i++)
+ {
+ #pragma acc loop reduction(+:sum)
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(+:sum)
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ #pragma acc loop reduction(-:diff)
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(-:diff)
+ for (k = 0; k < 10; k++)
+ diff = 1;
+ }
+
+ #pragma acc loop reduction(+:sum)
+ for (i = 0; i < 10; i++)
+ {
+ #pragma acc loop reduction(+:sum)
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(+:sum)
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ #pragma acc loop reduction(-:diff)
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(-:diff)
+ for (k = 0; k < 10; k++)
+ diff = 1;
+ }
+
+ #pragma acc loop reduction(+:sum)
+ for (i = 0; i < 10; i++)
+ {
+ #pragma acc loop reduction(+:sum)
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(+:sum)
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ #pragma acc loop
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(-:diff)
+ for (k = 0; k < 10; k++)
+ diff = 1;
+ }
+ }
+}
+
+/* The same tests as above, but using a combined parallel loop construct, and
+ the outermost reduction clause is on that one, not the outermost loop. */
+void acc_parallel_loop_reduction (void)
+{
+ int i, j, k, sum, diff;
+
+ #pragma acc parallel loop reduction(+:sum)
+ for (int h = 0; h < 10; ++h)
+ {
+ for (i = 0; i < 10; i++)
+ for (j = 0; j < 10; j++)
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ for (i = 0; i < 10; i++)
+ #pragma acc loop
+ for (j = 0; j < 10; j++)
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ #pragma acc loop reduction(+:sum)
+ for (i = 0; i < 10; i++)
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(+:sum)
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ for (i = 0; i < 10; i++)
+ for (j = 0; j < 10; j++)
+ #pragma acc loop
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ #pragma acc loop reduction(+:sum)
+ for (i = 0; i < 10; i++)
+ #pragma acc loop reduction(+:sum) // { dg-warning "insufficient partitioning available to parallelize loop" }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(+:sum)
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ #pragma acc loop reduction(+:sum) reduction(-:diff)
+ for (i = 0; i < 10; i++)
+ {
+ #pragma acc loop reduction(+:sum) // { dg-warning "insufficient partitioning available to parallelize loop" }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(+:sum)
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ #pragma acc loop reduction(-:diff) // { dg-warning "insufficient partitioning available to parallelize loop" }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(-:diff)
+ for (k = 0; k < 10; k++)
+ diff = 1;
+ }
+
+ #pragma acc loop reduction(+:sum)
+ for (i = 0; i < 10; i++)
+ {
+ #pragma acc loop reduction(+:sum) // { dg-warning "insufficient partitioning available to parallelize loop" }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(+:sum)
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ #pragma acc loop reduction(-:diff) // { dg-warning "insufficient partitioning available to parallelize loop" }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(-:diff)
+ for (k = 0; k < 10; k++)
+ diff = 1;
+ }
+
+ #pragma acc loop reduction(+:sum)
+ for (i = 0; i < 10; i++)
+ {
+ #pragma acc loop reduction(+:sum) // { dg-warning "insufficient partitioning available to parallelize loop" }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(+:sum)
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ #pragma acc loop // { dg-warning "insufficient partitioning available to parallelize loop" }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(-:diff)
+ for (k = 0; k < 10; k++)
+ diff = 1;
+ }
+ }
+}
+
+/* The same tests as above, but inside a routine construct. */
+#pragma acc routine gang
+void acc_routine (void)
+{
+ int i, j, k, sum, diff;
+
+ {
+ #pragma acc loop reduction(+:sum)
+ for (i = 0; i < 10; i++)
+ for (j = 0; j < 10; j++)
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ #pragma acc loop collapse(2) reduction(+:sum)
+ for (i = 0; i < 10; i++)
+ for (j = 0; j < 10; j++)
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ #pragma acc loop reduction(+:sum)
+ for (i = 0; i < 10; i++)
+ #pragma acc loop reduction(+:sum)
+ for (j = 0; j < 10; j++)
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ #pragma acc loop reduction(+:sum)
+ for (i = 0; i < 10; i++)
+ #pragma acc loop collapse(2) reduction(+:sum)
+ for (j = 0; j < 10; j++)
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ #pragma acc loop reduction(+:sum)
+ for (i = 0; i < 10; i++)
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(+:sum)
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ #pragma acc loop reduction(+:sum)
+ for (i = 0; i < 10; i++)
+ #pragma acc loop reduction(+:sum)
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(+:sum)
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ #pragma acc loop reduction(+:sum) reduction(-:diff)
+ for (i = 0; i < 10; i++)
+ {
+ #pragma acc loop reduction(+:sum)
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(+:sum)
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ #pragma acc loop reduction(-:diff)
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(-:diff)
+ for (k = 0; k < 10; k++)
+ diff = 1;
+ }
+ }
+}
+
+void acc_kernels (void)
+{
+ int i, j, k, sum, diff;
+
+ /* FIXME: These tests are not meaningful yet because reductions in
+ kernels regions are not supported yet. */
+ #pragma acc kernels
+ {
+ #pragma acc loop reduction(+:sum)
+ for (i = 0; i < 10; i++)
+ for (j = 0; j < 10; j++)
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ #pragma acc loop reduction(+:sum)
+ for (i = 0; i < 10; i++)
+ #pragma acc loop reduction(+:sum)
+ for (j = 0; j < 10; j++)
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ #pragma acc loop reduction(+:sum)
+ for (i = 0; i < 10; i++)
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(+:sum)
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ #pragma acc loop reduction(+:sum)
+ for (i = 0; i < 10; i++)
+ #pragma acc loop reduction(+:sum)
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(+:sum)
+ for (k = 0; k < 10; k++)
+ sum = 1;
+ }
+}
diff --git a/gcc/testsuite/c-c++-common/goacc/reduction-6.c b/gcc/testsuite/c-c++-common/goacc/reduction-6.c
index 619f82b9d8b..3c10b4dddaf 100644
--- a/gcc/testsuite/c-c++-common/goacc/reduction-6.c
+++ b/gcc/testsuite/c-c++-common/goacc/reduction-6.c
@@ -16,17 +16,6 @@ int foo (int N)
}
}
- #pragma acc parallel
- {
- #pragma acc loop reduction(+:b)
- for (int i = 0; i < N; i++)
- {
- #pragma acc loop
- for (int j = 0; j < N; j++)
- b += 1;
- }
- }
-
#pragma acc parallel
{
#pragma acc loop reduction(+:c)
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-1.c
index 5e82e1d350c..91fe772045d 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-1.c
@@ -15,7 +15,7 @@ main (int argc, char *argv[])
#pragma acc parallel num_gangs(32) num_workers(32) vector_length(32) \
reduction(+:res) copy(res)
{
- #pragma acc loop gang
+ #pragma acc loop gang reduction(+:res)
for (j = 0; j < 32; j++)
{
#pragma acc loop worker reduction(+:res)
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-2.c
index a339f327956..8a7bc724070 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-2.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-2.c
@@ -14,7 +14,7 @@ main (int argc, char *argv[])
#pragma acc parallel num_gangs(32) num_workers(32) vector_length(32) \
reduction(^:res)
{
- #pragma acc loop gang
+ #pragma acc loop gang reduction(^:res)
for (j = 0; j < 32; j++)
{
#pragma acc loop worker vector reduction(^:res)
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-3.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-3.c
index 6369d7fbb33..eba5c6a544e 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-3.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-3.c
@@ -16,7 +16,7 @@ main (int argc, char *argv[])
#pragma acc parallel num_gangs(32) num_workers(32) vector_length(32) \
reduction(+:res) copy(res)
{
- #pragma acc loop gang
+ #pragma acc loop gang reduction(+:res)
for (j = 0; j < 32; j++)
{
#pragma acc loop worker vector reduction(+:res)
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-4.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-4.c
index 140c3226327..12b823f33ab 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-4.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-4.c
@@ -16,7 +16,7 @@ main (int argc, char *argv[])
#pragma acc parallel num_gangs(32) num_workers(32) vector_length(32) \
reduction(+:res) reduction(max:mres) copy(res, mres)
{
- #pragma acc loop gang
+ #pragma acc loop gang reduction(+:res) reduction(max:mres)
for (j = 0; j < 32; j++)
{
#pragma acc loop worker vector reduction(+:res) reduction(max:mres)
--
2.17.1
