[committed] Warn about inconsistent OpenACC nested reduction clauses

Message ID	20191106124148.26041-1-frederik@codesourcery.com
State	New
Headers	show Return-Path: <gcc-patches-return-512596-incoming=patchwork.ozlabs.org@gcc.gnu.org> DomainKey-Signature: a=rsa-sha1; c=nofws; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender:from :to:cc:subject:date:message-id:in-reply-to:references :mime-version:content-type:content-transfer-encoding; q=dns; s= default; b=ir1QkWZeqOmEeSpf4OFN/A0JC8MNbGbSFZssrMz7gQ1HSjhOdnwI7 fR4BUirDAyA1LjNNgvlc5adJ28TbQINuSBoK6VQZJ362hcthfI+C3hJT4jqD3u4O IyGvakmdIGi1NhIYOcLvmQI1ESd+OeOGM/863WJt9ntHWrzJUSYQ/U= Mailing-List: contact gcc-patches-help@gcc.gnu.org; run by ezmlm Precedence: bulk Sender: gcc-patches-owner@gcc.gnu.org IronPort-SDR: J5nCE6+suLMkqKdyhs56zp0pGjs5NMTMg/XLc2zPQBZaq0Q8dSX628DFgnK8e08LzfFmORXhYd RA5E9yq19uujW1Rgh6S9WdSkLEZXSy++2lipgrtx6Te0JhnpNLOKqjv2pPhSn7K/OHuVqrFyxG HP99mfJGdzrIdqBeQzM9ihsAkawvfeRvM3Ei6pUoUevfstGLeT5p/wR5km8fHw4lnFvUDUJrgf b0fM4q4A10tMZ9JX0P936ZQ/tQaSGuix9038+Jh6YgGnq5XXJSQUNpZgw+G6UdLAfbSNeTwBgz gq0= IronPort-SDR: hgWwZV0MYFXTAOvLxVz0ZIu1nxQmjW8Il3+jYGzLcGWKtbumeRq99hWpCGIDJmKAyfLgaXZm43 kfC67DXDum6BcIZ/D0o8KB2EpZo4F7dxPGBtJSSNJIvVhuraU2UV/Ozl+hlxEbI9q8m5IX15wt kJLHGk4gPVs4vq/MUbYM/wMk3w+l2KtPpv0Zs/TXwRasBm+2xftIIQrv4TyfKBlJG09Z/eHbLv /IK4Vd7/pZFYJrLzLz7qY+UVa6ixj9xJKyMUxME7j/lUD1Fygff+dF3lV1ZyJUkXdUY1uH0i1p 5sY= From: <frederik@codesourcery.com> To: <gcc-patches@gcc.gnu.org>, <thomas@codesourcery.com> CC: <jakub@redhat.com> Subject: [PATCH][committed] Warn about inconsistent OpenACC nested reduction clauses Date: Wed, 6 Nov 2019 13:41:47 +0100 Message-ID: <20191106124148.26041-1-frederik@codesourcery.com> In-Reply-To: <87y2wuqu6h.fsf@euler.schwinge.homeip.net> References: <87y2wuqu6h.fsf@euler.schwinge.homeip.net> MIME-Version: 1.0 Content-Type: text/plain; charset="UTF-8" Content-Transfer-Encoding: 8bit
Series	[committed] Warn about inconsistent OpenACC nested reduction clauses \| expand [committed] Warn about inconsistent OpenACC nested reduction clauses

diff --git a/gcc/ChangeLog b/gcc/ChangeLog index 7fee0f37e9bf..38160dd631e9 100644 --- a/gcc/ChangeLog +++ b/gcc/ChangeLog @@ -1,3 +1,13 @@ +2019-11-06 Gergö Barany <gergo@codesourcery.com> + Frederik Harwath <frederik@codesourcery.com> + Thomas Schwinge <thomas@codesourcery.com> + + * omp-low.c (struct omp_context): New fields + local_reduction_clauses, outer_reduction_clauses. + (new_omp_context): Initialize these. + (scan_sharing_clauses): Record reduction clauses on OpenACC constructs. + (scan_omp_for): Check reduction clauses for incorrect nesting. + 2019-11-06 Jakub Jelinek <jakub@redhat.com> PR inline-asm/92352 diff --git a/gcc/omp-low.c b/gcc/omp-low.c index 122f42788813..fa76ceba33c6 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<tree, tree> *lastprivate_conditional_map; + /* A tree_list of the reduction clauses in this context. */ + tree local_reduction_clauses; + + /* A tree_list of the reduction clauses in outer contexts. */ + tree outer_reduction_clauses; + /* Nesting depth of this context. Used to beautify error messages re invalid gotos. The outermost ctx is depth 1, with depth 0 being reserved for the main body of the function. */ @@ -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 a0f52ec2c5cd..64568e5661af 100644 --- a/gcc/testsuite/ChangeLog +++ b/gcc/testsuite/ChangeLog @@ -1,3 +1,12 @@ +2019-11-06 Gergö Barany <gergo@codesourcery.com> + Frederik Harwath <frederik@codesourcery.com> + Thomas Schwinge <thomas@codesourcery.com> + + * 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 <jakub@redhat.com> 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 000000000000..e2af66e4fa3b --- /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 000000000000..15385c4da992 --- /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 000000000000..ec36bc9616e0 --- /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 000000000000..3becafabe622 --- /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 15521b48bd88..067d247da9ac 100644 --- a/libgomp/ChangeLog +++ b/libgomp/ChangeLog @@ -1,3 +1,14 @@ +2019-11-06 Thomas Schwinge <thomas@codesourcery.com> + + * 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 <tobias@codesourcery.com> * 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 5e82e1d350cc..8c08717fdf2c 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 a339f327956e..41042accbde5 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 6369d7fbb330..ace1e005e2e2 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 140c32263278..c3cc12fa953a 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)

[committed] Warn about inconsistent OpenACC nested reduction clauses

Commit Message

Comments

Patch