| Message ID | yxfp7eg4qnla.fsf@hertz.schwinge.homeip.net |
|---|---|
| State | New |
| Headers | show |
| Series | [og8] Report errors on missing OpenACC reduction clauses in nested reductions | expand |
Hi Frederik! As you've been the last one to work on this code and get it into GCC master branch (commit 5d183d1740d8d8b84991f186ce4d992ee799536f "Warn about inconsistent OpenACC nested reduction clauses"), you get to look into the following issue. ..., and yes, I'm aware that you have a follow-up patch which is reworking this code, and which still needs to be reviewed -- <http://mid.mail-archive.com/7ff53a86-d2ec-21ea-d0f4-4e0b13986771@codesourcery.com> -- but let's please first have a look at the current code in GCC master branch, and address what might be a bug. On 2018-12-20T15:28:33+0100, I wrote: > On behalf of Gergő (who doesn't have write access yet) I've pushed the > attached to openacc-gcc-8-branch. > From a9e48066198ffb1e7bc2b137167a61a6cb47748c Mon Sep 17 00:00:00 2001 > From: =?UTF-8?q?Gerg=C3=B6=20Barany?= <gergo@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". > --- 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; Via <https://gcc.gnu.org/PR94629> "10 issues located by the PVS-studio static analyzer" (so please reference that one on any patch submission), on <https://habr.com/en/company/pvs-studio/blog/497640/> in "Fragment N3, Assigning a variable to itself", we find this latter assignment qualified as "very strange to assign a variable to itself". Probably that should've been 'outer_ctx' instead of 'ctx'? But: why/how then does the current algorith still work despite this error? (I haven't looked at the code very much.) Grüße Thomas > } > 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 <gergo@codesourcery.com> > + Thomas Schwinge <thomas@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 <macro@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 <gergo@codesourcery.com> > + Thomas Schwinge <thomas@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 <julian@codesourcery.com> > Maciej W. Rozycki <macro@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) ----------------- Mentor Graphics (Deutschland) GmbH, Arnulfstraße 201, 80634 München / Germany Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Alexander Walter
Thomas Schwinge <thomas@codesourcery.com> writes: Hi Thomas, > Via <https://gcc.gnu.org/PR94629> "10 issues located by the PVS-studio > static analyzer" (so please reference that one on any patch submission), > on <https://habr.com/en/company/pvs-studio/blog/497640/> in "Fragment N3, > Assigning a variable to itself", we find this latter assignment qualified > as "very strange to assign a variable to itself". > > Probably that should've been 'outer_ctx' instead of 'ctx'? I agree that the original intention must have been to assign the outer_ctx's "outer_reduction_clauses" to the corresponding field of the inner "ctx". This would make sense, semantically. But this field is meant to be used by the function "scan_omp_for" only and ... > then does the current algorith still work despite this error? ... this function never requires the struct field to be intialized in that way. Before the field is used, it always copies the clauses from the outer context's outer_reduction_clauses to ctx->outer_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); Hence I found it preferrable to remove the assignment to the "outer_reduction_clauses" field and the "local_reduction_clauses" field from "new_omp_context" completely. (The fields are still zero intialized by the allocation of the struct which uses XCNEW.) That way the whole logic regarding the fields is now contained in "scan_omp_for". I have executed "make check" (on x86_64-linux-gnu) to verify that the change causes no regressions. Ok to push the commit to master? Best regards, Frederik ----------------- Mentor Graphics (Deutschland) GmbH, Arnulfstraße 201, 80634 München / Germany Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Alexander Walter
Hi Frederik! On 2020-04-21T14:13:33+0200, Frederik Harwath <frederik@codesourcery.com> wrote: > Thomas Schwinge <thomas@codesourcery.com> writes: >> Via <https://gcc.gnu.org/PR94629> "10 issues located by the PVS-studio >> static analyzer" (so please reference that one on any patch submission), >> on <https://habr.com/en/company/pvs-studio/blog/497640/> in "Fragment N3, >> Assigning a variable to itself", we find this latter assignment qualified >> as "very strange to assign a variable to itself". > [...] the original intention must have been [...] >> then does the current algorith still work despite this error? > > [...] Thanks for this analysis! > Hence I found it preferrable to remove the assignment to the > "outer_reduction_clauses" field and the "local_reduction_clauses" field > from "new_omp_context" completely. (The fields are still zero intialized > by the allocation of the struct which uses XCNEW.) That way the whole > logic regarding the fields is now contained in "scan_omp_for". > > I have executed "make check" (on x86_64-linux-gnu) to verify that the > change causes no regressions. Ok to push the commit to master? ACK, thanks! To record the review effort, please include "Reviewed-by: Thomas Schwinge <thomas@codesourcery.com>" in the commit log, see <https://gcc.gnu.org/wiki/Reviewed-by>. Grüße Thomas > From 2d60b374a44b212ff97c8b1fd6f8c39e478dc70f Mon Sep 17 00:00:00 2001 > From: Frederik Harwath <frederik@codesourcery.com> > Date: Tue, 21 Apr 2020 12:36:14 +0200 > Subject: [PATCH] Remove fishy self-assignment in omp-low.c [PR94629] > > The PR noticed that omp-low.c contains a self-assignment in the > function new_omp_context: > > if (outer_ctx) { > ... > ctx->outer_reduction_clauses = ctx->outer_reduction_clauses; > > This is obviously useless. The original intention might have been > to copy the field from the outer_ctx to ctx. Since this is done > (properly) in the only function where this field is actually used > (in function scan_omp_for) and the field is being initialized to zero > during the struct allocation, there is no need to attempt to do > anything to this field in new_omp_context. Thus this commit > removes any assignment to the field from new_omp_context. > > 2020-04-21 Frederik Harwath <frederik@codesourcery.com> > > PR other/94629 > * gcc/omp-low.c (new_omp_context): Remove assignments to > ctx->outer_reduction_clauses and ctx->local_reduction_clauses. > --- > gcc/omp-low.c | 14 ++++++++------ > 1 file changed, 8 insertions(+), 6 deletions(-) > > diff --git a/gcc/omp-low.c b/gcc/omp-low.c > index 67565d61400..88f23e60d34 100644 > --- a/gcc/omp-low.c > +++ b/gcc/omp-low.c > @@ -128,10 +128,16 @@ 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. */ > + /* A tree_list of the reduction clauses in this context. This is > + only used for checking the consistency of OpenACC reduction > + clauses in scan_omp_for and is not guaranteed to contain a valid > + value outside of this function. */ > tree local_reduction_clauses; > > - /* A tree_list of the reduction clauses in outer contexts. */ > + /* A tree_list of the reduction clauses in outer contexts. This is > + only used for checking the consistency of OpenACC reduction > + clauses in scan_omp_for and is not guaranteed to contain a valid > + value outside of this function. */ > tree outer_reduction_clauses; > > /* Nesting depth of this context. Used to beautify error messages re > @@ -931,8 +937,6 @@ new_omp_context (gimple *stmt, omp_context *outer_ctx) > ctx->outer = outer_ctx; > ctx->cb = outer_ctx->cb; > ctx->cb.block = NULL; > - ctx->local_reduction_clauses = NULL; > - ctx->outer_reduction_clauses = ctx->outer_reduction_clauses; > ctx->depth = outer_ctx->depth + 1; > } > else > @@ -948,8 +952,6 @@ new_omp_context (gimple *stmt, omp_context *outer_ctx) > ctx->cb.transform_call_graph_edges = CB_CGE_MOVE; > ctx->cb.adjust_array_error_bounds = true; > ctx->cb.dont_remap_vla_if_no_change = true; > - ctx->local_reduction_clauses = NULL; > - ctx->outer_reduction_clauses = NULL; > ctx->depth = 1; > } > ----------------- Mentor Graphics (Deutschland) GmbH, Arnulfstraße 201, 80634 München / Germany Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Alexander Walter
Hi! On 2018-12-20T15:28:33+0100, I wrote: > On behalf of Gergő (who doesn't have write access yet) I've pushed the > attached to openacc-gcc-8-branch. (Which then eventually got into master branch via Frederik.) > --- /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; > +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; For ongoing maintenance, I find it easier if such repetitive testing (here: to cover all different compute constructs and 'routine') is split into separate files: easy to diff, etc. > --- /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. */ Likewise. I've pushed "[OpenACC] Split up testcases for inconsistent nested 'reduction' clauses checking" to master branch in commit fedf3e94efe774b8c0539d344130a7b25f50a881, and backported to releases/gcc-10 branch in commit eeeb6833d2c6bcc0e675928f17a75efb41eeaf13, see attached. Grüße Thomas ----------------- Mentor Graphics (Deutschland) GmbH, Arnulfstraße 201, 80634 München / Germany Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Alexander Walter
Hi! On 2020-11-03T09:39:54+0100, I wrote: > On 2018-12-20T15:28:33+0100, I wrote: >> On behalf of Gergő (who doesn't have write access yet) I've pushed the >> attached to openacc-gcc-8-branch. > > (Which then eventually got into master branch via Frederik.) >> --- /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; > >> +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; Getting these diagnostics consistent is easy enough, however. I've pushed "[OpenACC] Enable inconsistent nested 'reduction' clauses checking for OpenACC 'kernels'" to master branch in commit 64dc14b1a764bd3059170431c9b43c6192dbd48f, and backported to releases/gcc-10 branch in commit 217fb4d4e59e7d6e03a3704f80f401e2a641dbe5, see attached. Grüße Thomas ----------------- Mentor Graphics (Deutschland) GmbH, Arnulfstraße 201, 80634 München / Germany Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Alexander Walter
From a9e48066198ffb1e7bc2b137167a61a6cb47748c Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Gerg=C3=B6=20Barany?= <gergo@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 <gergo@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 <julian@codesourcery.com> Maciej W. Rozycki <macro@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 <gergo@codesourcery.com> + Thomas Schwinge <thomas@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 <macro@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 <gergo@codesourcery.com> + Thomas Schwinge <thomas@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 <julian@codesourcery.com> Maciej W. Rozycki <macro@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