Hi!
On behalf of Gergő (who doesn't have write access yet) I've pushed the
attached to openacc-gcc-8-branch.
Grüße
Thomas
From a9e48066198ffb1e7bc2b137167a61a6cb47748c Mon Sep 17 00:00:00 2001
From: =?UTF-8?q?Gerg=C3=B6=20Barany?= <ge...@codesourcery.com>
Date: Thu, 20 Dec 2018 15:07:34 +0100
Subject: [PATCH] Report errors on missing OpenACC reduction clauses in nested
reductions
..., as suggested by OpenACC 2.6, 2.9.11. "reduction clause".
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.
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/ChangeLog.openacc | 9 +
gcc/omp-low.c | 105 ++++
gcc/testsuite/ChangeLog.openacc | 7 +
.../goacc/nested-reductions-fail.c | 492 ++++++++++++++++++
.../c-c++-common/goacc/nested-reductions.c | 420 +++++++++++++++
.../c-c++-common/goacc/reduction-6.c | 11 -
libgomp/ChangeLog.openacc | 12 +
.../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 +-
11 files changed, 1049 insertions(+), 15 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/ChangeLog.openacc b/gcc/ChangeLog.openacc
index e57971ce4cf2..597362505bb3 100644
--- a/gcc/ChangeLog.openacc
+++ b/gcc/ChangeLog.openacc
@@ -1,3 +1,12 @@
+2018-12-20 Gergö Barany <ge...@codesourcery.com>
+
+ * 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.
+
2018-12-20 Julian Brown <jul...@codesourcery.com>
Maciej W. Rozycki <ma...@codesourcery.com>
diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index 967916521001..6b7b23e6909b 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -129,6 +129,12 @@ struct omp_context
/* Hash map of dynamic arrays in this context. */
hash_map<tree_operand_hash, tree> *dynamic_arrays;
+
+ /* 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;
};
static splay_tree all_contexts;
@@ -1040,6 +1046,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
{
@@ -1053,6 +1061,8 @@ new_omp_context (gimple *stmt, omp_context *outer_ctx)
ctx->cb.eh_lp_nr = 0;
ctx->cb.transform_call_graph_edges = CB_CGE_MOVE;
ctx->depth = 1;
+ ctx->local_reduction_clauses = NULL;
+ ctx->outer_reduction_clauses = NULL;
}
ctx->cb.decl_map = new hash_map<tree, tree>;
@@ -1276,6 +1286,9 @@ scan_sharing_clauses (tree clauses, omp_context *ctx,
goto do_private;
case OMP_CLAUSE_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 (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_REDUCTION
&& TREE_CODE (decl) == MEM_REF)
@@ -2458,6 +2471,98 @@ 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/ChangeLog.openacc b/gcc/testsuite/ChangeLog.openacc
index b1758eca0b21..4af31e5e1060 100644
--- a/gcc/testsuite/ChangeLog.openacc
+++ b/gcc/testsuite/ChangeLog.openacc
@@ -1,3 +1,10 @@
+2018-12-20 Gergö Barany <ge...@codesourcery.com>
+ Thomas Schwinge <tho...@codesourcery.com>
+
+ * 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.
+
2018-12-20 Maciej W. Rozycki <ma...@codesourcery.com>
* c-c++-common/goacc/serial-dims.c: New test.
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 000000000000..a642dd038709
--- /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 000000000000..bff66527df58
--- /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) // { dg-bogus "region is gang partitioned but does not contain gang partitioned code" "TODO" { xfail *-*-* } }
+{
+ 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) // { dg-bogus "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } }
+ 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-bogus "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } }
+ 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-bogus "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } }
+ 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 619f82b9d8b3..3c10b4dddaf2 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/ChangeLog.openacc b/libgomp/ChangeLog.openacc
index b16008a77d19..b48453b3066a 100644
--- a/libgomp/ChangeLog.openacc
+++ b/libgomp/ChangeLog.openacc
@@ -1,3 +1,15 @@
+2018-12-20 Gergö Barany <ge...@codesourcery.com>
+ Thomas Schwinge <tho...@codesourcery.com>
+
+ * 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.
+
2018-12-20 Julian Brown <jul...@codesourcery.com>
Maciej W. Rozycki <ma...@codesourcery.com>
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 5e82e1d350cc..91fe772045d9 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 a339f327956e..8a7bc7240708 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 6369d7fbb330..eba5c6a544e8 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 140c32263278..12b823f33ab1 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
