From: Gergö Barany Date: Thu, 20 Dec 2018 14:07:34 +0000 (+0100) Subject: Report errors on missing OpenACC reduction clauses in nested reductions X-Git-Url: http://git.ipfire.org/cgi-bin/gitweb.cgi?a=commitdiff_plain;h=edb782d825ee342289cafbd2f904b26b14a59011;p=thirdparty%2Fgcc.git 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. (cherry picked from openacc-gcc-9-branch commit ffa02cac81b32e526c1b901bc9dc8a0123fcb912) --- diff --git a/gcc/ChangeLog.omp b/gcc/ChangeLog.omp index 10138ae1b000..797a62e1af29 100644 --- a/gcc/ChangeLog.omp +++ b/gcc/ChangeLog.omp @@ -1,3 +1,12 @@ +2018-12-20 Gergö Barany + + * 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. + 2019-01-31 Kwok Cheung Yeung * omp-low.c (lower_omp_target): For use_device clauses, generate diff --git a/gcc/omp-low.c b/gcc/omp-low.c index 1fd29e88ef93..c50b8ff6eec8 100644 --- a/gcc/omp-low.c +++ b/gcc/omp-low.c @@ -138,6 +138,12 @@ struct omp_context /* Hash map of dynamic arrays in this context. */ hash_map *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; + /* The number of levels of OpenACC partitioning invoked in this context. */ unsigned oacc_partitioning_levels; @@ -1001,6 +1007,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 { @@ -1016,6 +1024,8 @@ new_omp_context (gimple *stmt, omp_context *outer_ctx) ctx->cb.adjust_array_error_bounds = true; ctx->cb.dont_remap_vla_if_no_change = true; ctx->depth = 1; + ctx->local_reduction_clauses = NULL; + ctx->outer_reduction_clauses = NULL; } ctx->cb.decl_map = new hash_map; @@ -1236,6 +1246,9 @@ scan_sharing_clauses (tree clauses, omp_context *ctx, case OMP_CLAUSE_REDUCTION: case OMP_CLAUSE_IN_REDUCTION: + if (is_oacc_parallel (ctx) || is_oacc_kernels (ctx)) + ctx->local_reduction_clauses + = tree_cons (NULL, c, ctx->local_reduction_clauses); decl = OMP_CLAUSE_DECL (c); if (TREE_CODE (decl) == MEM_REF) { @@ -2534,6 +2547,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.omp b/gcc/testsuite/ChangeLog.omp index 75aa8585cce6..10b52bee15d0 100644 --- a/gcc/testsuite/ChangeLog.omp +++ b/gcc/testsuite/ChangeLog.omp @@ -1,3 +1,9 @@ +2018-12-20 Gergö Barany + + * 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-10-02 Cesar Philippidis * gfortran.dg/goacc/pr72741.f90: Update test for current diagnostics. 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.omp b/libgomp/ChangeLog.omp index 57e959aa99af..7f1642187843 100644 --- a/libgomp/ChangeLog.omp +++ b/libgomp/ChangeLog.omp @@ -1,3 +1,14 @@ +2018-12-20 Gergö Barany + + * 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. + 2019-01-30 Kwok Cheung Yeung * testsuite/libgomp.oacc-fortran/optional-cache.f95 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)