From patchwork Wed Nov 6 12:41:47 2019 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 8bit X-Patchwork-Submitter: Frederik Harwath X-Patchwork-Id: 1190380 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@bilbo.ozlabs.org Authentication-Results: ozlabs.org; spf=pass (sender SPF authorized) smtp.mailfrom=gcc.gnu.org (client-ip=209.132.180.131; helo=sourceware.org; envelope-from=gcc-patches-return-512596-incoming=patchwork.ozlabs.org@gcc.gnu.org; receiver=) Authentication-Results: ozlabs.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: ozlabs.org; dkim=pass (1024-bit key; unprotected) header.d=gcc.gnu.org header.i=@gcc.gnu.org header.b="EzS5EAFh"; dkim-atps=neutral Received: from sourceware.org (server1.sourceware.org [209.132.180.131]) (using TLSv1.2 with cipher ECDHE-RSA-AES256-GCM-SHA384 (256/256 bits)) (No client certificate requested) by ozlabs.org (Postfix) with ESMTPS id 477R586YTcz9sNx for ; Wed, 6 Nov 2019 23:42:36 +1100 (AEDT) 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= DKIM-Signature: v=1; a=rsa-sha1; c=relaxed; 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; s=default; bh=9Fsdo3f8JiN3zPJ4A4SLRcwrKsI=; b=EzS5EAFhfAzY440vkU/iV1Vmx13+ ROkgvBIpxdNacW1Mr03D+mHRi4+rs31XPJFh775hXaNvjuTV5WhYa9rp2khjl17Z NE6lOXMLYedvcXzDVMOYsM0tkA6wvfjJn0i/PKLfR9xkDj141/ulHHBtDS9yR+2m P7ycmsx8QsknJe8= Received: (qmail 38990 invoked by alias); 6 Nov 2019 12:42:25 -0000 Mailing-List: contact gcc-patches-help@gcc.gnu.org; run by ezmlm Precedence: bulk List-Id: List-Unsubscribe: List-Archive: List-Post: List-Help: Sender: gcc-patches-owner@gcc.gnu.org Delivered-To: mailing list gcc-patches@gcc.gnu.org Received: (qmail 38975 invoked by uid 89); 6 Nov 2019 12:42:24 -0000 Authentication-Results: sourceware.org; auth=none X-Spam-SWARE-Status: No, score=-18.1 required=5.0 tests=AWL, BAYES_00, GIT_PATCH_0, GIT_PATCH_1, GIT_PATCH_2, GIT_PATCH_3 autolearn=ham version=3.3.1 spammy=collapse, 1286, sk:frederi X-HELO: esa3.mentor.iphmx.com Received: from esa3.mentor.iphmx.com (HELO esa3.mentor.iphmx.com) (68.232.137.180) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with ESMTP; Wed, 06 Nov 2019 12:42:15 +0000 IronPort-SDR: J5nCE6+suLMkqKdyhs56zp0pGjs5NMTMg/XLc2zPQBZaq0Q8dSX628DFgnK8e08LzfFmORXhYd RA5E9yq19uujW1Rgh6S9WdSkLEZXSy++2lipgrtx6Te0JhnpNLOKqjv2pPhSn7K/OHuVqrFyxG HP99mfJGdzrIdqBeQzM9ihsAkawvfeRvM3Ei6pUoUevfstGLeT5p/wR5km8fHw4lnFvUDUJrgf b0fM4q4A10tMZ9JX0P936ZQ/tQaSGuix9038+Jh6YgGnq5XXJSQUNpZgw+G6UdLAfbSNeTwBgz gq0= Received: from orw-gwy-02-in.mentorg.com ([192.94.38.167]) by esa3.mentor.iphmx.com with ESMTP; 06 Nov 2019 04:42:13 -0800 IronPort-SDR: hgWwZV0MYFXTAOvLxVz0ZIu1nxQmjW8Il3+jYGzLcGWKtbumeRq99hWpCGIDJmKAyfLgaXZm43 kfC67DXDum6BcIZ/D0o8KB2EpZo4F7dxPGBtJSSNJIvVhuraU2UV/Ozl+hlxEbI9q8m5IX15wt kJLHGk4gPVs4vq/MUbYM/wMk3w+l2KtPpv0Zs/TXwRasBm+2xftIIQrv4TyfKBlJG09Z/eHbLv /IK4Vd7/pZFYJrLzLz7qY+UVa6ixj9xJKyMUxME7j/lUD1Fygff+dF3lV1ZyJUkXdUY1uH0i1p 5sY= From: To: , CC: 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 X-IsSubscribed: yes From: frederik OpenACC (cf. OpenACC 2.7, section 2.9.11. "reduction clause"; this was first clarified by OpenACC 2.6) requires that, if a variable is used in reduction clauses on two nested loops, then there must be reduction clauses for that variable on all loops that are nested in between the two loops and all these reduction clauses must use the same operator. This commit introduces a check for that property which reports warnings if it is violated. 2019-11-06 Gergö Barany Frederik Harwath Thomas Schwinge gcc/ * omp-low.c (struct omp_context): New fields local_reduction_clauses, outer_reduction_clauses. (new_omp_context): Initialize these. (scan_sharing_clauses): Record reduction clauses on OpenACC constructs. (scan_omp_for): Check reduction clauses for incorrect nesting. gcc/testsuite/ * c-c++-common/goacc/nested-reductions-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. libgomp/ * 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. Reviewed-by: Thomas Schwinge git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/trunk@277875 138bc75d-0d04-0410-961f-82ee72b054a4 --- gcc/ChangeLog | 10 + gcc/omp-low.c | 97 +++ gcc/testsuite/ChangeLog | 9 + .../goacc/nested-reductions-warn.c | 525 ++++++++++++++ .../c-c++-common/goacc/nested-reductions.c | 420 +++++++++++ .../goacc/nested-reductions-warn.f90 | 674 ++++++++++++++++++ .../gfortran.dg/goacc/nested-reductions.f90 | 540 ++++++++++++++ libgomp/ChangeLog | 11 + .../par-loop-comb-reduction-1.c | 2 +- .../par-loop-comb-reduction-2.c | 2 +- .../par-loop-comb-reduction-3.c | 2 +- .../par-loop-comb-reduction-4.c | 2 +- 12 files changed, 2290 insertions(+), 4 deletions(-) create mode 100644 gcc/testsuite/c-c++-common/goacc/nested-reductions-warn.c create mode 100644 gcc/testsuite/c-c++-common/goacc/nested-reductions.c create mode 100644 gcc/testsuite/gfortran.dg/goacc/nested-reductions-warn.f90 create mode 100644 gcc/testsuite/gfortran.dg/goacc/nested-reductions.f90 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 + 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 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 *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 + 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 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 + + * 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 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)