diff --git a/gcc/ChangeLog b/gcc/ChangeLog index 7fee0f37e9b..38160dd631e 100644 --- a/gcc/ChangeLog +++ b/gcc/ChangeLog @@ -1,3 +1,13 @@ +2019-11-06 Gergö Barany + Frederik Harwath + Thomas Schwinge + + * 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-11-06 Jakub Jelinek PR inline-asm/92352 diff --git a/gcc/omp-low.c b/gcc/omp-low.c index 122f4278881..fa76ceba33c 100644 --- a/gcc/omp-low.c +++ b/gcc/omp-low.c @@ -128,6 +128,12 @@ struct omp_context corresponding tracking loop iteration variables. */ hash_map *lastprivate_conditional_map; + /* A tree_list of the reduction clauses in this context. */ + tree local_reduction_clauses; + + /* A tree_list of the reduction clauses in outer contexts. */ + tree outer_reduction_clauses; + /* Nesting depth of this context. Used to beautify error messages re invalid gotos. The outermost ctx is depth 1, with depth 0 being reserved for the main body of the function. */ @@ -910,6 +916,8 @@ 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 @@ -925,6 +933,8 @@ 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; } @@ -1139,6 +1149,11 @@ 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); + /* FALLTHRU */ + case OMP_CLAUSE_IN_REDUCTION: decl = OMP_CLAUSE_DECL (c); if (TREE_CODE (decl) == MEM_REF) @@ -2423,6 +2438,88 @@ 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) + { + warning_at (gimple_location (stmt), 0, + "conflicting reduction operations for %qE", + local_var); + inform (OMP_CLAUSE_LOCATION (outer_clause), + "location of the previous reduction for %qE", + outer_var); + } + 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) + warning_at (gimple_location (curr_loop->stmt), 0, + "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 b/gcc/testsuite/ChangeLog index a0f52ec2c5c..64568e5661a 100644 --- a/gcc/testsuite/ChangeLog +++ b/gcc/testsuite/ChangeLog @@ -1,3 +1,12 @@ +2019-11-06 Gergö Barany + Frederik Harwath + Thomas Schwinge + + * c-c++-common/goacc/nested-reductions-warn.c: New test. + * c-c++-common/goacc/nested-reductions.c: New test. + * gfortran.dg/goacc/nested-reductions-warn.f90: New test. + * gfortran.dg/goacc/nested-reductions.f90: New test. + 2019-11-06 Jakub Jelinek PR inline-asm/92352 diff --git a/gcc/testsuite/c-c++-common/goacc/nested-reductions-warn.c b/gcc/testsuite/c-c++-common/goacc/nested-reductions-warn.c new file mode 100644 index 00000000000..e2af66e4fa3 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/nested-reductions-warn.c @@ -0,0 +1,525 @@ +/* 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-warning "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-warning "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-warning "nested loop in reduction needs reduction clause for .sum." } + for (j = 0; j < 10; j++) + #pragma acc loop // { dg-warning "nested loop in reduction needs reduction clause for .sum." } + // { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 } + 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-warning "conflicting reduction operations for .sum." } + for (j = 0; j < 10; j++) + #pragma acc loop reduction(+:sum) // { dg-warning "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-warning "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-warning "conflicting reduction operations for .sum." } + for (j = 0; j < 10; j++) + #pragma acc loop // { dg-warning "nested loop in reduction needs reduction clause for .sum." } + // { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 } + for (k = 0; k < 10; k++) + #pragma acc loop reduction(*:sum) // { dg-warning "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-warning "conflicting reduction operations for .sum." } + for (j = 0; j < 10; j++) + #pragma acc loop reduction(+:sum) // { dg-warning "conflicting reduction operations for .sum." } + // { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 } + for (k = 0; k < 10; k++) + #pragma acc loop reduction(*:sum) // { dg-warning "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-warning "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-warning "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-warning "nested loop in reduction needs reduction clause for .sum." } + // { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 } + 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-warning "nested loop in reduction needs reduction clause for .sum." } + // { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 } + 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-warning "nested loop in reduction needs reduction clause for .sum." } + // { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 } + for (j = 0; j < 10; j++) + #pragma acc loop // { dg-warning "nested loop in reduction needs reduction clause for .sum." } + // { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 } + 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-warning "conflicting reduction operations for .sum." } + // { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 } + for (j = 0; j < 10; j++) + #pragma acc loop reduction(+:sum) // { dg-warning "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-warning "conflicting reduction operations for .sum." } + // { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 } + 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 "conflicting reduction operations for .sum." } + // { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 } + for (j = 0; j < 10; j++) + #pragma acc loop // { dg-warning "nested loop in reduction needs reduction clause for .sum." } + // { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 } + for (k = 0; k < 10; k++) + #pragma acc loop reduction(*:sum) // { dg-warning "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-warning "conflicting reduction operations for .sum." } + // { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 } + for (j = 0; j < 10; j++) + #pragma acc loop reduction(+:sum) // { dg-warning "conflicting reduction operations for .sum." } + // { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 } + for (k = 0; k < 10; k++) + #pragma acc loop reduction(*:sum) // { dg-warning "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-warning "nested loop in reduction needs reduction clause for .sum." } + // { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 } + 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-warning "nested loop in reduction needs reduction clause for .diff." } + // { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 } + 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-warning "nested loop in reduction needs reduction clause for .sum." } + for (i = 0; i < 10; i++) + #pragma acc loop // { dg-warning "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-warning "nested loop in reduction needs reduction clause for .sum." } + for (i = 0; i < 10; i++) + #pragma acc loop collapse(2) // { dg-warning "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-warning "nested loop in reduction needs reduction clause for .sum." } + for (i = 0; i < 10; i++) + #pragma acc loop // { dg-warning "nested loop in reduction needs reduction clause for .sum." } + for (j = 0; j < 10; j++) + #pragma acc loop // { dg-warning "nested loop in reduction needs reduction clause for .sum." } + // { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 } + for (k = 0; k < 10; k++) + #pragma acc loop reduction(+:sum) + for (l = 0; l < 10; l++) + sum = 1; + + #pragma acc loop // { dg-warning "nested loop in reduction needs reduction clause for .sum." } + for (i = 0; i < 10; i++) + #pragma acc loop reduction(-:sum) // { dg-warning "conflicting reduction operations for .sum." } + for (j = 0; j < 10; j++) + #pragma acc loop reduction(+:sum) // { dg-warning "conflicting reduction operations for .sum." } + for (k = 0; k < 10; k++) + sum = 1; + + #pragma acc loop // { dg-warning "nested loop in reduction needs reduction clause for .sum." } + for (i = 0; i < 10; i++) + #pragma acc loop reduction(-:sum) // { dg-warning "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(max:sum) // { dg-warning "conflicting reduction operations for .sum." } + for (i = 0; i < 10; i++) + #pragma acc loop reduction(-:sum) // { dg-warning "conflicting reduction operations for .sum." } + for (j = 0; j < 10; j++) + #pragma acc loop // { dg-warning "nested loop in reduction needs reduction clause for .sum." } + // { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 } + for (k = 0; k < 10; k++) + #pragma acc loop reduction(*:sum) // { dg-warning "conflicting reduction operations for .sum." } + for (l = 0; l < 10; l++) + sum = 1; + + #pragma acc loop reduction(max:sum) // { dg-warning "conflicting reduction operations for .sum." } + for (i = 0; i < 10; i++) + #pragma acc loop reduction(-:sum) // { dg-warning "conflicting reduction operations for .sum." } + for (j = 0; j < 10; j++) + #pragma acc loop reduction(+:sum) // { dg-warning "conflicting reduction operations for .sum." }) + // { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 } + for (k = 0; k < 10; k++) + #pragma acc loop reduction(*:sum) // { dg-warning "conflicting reduction operations for .sum." } + for (l = 0; l < 10; l++) + sum = 1; + + #pragma acc loop reduction(-:diff) // { dg-warning "nested loop in reduction needs reduction clause for .sum." } + for (i = 0; i < 10; i++) + { + #pragma acc loop reduction(-:diff) // { dg-warning "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-warning "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-warning "nested loop in reduction needs reduction clause for .sum." } + for (i = 0; i < 10; i++) + #pragma acc loop // { dg-warning "nested loop in reduction needs reduction clause for .sum." } + // { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 } + for (j = 0; j < 10; j++) + #pragma acc loop reduction(+:sum) + for (k = 0; k < 10; k++) + sum = 1; + + #pragma acc loop // { dg-warning "nested loop in reduction needs reduction clause for .sum." } + for (i = 0; i < 10; i++) + #pragma acc loop collapse(2) // { dg-warning "nested loop in reduction needs reduction clause for .sum." } + // { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 } + 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-warning "nested loop in reduction needs reduction clause for .sum." } + for (i = 0; i < 10; i++) + #pragma acc loop // { dg-warning "nested loop in reduction needs reduction clause for .sum." } + // { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 } + for (j = 0; j < 10; j++) + #pragma acc loop // { dg-warning "nested loop in reduction needs reduction clause for .sum." } + // { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 } + for (k = 0; k < 10; k++) + #pragma acc loop reduction(+:sum) + for (l = 0; l < 10; l++) + sum = 1; + + #pragma acc loop // { dg-warning "nested loop in reduction needs reduction clause for .sum." } + for (i = 0; i < 10; i++) + #pragma acc loop reduction(-:sum) // { dg-warning "conflicting reduction operations for .sum." } + // { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 } + for (j = 0; j < 10; j++) + #pragma acc loop reduction(+:sum) // { dg-warning "conflicting reduction operations for .sum." } + for (k = 0; k < 10; k++) + sum = 1; + + #pragma acc loop // { dg-warning "nested loop in reduction needs reduction clause for .sum." } + for (i = 0; i < 10; i++) + #pragma acc loop reduction(-:sum) // { dg-warning "conflicting reduction operations for .sum." } + // { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 } + for (j = 0; j < 10; j++) + #pragma acc loop reduction(-:sum) + for (k = 0; k < 10; k++) + sum = 1; + + #pragma acc loop reduction(max:sum) // { dg-warning "conflicting reduction operations for .sum." } + for (i = 0; i < 10; i++) + #pragma acc loop reduction(-:sum) // { dg-warning "conflicting reduction operations for .sum." } + // { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 } + for (j = 0; j < 10; j++) + #pragma acc loop // { dg-warning "nested loop in reduction needs reduction clause for .sum." } + // { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 } + for (k = 0; k < 10; k++) + #pragma acc loop reduction(*:sum) // { dg-warning "conflicting reduction operations for .sum." } + for (l = 0; l < 10; l++) + sum = 1; + + #pragma acc loop reduction(max:sum) // { dg-warning "conflicting reduction operations for .sum." } + for (i = 0; i < 10; i++) + #pragma acc loop reduction(-:sum) // { dg-warning "conflicting reduction operations for .sum." } + // { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 } + for (j = 0; j < 10; j++) + #pragma acc loop reduction(+:sum) // { dg-warning "conflicting reduction operations for .sum." }) + // { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 } + for (k = 0; k < 10; k++) + #pragma acc loop reduction(*:sum) // { dg-warning "conflicting reduction operations for .sum." } + for (l = 0; l < 10; l++) + sum = 1; + + #pragma acc loop reduction(-:diff) // { dg-warning "nested loop in reduction needs reduction clause for .sum." } + for (i = 0; i < 10; i++) + { + #pragma acc loop reduction(-:diff) // { dg-warning "nested loop in reduction needs reduction clause for .sum." } + // { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 } + 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-warning "nested loop in reduction needs reduction clause for .diff." } + // { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 } + 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-warning "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-warning "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-warning "nested loop in reduction needs reduction clause for .sum." } + for (j = 0; j < 10; j++) + #pragma acc loop // { dg-warning "nested loop in reduction needs reduction clause for .sum." } + // { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 } + 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-warning "conflicting reduction operations for .sum." } + for (j = 0; j < 10; j++) + #pragma acc loop reduction(+:sum) // { dg-warning "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-warning "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-warning "conflicting reduction operations for .sum." } + for (j = 0; j < 10; j++) + #pragma acc loop // { dg-warning "nested loop in reduction needs reduction clause for .sum." } + // { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 } + for (k = 0; k < 10; k++) + #pragma acc loop reduction(*:sum) // { dg-warning "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-warning "conflicting reduction operations for .sum." } + for (j = 0; j < 10; j++) + #pragma acc loop reduction(+:sum) // { dg-warning "conflicting reduction operations for .sum." }) + // { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 } + for (k = 0; k < 10; k++) + #pragma acc loop reduction(*:sum) // { dg-warning "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-warning "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-warning "nested loop in reduction needs reduction clause for .diff." } + for (j = 0; j < 10; j++) + #pragma acc loop reduction(-:diff) + for (k = 0; k < 10; k++) + diff = 1; + } + } +} + +void acc_kernels (void) +{ + int i, j, k, sum, diff; + + /* FIXME: No diagnostics are produced for these loops because reductions + in kernels regions are not supported yet. */ + #pragma acc kernels + { + #pragma acc loop reduction(+:sum) + for (i = 0; i < 10; i++) + for (j = 0; j < 10; j++) + for (k = 0; k < 10; k++) + sum = 1; + + #pragma acc loop reduction(+:sum) + for (i = 0; i < 10; i++) + #pragma acc loop + for (j = 0; j < 10; j++) + for (k = 0; k < 10; k++) + sum = 1; + + #pragma acc loop reduction(+:sum) + for (i = 0; i < 10; i++) + #pragma acc loop reduction(-:diff) + for (j = 0; j < 10; j++) + #pragma acc loop + for (k = 0; k < 10; k++) + sum = 1; + + #pragma acc loop reduction(+:sum) + for (i = 0; i < 10; i++) + #pragma acc loop + for (j = 0; j < 10; j++) + #pragma acc loop reduction(+:sum) + for (k = 0; k < 10; k++) + sum = 1; + + #pragma acc loop reduction(+:sum) + for (i = 0; i < 10; i++) + #pragma acc loop reduction(-:sum) + for (j = 0; j < 10; j++) + #pragma acc loop reduction(+:sum) + for (k = 0; k < 10; k++) + sum = 1; + } +} diff --git a/gcc/testsuite/c-c++-common/goacc/nested-reductions.c b/gcc/testsuite/c-c++-common/goacc/nested-reductions.c new file mode 100644 index 00000000000..15385c4da99 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/nested-reductions.c @@ -0,0 +1,420 @@ +/* Test cases of nested reduction loops that should compile cleanly. */ + +void acc_parallel (void) +{ + int i, j, k, sum, diff; + + #pragma acc parallel + { + #pragma acc loop reduction(+:sum) + for (i = 0; i < 10; i++) + for (j = 0; j < 10; j++) + for (k = 0; k < 10; k++) + sum = 1; + + #pragma acc loop collapse(2) reduction(+:sum) + for (i = 0; i < 10; i++) + for (j = 0; j < 10; j++) + for (k = 0; k < 10; k++) + sum = 1; + + #pragma acc loop reduction(+:sum) + for (i = 0; i < 10; i++) + #pragma acc loop reduction(+:sum) + for (j = 0; j < 10; j++) + for (k = 0; k < 10; k++) + sum = 1; + + #pragma acc loop reduction(+:sum) + for (i = 0; i < 10; i++) + #pragma acc loop collapse(2) reduction(+:sum) + for (j = 0; j < 10; j++) + for (k = 0; k < 10; k++) + sum = 1; + + #pragma acc loop reduction(+:sum) + for (i = 0; i < 10; i++) + for (j = 0; j < 10; j++) + #pragma acc loop reduction(+:sum) + for (k = 0; k < 10; k++) + sum = 1; + + #pragma acc loop reduction(+:sum) + for (i = 0; i < 10; i++) + #pragma acc loop reduction(+:sum) + for (j = 0; j < 10; j++) + #pragma acc loop reduction(+:sum) + for (k = 0; k < 10; k++) + sum = 1; + + #pragma acc loop reduction(+:sum) reduction(-:diff) + for (i = 0; i < 10; i++) + { + #pragma acc loop reduction(+:sum) + for (j = 0; j < 10; j++) + #pragma acc loop reduction(+:sum) + for (k = 0; k < 10; k++) + sum = 1; + + #pragma acc loop reduction(-:diff) + for (j = 0; j < 10; j++) + #pragma acc loop reduction(-:diff) + for (k = 0; k < 10; k++) + diff = 1; + } + } +} + +/* The same tests as above, but using a combined parallel loop construct. */ + +void acc_parallel_loop (void) +{ + int i, j, k, l, sum, diff; + + #pragma acc parallel loop + for (int h = 0; h < 10; ++h) + { + #pragma acc loop reduction(+:sum) + for (i = 0; i < 10; i++) + for (j = 0; j < 10; j++) + for (k = 0; k < 10; k++) + sum = 1; + + #pragma acc loop collapse(2) reduction(+:sum) + for (i = 0; i < 10; i++) + for (j = 0; j < 10; j++) + for (k = 0; k < 10; k++) + sum = 1; + + #pragma acc loop reduction(+:sum) + for (i = 0; i < 10; i++) + #pragma acc loop reduction(+:sum) + for (j = 0; j < 10; j++) + for (k = 0; k < 10; k++) + sum = 1; + + #pragma acc loop reduction(+:sum) + for (i = 0; i < 10; i++) + #pragma acc loop collapse(2) reduction(+:sum) + for (j = 0; j < 10; j++) + for (k = 0; k < 10; k++) + sum = 1; + + #pragma acc loop reduction(+:sum) + for (i = 0; i < 10; i++) + for (j = 0; j < 10; j++) + #pragma acc loop reduction(+:sum) + for (k = 0; k < 10; k++) + sum = 1; + + #pragma acc loop reduction(+:sum) + for (i = 0; i < 10; i++) + #pragma acc loop reduction(+:sum) // { dg-warning "insufficient partitioning available to parallelize loop" } + for (j = 0; j < 10; j++) + #pragma acc loop reduction(+:sum) + for (k = 0; k < 10; k++) + sum = 1; + + #pragma acc loop reduction(+:sum) reduction(-:diff) + for (i = 0; i < 10; i++) + { + #pragma acc loop reduction(+:sum) // { dg-warning "insufficient partitioning available to parallelize loop" } + for (j = 0; j < 10; j++) + #pragma acc loop reduction(+:sum) + for (k = 0; k < 10; k++) + sum = 1; + + #pragma acc loop reduction(-:diff) // { dg-warning "insufficient partitioning available to parallelize loop" } + for (j = 0; j < 10; j++) + #pragma acc loop reduction(-:diff) + for (k = 0; k < 10; k++) + diff = 1; + } + } +} + +/* The same tests as above, but now the outermost reduction clause is on + the parallel region, not the outermost loop. */ + +void acc_parallel_reduction (void) +{ + int i, j, k, sum, diff; + + #pragma acc parallel reduction(+:sum) + { + for (i = 0; i < 10; i++) + for (j = 0; j < 10; j++) + for (k = 0; k < 10; k++) + sum = 1; + + for (i = 0; i < 10; i++) + #pragma acc loop + for (j = 0; j < 10; j++) + for (k = 0; k < 10; k++) + sum = 1; + + #pragma acc loop reduction(+:sum) + for (i = 0; i < 10; i++) + for (j = 0; j < 10; j++) + #pragma acc loop reduction(+:sum) + for (k = 0; k < 10; k++) + sum = 1; + + for (i = 0; i < 10; i++) + for (j = 0; j < 10; j++) + #pragma acc loop + for (k = 0; k < 10; k++) + sum = 1; + + #pragma acc loop reduction(+:sum) + for (i = 0; i < 10; i++) + #pragma acc loop reduction(+:sum) + for (j = 0; j < 10; j++) + #pragma acc loop reduction(+:sum) + for (k = 0; k < 10; k++) + sum = 1; + + #pragma acc loop reduction(+:sum) reduction(-:diff) + for (i = 0; i < 10; i++) + { + #pragma acc loop reduction(+:sum) + for (j = 0; j < 10; j++) + #pragma acc loop reduction(+:sum) + for (k = 0; k < 10; k++) + sum = 1; + + #pragma acc loop reduction(-:diff) + for (j = 0; j < 10; j++) + #pragma acc loop reduction(-:diff) + for (k = 0; k < 10; k++) + diff = 1; + } + + #pragma acc loop reduction(+:sum) + for (i = 0; i < 10; i++) + { + #pragma acc loop reduction(+:sum) + for (j = 0; j < 10; j++) + #pragma acc loop reduction(+:sum) + for (k = 0; k < 10; k++) + sum = 1; + + #pragma acc loop reduction(-:diff) + for (j = 0; j < 10; j++) + #pragma acc loop reduction(-:diff) + for (k = 0; k < 10; k++) + diff = 1; + } + + #pragma acc loop reduction(+:sum) + for (i = 0; i < 10; i++) + { + #pragma acc loop reduction(+:sum) + for (j = 0; j < 10; j++) + #pragma acc loop reduction(+:sum) + for (k = 0; k < 10; k++) + sum = 1; + + #pragma acc loop + for (j = 0; j < 10; j++) + #pragma acc loop reduction(-:diff) + for (k = 0; k < 10; k++) + diff = 1; + } + } +} + +/* The same tests as above, but using a combined parallel loop construct, and + the outermost reduction clause is on that one, not the outermost loop. */ +void acc_parallel_loop_reduction (void) +{ + int i, j, k, sum, diff; + + #pragma acc parallel loop reduction(+:sum) + for (int h = 0; h < 10; ++h) + { + for (i = 0; i < 10; i++) + for (j = 0; j < 10; j++) + for (k = 0; k < 10; k++) + sum = 1; + + for (i = 0; i < 10; i++) + #pragma acc loop + for (j = 0; j < 10; j++) + for (k = 0; k < 10; k++) + sum = 1; + + #pragma acc loop reduction(+:sum) + for (i = 0; i < 10; i++) + for (j = 0; j < 10; j++) + #pragma acc loop reduction(+:sum) + for (k = 0; k < 10; k++) + sum = 1; + + for (i = 0; i < 10; i++) + for (j = 0; j < 10; j++) + #pragma acc loop + for (k = 0; k < 10; k++) + sum = 1; + + #pragma acc loop reduction(+:sum) + for (i = 0; i < 10; i++) + #pragma acc loop reduction(+:sum) // { dg-warning "insufficient partitioning available to parallelize loop" } + for (j = 0; j < 10; j++) + #pragma acc loop reduction(+:sum) + for (k = 0; k < 10; k++) + sum = 1; + + #pragma acc loop reduction(+:sum) reduction(-:diff) + for (i = 0; i < 10; i++) + { + #pragma acc loop reduction(+:sum) // { dg-warning "insufficient partitioning available to parallelize loop" } + for (j = 0; j < 10; j++) + #pragma acc loop reduction(+:sum) + for (k = 0; k < 10; k++) + sum = 1; + + #pragma acc loop reduction(-:diff) // { dg-warning "insufficient partitioning available to parallelize loop" } + for (j = 0; j < 10; j++) + #pragma acc loop reduction(-:diff) + for (k = 0; k < 10; k++) + diff = 1; + } + + #pragma acc loop reduction(+:sum) + for (i = 0; i < 10; i++) + { + #pragma acc loop reduction(+:sum) // { dg-warning "insufficient partitioning available to parallelize loop" } + for (j = 0; j < 10; j++) + #pragma acc loop reduction(+:sum) + for (k = 0; k < 10; k++) + sum = 1; + + #pragma acc loop reduction(-:diff) // { dg-warning "insufficient partitioning available to parallelize loop" } + for (j = 0; j < 10; j++) + #pragma acc loop reduction(-:diff) + for (k = 0; k < 10; k++) + diff = 1; + } + + #pragma acc loop reduction(+:sum) + for (i = 0; i < 10; i++) + { + #pragma acc loop reduction(+:sum) // { dg-warning "insufficient partitioning available to parallelize loop" } + for (j = 0; j < 10; j++) + #pragma acc loop reduction(+:sum) + for (k = 0; k < 10; k++) + sum = 1; + + #pragma acc loop // { dg-warning "insufficient partitioning available to parallelize loop" } + for (j = 0; j < 10; j++) + #pragma acc loop reduction(-:diff) + for (k = 0; k < 10; k++) + diff = 1; + } + } +} + +/* The same tests as above, but inside a routine construct. */ +#pragma acc routine gang +void acc_routine (void) +{ + int i, j, k, sum, diff; + + { + #pragma acc loop reduction(+:sum) + for (i = 0; i < 10; i++) + for (j = 0; j < 10; j++) + for (k = 0; k < 10; k++) + sum = 1; + + #pragma acc loop collapse(2) reduction(+:sum) + for (i = 0; i < 10; i++) + for (j = 0; j < 10; j++) + for (k = 0; k < 10; k++) + sum = 1; + + #pragma acc loop reduction(+:sum) + for (i = 0; i < 10; i++) + #pragma acc loop reduction(+:sum) + for (j = 0; j < 10; j++) + for (k = 0; k < 10; k++) + sum = 1; + + #pragma acc loop reduction(+:sum) + for (i = 0; i < 10; i++) + #pragma acc loop collapse(2) reduction(+:sum) + for (j = 0; j < 10; j++) + for (k = 0; k < 10; k++) + sum = 1; + + #pragma acc loop reduction(+:sum) + for (i = 0; i < 10; i++) + for (j = 0; j < 10; j++) + #pragma acc loop reduction(+:sum) + for (k = 0; k < 10; k++) + sum = 1; + + #pragma acc loop reduction(+:sum) + for (i = 0; i < 10; i++) + #pragma acc loop reduction(+:sum) + for (j = 0; j < 10; j++) + #pragma acc loop reduction(+:sum) + for (k = 0; k < 10; k++) + sum = 1; + + #pragma acc loop reduction(+:sum) reduction(-:diff) + for (i = 0; i < 10; i++) + { + #pragma acc loop reduction(+:sum) + for (j = 0; j < 10; j++) + #pragma acc loop reduction(+:sum) + for (k = 0; k < 10; k++) + sum = 1; + + #pragma acc loop reduction(-:diff) + for (j = 0; j < 10; j++) + #pragma acc loop reduction(-:diff) + for (k = 0; k < 10; k++) + diff = 1; + } + } +} + +void acc_kernels (void) +{ + int i, j, k, sum, diff; + + /* FIXME: These tests are not meaningful yet because reductions in + kernels regions are not supported yet. */ + #pragma acc kernels + { + #pragma acc loop reduction(+:sum) + for (i = 0; i < 10; i++) + for (j = 0; j < 10; j++) + for (k = 0; k < 10; k++) + sum = 1; + + #pragma acc loop reduction(+:sum) + for (i = 0; i < 10; i++) + #pragma acc loop reduction(+:sum) + for (j = 0; j < 10; j++) + for (k = 0; k < 10; k++) + sum = 1; + + #pragma acc loop reduction(+:sum) + for (i = 0; i < 10; i++) + for (j = 0; j < 10; j++) + #pragma acc loop reduction(+:sum) + for (k = 0; k < 10; k++) + sum = 1; + + #pragma acc loop reduction(+:sum) + for (i = 0; i < 10; i++) + #pragma acc loop reduction(+:sum) + for (j = 0; j < 10; j++) + #pragma acc loop reduction(+:sum) + for (k = 0; k < 10; k++) + sum = 1; + } +} diff --git a/gcc/testsuite/gfortran.dg/goacc/nested-reductions-warn.f90 b/gcc/testsuite/gfortran.dg/goacc/nested-reductions-warn.f90 new file mode 100644 index 00000000000..ec36bc9616e --- /dev/null +++ b/gcc/testsuite/gfortran.dg/goacc/nested-reductions-warn.f90 @@ -0,0 +1,674 @@ +! Test erroneous cases of nested reduction loops. + +subroutine acc_parallel () + implicit none (type, external) + integer :: i, j, k, l, sum, diff + + !$acc parallel + !$acc loop reduction(+:sum) + do i = 1, 10 + !$acc loop ! { dg-warning "nested loop in reduction needs reduction clause for .sum." } + do j = 1, 10 + !$acc loop reduction(+:sum) + do k = 1, 10 + sum = 1 + end do + end do + end do + + !$acc loop reduction(+:sum) + do i = 1, 10 + !$acc loop collapse(2) ! { dg-warning "nested loop in reduction needs reduction clause for .sum." } + do j = 1, 10 + do k = 1, 10 + !$acc loop reduction(+:sum) + do l = 1, 10 + sum = 1 + end do + end do + end do + end do + + !$acc loop reduction(+:sum) + do i = 1, 10 + !$acc loop ! { dg-warning "nested loop in reduction needs reduction clause for .sum." } + do j = 1, 10 + !$acc loop ! { dg-warning "nested loop in reduction needs reduction clause for .sum." } + ! { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 } + do k = 1, 10 + !$acc loop reduction(+:sum) + do l = 1, 10 + sum = 1 + end do + end do + end do + end do + + !$acc loop reduction(+:sum) + do i = 1, 10 + !$acc loop reduction(-:sum) ! { dg-warning "conflicting reduction operations for .sum." } + do j = 1, 10 + !$acc loop reduction(+:sum) ! { dg-warning "conflicting reduction operations for .sum." } + do k = 1, 10 + sum = 1 + end do + end do + end do + + !$acc loop reduction(+:sum) + do i = 1, 10 + !$acc loop reduction(-:sum) ! { dg-warning "conflicting reduction operations for .sum." } + do j = 1, 10 + !$acc loop reduction(-:sum) + do k = 1, 10 + sum = 1 + end do + end do + end do + + !$acc loop reduction(+:sum) + do i = 1, 10 + !$acc loop reduction(-:sum) ! { dg-warning "conflicting reduction operations for .sum." } + do j = 1, 10 + !$acc loop ! { dg-warning "nested loop in reduction needs reduction clause for .sum." } + ! { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 } + do k = 1, 10 + !$acc loop reduction(*:sum) ! { dg-warning "conflicting reduction operations for .sum." } + do l = 1, 10 + sum = 1 + end do + end do + end do + end do + + !$acc loop reduction(+:sum) + do i = 1, 10 + !$acc loop reduction(-:sum) ! { dg-warning "conflicting reduction operations for .sum." } + do j = 1, 10 + !$acc loop reduction(+:sum) ! { dg-warning "conflicting reduction operations for .sum." } + ! { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 } + do k = 1, 10 + !$acc loop reduction(*:sum) ! { dg-warning "conflicting reduction operations for .sum." } + do l = 1, 10 + sum = 1 + end do + end do + end do + end do + + !$acc loop reduction(+:sum) reduction(-:diff) + do i = 1, 10 + !$acc loop reduction(-:diff) ! { dg-warning "nested loop in reduction needs reduction clause for .sum." } + do j = 1, 10 + !$acc loop reduction(+:sum) + do k = 1, 10 + sum = 1 + end do + end do + + !$acc loop reduction(+:sum) ! { dg-warning "nested loop in reduction needs reduction clause for .diff." } + do j = 1, 10 + !$acc loop reduction(-:diff) + do k = 1, 10 + diff = 1 + end do + end do + end do + !$acc end parallel +end subroutine acc_parallel + +! The same tests as above, but using a combined parallel loop construct. + +subroutine acc_parallel_loop () + implicit none (type, external) + integer :: h, i, j, k, l, sum, diff + + !$acc parallel loop + do h = 1, 10 + !$acc loop reduction(+:sum) + do i = 1, 10 + !$acc loop ! { dg-warning "nested loop in reduction needs reduction clause for .sum." } + ! { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 } + do j = 1, 10 + !$acc loop reduction(+:sum) + do k = 1, 10 + sum = 1 + end do + end do + end do + + !$acc loop reduction(+:sum) + do i = 1, 10 + !$acc loop collapse(2) ! { dg-warning "nested loop in reduction needs reduction clause for .sum." } + ! { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 } + do j = 1, 10 + do k = 1, 10 + !$acc loop reduction(+:sum) + do l = 1, 10 + sum = 1 + end do + end do + end do + end do + + + !$acc loop reduction(+:sum) + do i = 1, 10 + !$acc loop ! { dg-warning "nested loop in reduction needs reduction clause for .sum." } + ! { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 } + do j = 1, 10 + !$acc loop ! { dg-warning "nested loop in reduction needs reduction clause for .sum." } + ! { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 } + do k = 1, 10 + !$acc loop reduction(+:sum) + do l = 1, 10 + sum = 1 + end do + end do + end do + end do + + !$acc loop reduction(+:sum) + do i = 1, 10 + !$acc loop reduction(-:sum) ! { dg-warning "conflicting reduction operations for .sum." } + ! { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 } + do j = 1, 10 + !$acc loop reduction(+:sum) ! { dg-warning "conflicting reduction operations for .sum." } + do k = 1, 10 + sum = 1 + end do + end do + end do + + !$acc loop reduction(+:sum) + do i = 1, 10 + !$acc loop reduction(-:sum) ! { dg-warning "conflicting reduction operations for .sum." } + ! { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 } + do j = 1, 10 + !$acc loop reduction(-:sum) + do k = 1, 10 + sum = 1 + end do + end do + end do + + !$acc loop reduction(+:sum) + do i = 1, 10 + !$acc loop reduction(-:sum) ! { dg-warning "conflicting reduction operations for .sum." } + ! { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 } + do j = 1, 10 + !$acc loop ! { dg-warning "nested loop in reduction needs reduction clause for .sum." } + ! { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 } + do k = 1, 10 + !$acc loop reduction(*:sum) ! { dg-warning "conflicting reduction operations for .sum." } + do l = 1, 10 + sum = 1 + end do + end do + end do + end do + + !$acc loop reduction(+:sum) + do i = 1, 10 + !$acc loop reduction(-:sum) ! { dg-warning "conflicting reduction operations for .sum." } + ! { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 } + do j = 1, 10 + !$acc loop reduction(+:sum) ! { dg-warning "conflicting reduction operations for .sum." } + ! { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 } + do k = 1, 10 + !$acc loop reduction(*:sum) ! { dg-warning "conflicting reduction operations for .sum." } + do l = 1, 10 + sum = 1 + end do + end do + end do + end do + + !$acc loop reduction(+:sum) reduction(-:diff) + do i = 1, 10 + !$acc loop reduction(-:diff) ! { dg-warning "nested loop in reduction needs reduction clause for .sum." } + ! { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 } + do j = 1, 10 + !$acc loop reduction(+:sum) + do k = 1, 10 + sum = 1 + end do + end do + + !$acc loop reduction(+:sum) ! { dg-warning "nested loop in reduction needs reduction clause for .diff." } + ! { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 } + do j = 1, 10 + !$acc loop reduction(-:diff) + do k = 1, 10 + diff = 1 + end do + end do + end do + end do +end subroutine acc_parallel_loop + +! The same tests as above, but now the outermost reduction clause is on +! the parallel region, not the outermost loop. + +subroutine acc_parallel_reduction () + implicit none (type, external) + integer :: i, j, k, l, sum, diff + + !$acc parallel reduction(+:sum) + !$acc loop ! { dg-warning "nested loop in reduction needs reduction clause for .sum." } + do i = 1, 10 + !$acc loop ! { dg-warning "nested loop in reduction needs reduction clause for .sum." } + do j = 1, 10 + !$acc loop reduction(+:sum) + do k = 1, 10 + sum = 1 + end do + end do + end do + + !$acc loop ! { dg-warning "nested loop in reduction needs reduction clause for .sum." } + do i = 1, 10 + !$acc loop collapse(2) ! { dg-warning "nested loop in reduction needs reduction clause for .sum." } + do j = 1, 10 + do k = 1, 10 + !$acc loop reduction(+:sum) + do l = 1, 10 + sum = 1 + end do + end do + end do + end do + + !$acc loop ! { dg-warning "nested loop in reduction needs reduction clause for .sum." } + do i = 1, 10 + !$acc loop ! { dg-warning "nested loop in reduction needs reduction clause for .sum." } + do j = 1, 10 + !$acc loop ! { dg-warning "nested loop in reduction needs reduction clause for .sum." } + ! { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 } + do k = 1, 10 + !$acc loop reduction(+:sum) + do l = 1, 10 + sum = 1 + end do + end do + end do + end do + + !$acc loop ! { dg-warning "nested loop in reduction needs reduction clause for .sum." } + do i = 1, 10 + !$acc loop reduction(-:sum) ! { dg-warning "conflicting reduction operations for .sum." } + do j = 1, 10 + !$acc loop reduction(+:sum) ! { dg-warning "conflicting reduction operations for .sum." } + do k = 1, 10 + sum = 1 + end do + end do + end do + + !$acc loop ! { dg-warning "nested loop in reduction needs reduction clause for .sum." } + do i = 1, 10 + !$acc loop reduction(-:sum) ! { dg-warning "conflicting reduction operations for .sum." } + do j = 1, 10 + !$acc loop reduction(-:sum) + do k = 1, 10 + sum = 1 + end do + end do + end do + + !$acc loop reduction(max:sum) ! { dg-warning "conflicting reduction operations for .sum." } + do i = 1, 10 + !$acc loop reduction(-:sum) ! { dg-warning "conflicting reduction operations for .sum." } + do j = 1, 10 + !$acc loop ! { dg-warning "nested loop in reduction needs reduction clause for .sum." } + ! { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 } + do k = 1, 10 + !$acc loop reduction(*:sum) ! { dg-warning "conflicting reduction operations for .sum." } + do l = 1, 10 + sum = 1 + end do + end do + end do + end do + + !$acc loop reduction(max:sum) ! { dg-warning "conflicting reduction operations for .sum." } + do i = 1, 10 + !$acc loop reduction(-:sum) ! { dg-warning "conflicting reduction operations for .sum." } + do j = 1, 10 + !$acc loop reduction(+:sum) ! { dg-warning "conflicting reduction operations for .sum." } + ! { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 } + do k = 1, 10 + !$acc loop reduction(*:sum) ! { dg-warning "conflicting reduction operations for .sum." } + do l = 1, 10 + sum = 1 + end do + end do + end do + end do + + !$acc loop reduction(-:diff) ! { dg-warning "nested loop in reduction needs reduction clause for .sum." } + do i = 1, 10 + !$acc loop reduction(-:diff) ! { dg-warning "nested loop in reduction needs reduction clause for .sum." } + do j = 1, 10 + !$acc loop reduction(+:sum) + do k = 1, 10 + sum = 1 + end do + end do + + !$acc loop reduction(+:sum) ! { dg-warning "nested loop in reduction needs reduction clause for .diff." } + do j = 1, 10 + !$acc loop reduction(-:diff) + do k = 1, 10 + diff = 1 + end do + end do + end do + !$acc end parallel +end subroutine acc_parallel_reduction + +! 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. */ +subroutine acc_parallel_loop_reduction () + implicit none (type, external) + integer :: h, i, j, k, l, sum, diff + + !$acc parallel loop reduction(+:sum) + do h = 1, 10 + !$acc loop ! { dg-warning "nested loop in reduction needs reduction clause for .sum." } + do i = 1, 10 + !$acc loop ! { dg-warning "nested loop in reduction needs reduction clause for .sum." } + ! { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 } + do j = 1, 10 + !$acc loop reduction(+:sum) + do k = 1, 10 + sum = 1 + end do + end do + end do + + !$acc loop ! { dg-warning "nested loop in reduction needs reduction clause for .sum." } + do i = 1, 10 + !$acc loop collapse(2) ! { dg-warning "nested loop in reduction needs reduction clause for .sum." } + ! { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 } + do j = 1, 10 + do k = 1, 10 + !$acc loop reduction(+:sum) + do l = 1, 10 + sum = 1 + end do + end do + end do + end do + + !$acc loop ! { dg-warning "nested loop in reduction needs reduction clause for .sum." } + do i = 1, 10 + !$acc loop ! { dg-warning "nested loop in reduction needs reduction clause for .sum." } + ! { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 } + do j = 1, 10 + !$acc loop ! { dg-warning "nested loop in reduction needs reduction clause for .sum." } + ! { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 } + do k = 1, 10 + !$acc loop reduction(+:sum) + do l = 1, 10 + sum = 1 + end do + end do + end do + end do + + !$acc loop ! { dg-warning "nested loop in reduction needs reduction clause for .sum." } + do i = 1, 10 + !$acc loop reduction(-:sum) ! { dg-warning "conflicting reduction operations for .sum." } + ! { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 } + do j = 1, 10 + !$acc loop reduction(+:sum) ! { dg-warning "conflicting reduction operations for .sum." } + do k = 1, 10 + sum = 1 + end do + end do + end do + + !$acc loop ! { dg-warning "nested loop in reduction needs reduction clause for .sum." } + do i = 1, 10 + !$acc loop reduction(-:sum) ! { dg-warning "conflicting reduction operations for .sum." } + ! { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 } + do j = 1, 10 + !$acc loop reduction(-:sum) + do k = 1, 10 + sum = 1 + end do + end do + end do + + !$acc loop reduction(max:sum) ! { dg-warning "conflicting reduction operations for .sum." } + do i = 1, 10 + !$acc loop reduction(-:sum) ! { dg-warning "conflicting reduction operations for .sum." } + ! { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 } + do j = 1, 10 + !$acc loop ! { dg-warning "nested loop in reduction needs reduction clause for .sum." } + ! { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 } + do k = 1, 10 + !$acc loop reduction(*:sum) ! { dg-warning "conflicting reduction operations for .sum." } + do l = 1, 10 + sum = 1 + end do + end do + end do + end do + + !$acc loop reduction(max:sum) ! { dg-warning "conflicting reduction operations for .sum." } + do i = 1, 10 + !$acc loop reduction(-:sum) ! { dg-warning "conflicting reduction operations for .sum." } + ! { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 } + do j = 1, 10 + !$acc loop reduction(+:sum) ! { dg-warning "conflicting reduction operations for .sum." } + ! { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 } + do k = 1, 10 + !$acc loop reduction(*:sum) ! { dg-warning "conflicting reduction operations for .sum." } + do l = 1, 10 + sum = 1 + end do + end do + end do + end do + + !$acc loop reduction(-:diff) ! { dg-warning "nested loop in reduction needs reduction clause for .sum." } + do i = 1, 10 + !$acc loop reduction(-:diff) ! { dg-warning "nested loop in reduction needs reduction clause for .sum." } + ! { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 } + do j = 1, 10 + !$acc loop reduction(+:sum) + do k = 1, 10 + sum = 1 + end do + end do + + !$acc loop reduction(+:sum) ! { dg-warning "nested loop in reduction needs reduction clause for .diff." } + ! { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 } + do j = 1, 10 + !$acc loop reduction(-:diff) + do k = 1, 10 + diff = 1 + end do + end do + end do + end do +end subroutine acc_parallel_loop_reduction + +! The same tests as above, but inside a routine construct. +subroutine acc_routine () + implicit none (type, external) + !$acc routine gang + integer :: i, j, k, l, sum, diff + + !$acc loop reduction(+:sum) + do i = 1, 10 + !$acc loop ! { dg-warning "nested loop in reduction needs reduction clause for .sum." } + do j = 1, 10 + !$acc loop reduction(+:sum) + do k = 1, 10 + sum = 1 + end do + end do + end do + + !$acc loop reduction(+:sum) + do i = 1, 10 + !$acc loop collapse(2) ! { dg-warning "nested loop in reduction needs reduction clause for .sum." } + do j = 1, 10 + do k = 1, 10 + !$acc loop reduction(+:sum) + do l = 1, 10 + sum = 1 + end do + end do + end do + end do + + !$acc loop reduction(+:sum) + do i = 1, 10 + !$acc loop ! { dg-warning "nested loop in reduction needs reduction clause for .sum." } + do j = 1, 10 + !$acc loop ! { dg-warning "nested loop in reduction needs reduction clause for .sum." } + ! { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 } + do k = 1, 10 + !$acc loop reduction(+:sum) + do l = 1, 10 + sum = 1 + end do + end do + end do + end do + + !$acc loop reduction(+:sum) + do i = 1, 10 + !$acc loop reduction(-:sum) ! { dg-warning "conflicting reduction operations for .sum." } + do j = 1, 10 + !$acc loop reduction(+:sum) ! { dg-warning "conflicting reduction operations for .sum." } + do k = 1, 10 + sum = 1 + end do + end do + end do + + !$acc loop reduction(+:sum) + do i = 1, 10 + !$acc loop reduction(-:sum) ! { dg-warning "conflicting reduction operations for .sum." } + do j = 1, 10 + !$acc loop reduction(-:sum) + do k = 1, 10 + sum = 1 + end do + end do + end do + + !$acc loop reduction(+:sum) + do i = 1, 10 + !$acc loop reduction(-:sum) ! { dg-warning "conflicting reduction operations for .sum." } + do j = 1, 10 + !$acc loop ! { dg-warning "nested loop in reduction needs reduction clause for .sum." } + ! { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 } + do k = 1, 10 + !$acc loop reduction(*:sum) ! { dg-warning "conflicting reduction operations for .sum." } + do l = 1, 10 + sum = 1 + end do + end do + end do + end do + + !$acc loop reduction(+:sum) + do i = 1, 10 + !$acc loop reduction(-:sum) ! { dg-warning "conflicting reduction operations for .sum." } + do j = 1, 10 + !$acc loop reduction(+:sum) ! { dg-warning "conflicting reduction operations for .sum." } + ! { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 } + do k = 1, 10 + !$acc loop reduction(*:sum) ! { dg-warning "conflicting reduction operations for .sum." } + do l = 1, 10 + sum = 1 + end do + end do + end do + end do + + !$acc loop reduction(+:sum) reduction(-:diff) + do i = 1, 10 + !$acc loop reduction(-:diff) ! { dg-warning "nested loop in reduction needs reduction clause for .sum." } + do j = 1, 10 + !$acc loop reduction(+:sum) + do k = 1, 10 + sum = 1 + end do + end do + + !$acc loop reduction(+:sum) ! { dg-warning "nested loop in reduction needs reduction clause for .diff." } + do j = 1, 10 + !$acc loop reduction(-:diff) + do k = 1, 10 + diff = 1 + end do + end do + end do +end subroutine acc_routine + +subroutine acc_kernels () + integer :: i, j, k, sum, diff + + ! FIXME: No diagnostics are produced for these loops because reductions + ! in kernels regions are not supported yet. + !$acc kernels + !$acc loop reduction(+:sum) + do i = 1, 10 + do j = 1, 10 + do k = 1, 10 + sum = 1 + end do + end do + end do + + !$acc loop reduction(+:sum) + do i = 1, 10 + !$acc loop + do j = 1, 10 + do k = 1, 10 + sum = 1 + end do + end do + end do + + !$acc loop reduction(+:sum) + do i = 1, 10 + !$acc loop reduction(-:diff) + do j = 1, 10 + !$acc loop + do k = 1, 10 + sum = 1 + end do + end do + end do + + !$acc loop reduction(+:sum) + do i = 1, 10 + !$acc loop + do j = 1, 10 + !$acc loop reduction(+:sum) + do k = 1, 10 + sum = 1 + end do + end do + end do + + !$acc loop reduction(+:sum) + do i = 1, 10 + !$acc loop reduction(-:sum) + do j = 1, 10 + !$acc loop reduction(+:sum) + do k = 1, 10 + sum = 1 + end do + end do + end do + !$acc end kernels +end subroutine acc_kernels diff --git a/gcc/testsuite/gfortran.dg/goacc/nested-reductions.f90 b/gcc/testsuite/gfortran.dg/goacc/nested-reductions.f90 new file mode 100644 index 00000000000..3becafabe62 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/goacc/nested-reductions.f90 @@ -0,0 +1,540 @@ +! Test cases of nested reduction loops that should compile cleanly. + +subroutine acc_parallel () + implicit none (type, external) + integer :: i, j, k, sum, diff + + !$acc parallel + !$acc loop reduction(+:sum) + do i = 1, 10 + do j = 1, 10 + do k = 1, 10 + sum = 1 + end do + end do + end do + + !$acc loop collapse(2) reduction(+:sum) + do i = 1, 10 + do j = 1, 10 + do k = 1, 10 + sum = 1 + end do + end do + end do + + !$acc loop reduction(+:sum) + do i = 1, 10 + !$acc loop reduction(+:sum) + do j = 1, 10 + do k = 1, 10 + sum = 1 + end do + end do + end do + + !$acc loop reduction(+:sum) + do i = 1, 10 + !$acc loop collapse(2) reduction(+:sum) + do j = 1, 10 + do k = 1, 10 + sum = 1 + end do + end do + end do + + !$acc loop reduction(+:sum) + do i = 1, 10 + do j = 1, 10 + !$acc loop reduction(+:sum) + do k = 1, 10 + sum = 1 + end do + end do + end do + + !$acc loop reduction(+:sum) + do i = 1, 10 + !$acc loop reduction(+:sum) + do j = 1, 10 + !$acc loop reduction(+:sum) + do k = 1, 10 + sum = 1 + end do + end do + end do + + !$acc loop reduction(+:sum) reduction(-:diff) + do i = 1, 10 + !$acc loop reduction(+:sum) + do j = 1, 10 + !$acc loop reduction(+:sum) + do k = 1, 10 + sum = 1 + end do + end do + + !$acc loop reduction(-:diff) + do j = 1, 10 + !$acc loop reduction(-:diff) + do k = 1, 10 + diff = 1 + end do + end do + end do + !$acc end parallel +end subroutine acc_parallel + +! The same tests as above, but using a combined parallel loop construct. + +subroutine acc_parallel_loop () + implicit none (type, external) + integer :: h, i, j, k, l, sum, diff + + !$acc parallel loop + do h = 1, 10 + !$acc loop reduction(+:sum) + do i = 1, 10 + do j = 1, 10 + do k = 1, 10 + sum = 1 + end do + end do + end do + + !$acc loop collapse(2) reduction(+:sum) + do i = 1, 10 + do j = 1, 10 + do k = 1, 10 + sum = 1 + end do + end do + end do + + !$acc loop reduction(+:sum) + do i = 1, 10 + !$acc loop reduction(+:sum) + do j = 1, 10 + do k = 1, 10 + sum = 1 + end do + end do + end do + + !$acc loop reduction(+:sum) + do i = 1, 10 + !$acc loop collapse(2) reduction(+:sum) + do j = 1, 10 + do k = 1, 10 + sum = 1 + end do + end do + end do + + !$acc loop reduction(+:sum) + do i = 1, 10 + do j = 1, 10 + !$acc loop reduction(+:sum) + do k = 1, 10 + sum = 1 + end do + end do + end do + + !$acc loop reduction(+:sum) + do i = 1, 10 + !$acc loop reduction(+:sum) ! { dg-warning "insufficient partitioning available to parallelize loop" } + do j = 1, 10 + !$acc loop reduction(+:sum) + do k = 1, 10 + sum = 1 + end do + end do + end do + + !$acc loop reduction(+:sum) reduction(-:diff) + do i = 1, 10 + !$acc loop reduction(+:sum) ! { dg-warning "insufficient partitioning available to parallelize loop" } + do j = 1, 10 + !$acc loop reduction(+:sum) + do k = 1, 10 + sum = 1 + end do + end do + + !$acc loop reduction(-:diff) ! { dg-warning "insufficient partitioning available to parallelize loop" } + do j = 1, 10 + !$acc loop reduction(-:diff) + do k = 1, 10 + diff = 1 + end do + end do + end do + end do +end subroutine acc_parallel_loop + +! The same tests as above, but now the outermost reduction clause is on +! the parallel region, not the outermost loop. */ + +subroutine acc_parallel_reduction () + implicit none (type, external) + integer :: i, j, k, sum, diff + + !$acc parallel reduction(+:sum) + do i = 1, 10 + do j = 1, 10 + do k = 1, 10 + sum = 1 + end do + end do + end do + + do i = 1, 10 + !$acc loop + do j = 1, 10 + do k = 1, 10 + sum = 1 + end do + end do + end do + + !$acc loop reduction(+:sum) + do i = 1, 10 + do j = 1, 10 + !$acc loop reduction(+:sum) + do k = 1, 10 + sum = 1 + end do + end do + end do + + do i = 1, 10 + do j = 1, 10 + !$acc loop + do k = 1, 10 + sum = 1 + end do + end do + end do + + !$acc loop reduction(+:sum) + do i = 1, 10 + !$acc loop reduction(+:sum) + do j = 1, 10 + !$acc loop reduction(+:sum) + do k = 1, 10 + sum = 1 + end do + end do + end do + + !$acc loop reduction(+:sum) reduction(-:diff) + do i = 1, 10 + !$acc loop reduction(+:sum) + do j = 1, 10 + !$acc loop reduction(+:sum) + do k = 1, 10 + sum = 1 + end do + end do + + !$acc loop reduction(-:diff) + do j = 1, 10 + !$acc loop reduction(-:diff) + do k = 1, 10 + diff = 1 + end do + end do + end do + + !$acc loop reduction(+:sum) + do i = 1, 10 + !$acc loop reduction(+:sum) + do j = 1, 10 + !$acc loop reduction(+:sum) + do k = 1, 10 + sum = 1 + end do + end do + + !$acc loop reduction(-:diff) + do j = 1, 10 + !$acc loop reduction(-:diff) + do k = 1, 10 + diff = 1 + end do + end do + end do + + !$acc loop reduction(+:sum) + do i = 1, 10 + !$acc loop reduction(+:sum) + do j = 1, 10 + !$acc loop reduction(+:sum) + do k = 1, 10 + sum = 1 + end do + end do + + !$acc loop + do j = 1, 10 + !$acc loop reduction(-:diff) + do k = 1, 10 + diff = 1 + end do + end do + end do + !$acc end parallel +end subroutine acc_parallel_reduction + +! 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. */ +subroutine acc_parallel_loop_reduction () + implicit none (type, external) + integer :: h, i, j, k, sum, diff + + !$acc parallel loop reduction(+:sum) + do h = 1, 10 + do i = 1, 10 + do j = 1, 10 + do k = 1, 10 + sum = 1 + end do + end do + end do + + do i = 1, 10 + !$acc loop + do j = 1, 10 + do k = 1, 10 + sum = 1 + end do + end do + end do + + !$acc loop reduction(+:sum) + do i = 1, 10 + do j = 1, 10 + !$acc loop reduction(+:sum) + do k = 1, 10 + sum = 1 + end do + end do + end do + + do i = 1, 10 + do j = 1, 10 + !$acc loop + do k = 1, 10 + sum = 1 + end do + end do + end do + + !$acc loop reduction(+:sum) + do i = 1, 10 + !$acc loop reduction(+:sum) ! { dg-warning "insufficient partitioning available to parallelize loop" } + do j = 1, 10 + !$acc loop reduction(+:sum) + do k = 1, 10 + sum = 1 + end do + end do + end do + + !$acc loop reduction(+:sum) reduction(-:diff) + do i = 1, 10 + !$acc loop reduction(+:sum) ! { dg-warning "insufficient partitioning available to parallelize loop" } + do j = 1, 10 + !$acc loop reduction(+:sum) + do k = 1, 10 + sum = 1 + end do + end do + + !$acc loop reduction(-:diff) ! { dg-warning "insufficient partitioning available to parallelize loop" } + do j = 1, 10 + !$acc loop reduction(-:diff) + do k = 1, 10 + diff = 1 + end do + end do + end do + + !$acc loop reduction(+:sum) + do i = 1, 10 + !$acc loop reduction(+:sum) ! { dg-warning "insufficient partitioning available to parallelize loop" } + do j = 1, 10 + !$acc loop reduction(+:sum) + do k = 1, 10 + sum = 1 + end do + end do + + !$acc loop reduction(-:diff) ! { dg-warning "insufficient partitioning available to parallelize loop" } + do j = 1, 10 + !$acc loop reduction(-:diff) + do k = 1, 10 + diff = 1 + end do + end do + end do + + !$acc loop reduction(+:sum) + do i = 1, 10 + !$acc loop reduction(+:sum) ! { dg-warning "insufficient partitioning available to parallelize loop" } + do j = 1, 10 + !$acc loop reduction(+:sum) + do k = 1, 10 + sum = 1 + end do + end do + + !$acc loop ! { dg-warning "insufficient partitioning available to parallelize loop" } + do j = 1, 10 + !$acc loop reduction(-:diff) + do k = 1, 10 + diff = 1 + end do + end do + end do + end do +end subroutine acc_parallel_loop_reduction + +! The same tests as above, but inside a routine construct. +subroutine acc_routine () + implicit none (type, external) + !$acc routine gang + + integer :: i, j, k, sum, diff + + !$acc loop reduction(+:sum) + do i = 1, 10 + do j = 1, 10 + do k = 1, 10 + sum = 1 + end do + end do + end do + + !$acc loop collapse(2) reduction(+:sum) + do i = 1, 10 + do j = 1, 10 + do k = 1, 10 + sum = 1 + end do + end do + end do + + !$acc loop reduction(+:sum) + do i = 1, 10 + !$acc loop reduction(+:sum) + do j = 1, 10 + do k = 1, 10 + sum = 1 + end do + end do + end do + + !$acc loop reduction(+:sum) + do i = 1, 10 + !$acc loop collapse(2) reduction(+:sum) + do j = 1, 10 + do k = 1, 10 + sum = 1 + end do + end do + end do + + !$acc loop reduction(+:sum) + do i = 1, 10 + do j = 1, 10 + !$acc loop reduction(+:sum) + do k = 1, 10 + sum = 1 + end do + end do + end do + + !$acc loop reduction(+:sum) + do i = 1, 10 + !$acc loop reduction(+:sum) + do j = 1, 10 + !$acc loop reduction(+:sum) + do k = 1, 10 + sum = 1 + end do + end do + end do + + !$acc loop reduction(+:sum) reduction(-:diff) + do i = 1, 10 + !$acc loop reduction(+:sum) + do j = 1, 10 + !$acc loop reduction(+:sum) + do k = 1, 10 + sum = 1 + end do + end do + + !$acc loop reduction(-:diff) + do j = 1, 10 + !$acc loop reduction(-:diff) + do k = 1, 10 + diff = 1 + end do + end do + end do +end subroutine acc_routine + +subroutine acc_kernels () + implicit none (type, external) + integer :: i, j, k, sum, diff + + ! FIXME: These tests are not meaningful yet because reductions in + ! kernels regions are not supported yet. + !$acc kernels + !$acc loop reduction(+:sum) + do i = 1, 10 + do j = 1, 10 + do k = 1, 10 + sum = 1 + end do + end do + end do + + !$acc loop reduction(+:sum) + do i = 1, 10 + !$acc loop reduction(+:sum) + do j = 1, 10 + do k = 1, 10 + sum = 1 + end do + end do + end do + + + !$acc loop reduction(+:sum) + do i = 1, 10 + do j = 1, 10 + !$acc loop reduction(+:sum) + do k = 1, 10 + sum = 1 + end do + end do + end do + + + !$acc loop reduction(+:sum) + do i = 1, 10 + !$acc loop reduction(+:sum) + do j = 1, 10 + !$acc loop reduction(+:sum) + do k = 1, 10 + sum = 1 + end do + end do + end do + + !$acc end kernels +end subroutine acc_kernels diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog index 15521b48bd8..067d247da9a 100644 --- a/libgomp/ChangeLog +++ b/libgomp/ChangeLog @@ -1,3 +1,14 @@ +2019-11-06 Thomas Schwinge + + * testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-1.c: + Add expected warnings about 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-11-04 Tobias Burnus * testsuite/libgomp.fortran/pr66199-1.f90: Remove diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-1.c index 5e82e1d350c..8c08717fdf2 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 /* { dg-warning "nested loop in reduction needs reduction clause for 'res'" "TODO" } */ for (j = 0; j < 32; j++) { #pragma acc loop worker reduction(+:res) diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-2.c index a339f327956..41042accbde 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 /* { dg-warning "nested loop in reduction needs reduction clause for 'res'" "TODO" } */ for (j = 0; j < 32; j++) { #pragma acc loop worker vector reduction(^:res) diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-3.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-3.c index 6369d7fbb33..ace1e005e2e 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 /* { dg-warning "nested loop in reduction needs reduction clause for 'res'" "TODO" } */ for (j = 0; j < 32; j++) { #pragma acc loop worker vector reduction(+:res) diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-4.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-4.c index 140c3226327..c3cc12fa953 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 /* { dg-warning "nested loop in reduction needs reduction clause for 'm\?res'" "TODO" } */ for (j = 0; j < 32; j++) { #pragma acc loop worker vector reduction(+:res) reduction(max:mres)