From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: from esa4.mentor.iphmx.com (esa4.mentor.iphmx.com [68.232.137.252]) by sourceware.org (Postfix) with ESMTPS id B642D385DC0F; Mon, 20 Jul 2020 10:26:55 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.3.2 sourceware.org B642D385DC0F Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=Frederik_Harwath@mentor.com IronPort-SDR: 46h+8fD11VeXguT7r37FkkR2Pl2OGyN67SFrnLR2Ln5auxs81mhzmwPF2sVQIZSZQRjMPjPV19 rJcpJ6GQbR37kZ08z+gaYfXdfwB7ydiKcwqJOzTo/iEDtzNASLsR1WpNgRlvPrIBWPs2tu8mzY wn7Vd/tv23ang9PfJ6bi2YPN4i0Xy1hOgqqHDU9ecULcJOS0En5iJMuuz/iGr1TLCMR7g7eRSZ U4MWmzLk4Eu7h/xUdw5HlbOwiwDAj4dA8uR6Vbafd6lrqMwceT63S3F1Bj8Nbbpf7us5JxOBK4 saI= X-IronPort-AV: E=Sophos;i="5.75,374,1589270400"; d="scan'208,223";a="51196505" Received: from orw-gwy-02-in.mentorg.com ([192.94.38.167]) by esa4.mentor.iphmx.com with ESMTP; 20 Jul 2020 02:26:54 -0800 IronPort-SDR: WPdqUCYDiGD8A6YZyYScc2rjz1TUwNt23ia56r5tWVyEeIwVPBEYbyLjLpgev65EvXWUKskXwJ gWNkXqf4uLY3kRsFOqvBLw2AtTgj1g7lErt1Y+xgJ9hZB4fFCWhL7nTXUZdtbFrzt0Ly3oUBRB CavLfzACxS/fiFEx71iFeSIWOlidim+VbIQ+uQB7ZrznEUBySyZKYA0LLmOVurAtoTEqg0DX2k q9bAFz8lT0zAUyQCRJfco++buP3kRx6ffwOFPzuMLxo3xXY/oFYBw1lsh2fBWJnqcbYqapvMd0 PHU= From: Frederik Harwath To: Thomas Schwinge , Julian Brown , Kwok Cheung Yeung CC: , , Tobias Burnus , Jakub Jelinek Subject: Re: [PATCH] [og10] libgomp, Fortran: Fix OpenACC "gang reduction on an orphan loop" error message In-Reply-To: References: <91521931-373f-a3a6-7a3f-1752f08858a1@codesourcery.com> Date: Mon, 20 Jul 2020 12:26:48 +0200 Message-ID: MIME-Version: 1.0 Content-Type: multipart/mixed; boundary="=-=-=" X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: svr-ies-mbx-06.mgc.mentorg.com (139.181.222.6) To SVR-IES-MBX-04.mgc.mentorg.com (139.181.222.4) X-Spam-Status: No, score=-12.9 required=5.0 tests=BAYES_00, GIT_PATCH_0, HEADER_FROM_DIFFERENT_DOMAINS, KAM_DMARC_STATUS, SPF_HELO_PASS, SPF_PASS, TXREP autolearn=ham autolearn_force=no version=3.4.2 X-Spam-Checker-Version: SpamAssassin 3.4.2 (2018-09-13) on server2.sourceware.org X-BeenThere: fortran@gcc.gnu.org X-Mailman-Version: 2.1.29 Precedence: list List-Id: Fortran mailing list List-Unsubscribe: , List-Archive: List-Help: List-Subscribe: , X-List-Received-Date: Mon, 20 Jul 2020 10:26:58 -0000 --=-=-= Content-Type: text/plain Content-Transfer-Encoding: quoted-printable Thomas Schwinge writes: Hi Thomas, >> Can I include the patch in OG10? > > Unless Julian/Kwok speak up soon: OK, thanks. This has been delayed a bit by my vacation, but I have now committed the patch. > May want to remove "libgomp" from the first line of the commit log -- > this commit doesn't relate to libgomp specifically. > > (Ideally, we'd also test 'serial' construct in addition to 'kernels', > 'parallel', but we can add that later. I anyway have a WIP patch > waiting, adding more 'serial' construct testing, for a different reason, > so I'll include it there.) I forgot to remove "libgomp" from the commit message, sorry, but I have included the test cases for the "serial construct". Best regards, Frederik ----------------- Mentor Graphics (Deutschland) GmbH, Arnulfstra=DFe 201, 80634 M=FCnchen / G= ermany Registergericht M=FCnchen HRB 106955, Gesch=E4ftsf=FChrer: Thomas Heurung, = Alexander Walter --=-=-= Content-Type: text/x-diff Content-Disposition: inline; filename="0001-libgomp-Fortran-Fix-OpenACC-gang-reduction-on-an-orp.patch" >From 7c10ae450b95495dda362cb66770bb78b546592e Mon Sep 17 00:00:00 2001 From: Frederik Harwath Date: Mon, 20 Jul 2020 11:24:21 +0200 Subject: [PATCH] libgomp, Fortran: Fix OpenACC "gang reduction on an orphan loop" error message According to the OpenACC standard version 2.5 and later, reductions on orphaned gang loops are explicitly disallowed (cf. section "Changes from Version 2.0 to 2.5"). A loop is "orphaned" if it is not lexically contained in a compute construct (cf. section "Loop construct" of the OpenACC standard), i.e. in either a "parallel", a "serial", or a "kernels" construct. This commit fixes the check for reductions on orphaned gang loops in the Fortran frontend which (in contrast to the C, C++ frontends) erroneously rejects reductions on gang loops that are contained in "kernels" constructs. 2020-07-20 Frederik Harwath gcc/fortran/ * openmp.c (oacc_is_parallel_or_serial): Removed function. (oacc_is_kernels): New function. (oacc_is_compute_construct): New function. (resolve_oacc_loop_blocks): Use "oacc_is_compute_construct" instead of "oacc_is_parallel_or_serial" for checking that a loop is not orphaned. gcc/testsuite/ * gfortran.dg/goacc/orphan-reductions-2.f90: New test verifying that the "gang reduction on an orphan loop" error message is not emitted for non-orphaned loops. * c-c++-common/goacc/orphan-reductions-2.c: Likewise for C and C++. --- gcc/fortran/ChangeLog | 9 ++ gcc/fortran/openmp.c | 13 ++- gcc/testsuite/ChangeLog | 7 ++ .../c-c++-common/goacc/orphan-reductions-2.c | 103 ++++++++++++++++++ .../gfortran.dg/goacc/orphan-reductions-2.f90 | 87 +++++++++++++++ 5 files changed, 216 insertions(+), 3 deletions(-) create mode 100644 gcc/testsuite/c-c++-common/goacc/orphan-reductions-2.c create mode 100644 gcc/testsuite/gfortran.dg/goacc/orphan-reductions-2.f90 diff --git a/gcc/fortran/ChangeLog b/gcc/fortran/ChangeLog index e86279cb647..5a1f81c286e 100644 --- a/gcc/fortran/ChangeLog +++ b/gcc/fortran/ChangeLog @@ -1,3 +1,12 @@ +2020-07-20 Frederik Harwath + + * openmp.c (oacc_is_parallel_or_serial): Removed function. + (oacc_is_kernels): New function. + (oacc_is_compute_construct): New function. + (resolve_oacc_loop_blocks): Use "oacc_is_compute_construct" + instead of "oacc_is_parallel_or_serial" for checking that a + loop is not orphaned. + 2020-07-08 Harald Anlauf Backported from master: diff --git a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c index ab68e9f2173..706933c869a 100644 --- a/gcc/fortran/openmp.c +++ b/gcc/fortran/openmp.c @@ -5927,9 +5927,16 @@ oacc_is_serial (gfc_code *code) } static bool -oacc_is_parallel_or_serial (gfc_code *code) +oacc_is_kernels (gfc_code *code) { - return oacc_is_parallel (code) || oacc_is_serial (code); + return code->op == EXEC_OACC_KERNELS || code->op == EXEC_OACC_KERNELS_LOOP; +} + +static bool +oacc_is_compute_construct (gfc_code *code) +{ + return oacc_is_parallel (code) || oacc_is_serial (code) + || oacc_is_kernels (code); } static gfc_statement @@ -6223,7 +6230,7 @@ resolve_oacc_loop_blocks (gfc_code *code) for (c = omp_current_ctx; c; c = c->previous) if (!oacc_is_loop (c->code)) break; - if (c == NULL || !oacc_is_parallel_or_serial (c->code)) + if (c == NULL || !oacc_is_compute_construct (c->code)) gfc_error ("gang reduction on an orphan loop at %L", &code->loc); } diff --git a/gcc/testsuite/ChangeLog b/gcc/testsuite/ChangeLog index 59e6c93b07a..fa1937a4ea2 100644 --- a/gcc/testsuite/ChangeLog +++ b/gcc/testsuite/ChangeLog @@ -1,3 +1,10 @@ +2020-07-20 Frederik Harwath + + * gfortran.dg/goacc/orphan-reductions-2.f90: New test + verifying that the "gang reduction on an orphan loop" error message + is not emitted for non-orphaned loops. + * c-c++-common/goacc/orphan-reductions-2.c: Likewise for C and C++. + 2020-07-12 Jakub Jelinek Backported from master: diff --git a/gcc/testsuite/c-c++-common/goacc/orphan-reductions-2.c b/gcc/testsuite/c-c++-common/goacc/orphan-reductions-2.c new file mode 100644 index 00000000000..d30321710dd --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/orphan-reductions-2.c @@ -0,0 +1,103 @@ +/* Verify that the error message for gang reduction on orphaned OpenACC loops + is not reported for non-orphaned loops. */ + +#include + +int +kernels (int n) +{ + int i, s1 = 0, s2 = 0; +#pragma acc kernels + { +#pragma acc loop gang reduction(+:s1) /* { dg-bogus "gang reduction on an orphan loop" } */ + for (i = 0; i < n; i++) + s1 = s1 + 2; + +#pragma acc loop gang reduction(+:s2) /* { dg-bogus "gang reduction on an orphan loop" } */ + for (i = 0; i < n; i++) + s2 = s2 + 2; + } + return s1 + s2; +} + +int +parallel (int n) +{ + int i, s1 = 0, s2 = 0; +#pragma acc parallel + { +#pragma acc loop gang reduction(+:s1) /* { dg-bogus "gang reduction on an orphan loop" } */ + for (i = 0; i < n; i++) + s1 = s1 + 2; + +#pragma acc loop gang reduction(+:s2) /* { dg-bogus "gang reduction on an orphan loop" } */ + for (i = 0; i < n; i++) + s2 = s2 + 2; + } + return s1 + s2; +} + +int +serial (int n) +{ + int i, s1 = 0, s2 = 0; +#pragma acc serial /* { dg-warning "region contains gang partitioned code but is not gang partitioned" } */ + { +#pragma acc loop gang reduction(+:s1) /* { dg-bogus "gang reduction on an orphan loop" } */ + for (i = 0; i < n; i++) + s1 = s1 + 2; + +#pragma acc loop gang reduction(+:s2) /* { dg-bogus "gang reduction on an orphan loop" } */ + for (i = 0; i < n; i++) + s2 = s2 + 2; + } + return s1 + s2; +} + +int +serial_combined (int n) +{ + int i, s1 = 0, s2 = 0; +#pragma acc serial loop gang reduction(+:s1) /* { dg-bogus "gang reduction on an orphan loop" } */ + /* { dg-warning "region contains gang partitioned code but is not gang partitioned" "" { target *-*-* } .-1 } */ + for (i = 0; i < n; i++) + s1 = s1 + 2; + +#pragma acc serial loop gang reduction(+:s2) /* { dg-bogus "gang reduction on an orphan loop" } */ + /* { dg-warning "region contains gang partitioned code but is not gang partitioned" "" { target *-*-* } .-1 } */ + for (i = 0; i < n; i++) + s2 = s2 + 2; + + return s1 + s2; +} + +int +parallel_combined (int n) +{ + int i, s1 = 0, s2 = 0; +#pragma acc parallel loop gang reduction(+:s1) /* { dg-bogus "gang reduction on an orphan loop" } */ + for (i = 0; i < n; i++) + s1 = s1 + 2; + +#pragma acc parallel loop gang reduction(+:s2) /* { dg-bogus "gang reduction on an orphan loop" } */ + for (i = 0; i < n; i++) + s2 = s2 + 2; + + return s1 + s2; +} + +int +kernels_combined (int n) +{ + int i, s1 = 0, s2 = 0; +#pragma acc kernels loop gang reduction(+:s1) /* { dg-bogus "gang reduction on an orphan loop" } */ + for (i = 0; i < n; i++) + s1 = s1 + 2; + +#pragma acc kernels loop gang reduction(+:s2) /* { dg-bogus "gang reduction on an orphan loop" } */ + for (i = 0; i < n; i++) + s2 = s2 + 2; + + return s1 + s2; +} + diff --git a/gcc/testsuite/gfortran.dg/goacc/orphan-reductions-2.f90 b/gcc/testsuite/gfortran.dg/goacc/orphan-reductions-2.f90 new file mode 100644 index 00000000000..6ad38039696 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/goacc/orphan-reductions-2.f90 @@ -0,0 +1,87 @@ +! Verify that the error message for gang reductions on orphaned OpenACC loops +! is not reported for non-orphaned loops. + +subroutine kernels + implicit none + + integer, parameter :: n = 100 + integer :: i, sum + sum = 0 + + !$acc kernels + !$acc loop gang reduction(+:sum) ! { dg-bogus "gang reduction on an orphan loop" } + do i = 1, n + sum = sum + 1 + end do + !$acc end kernels +end subroutine kernels + +subroutine parallel + implicit none + + integer, parameter :: n = 100 + integer :: i, sum + sum = 0 + + !$acc parallel + !$acc loop gang reduction(+:sum) ! { dg-bogus "gang reduction on an orphan loop" } + do i = 1, n + sum = sum + 1 + end do + !$acc end parallel +end subroutine parallel + +subroutine serial + implicit none + + integer, parameter :: n = 100 + integer :: i, sum + sum = 0 + + !$acc serial ! { dg-warning "region contains gang partitioned code but is not gang partitioned" } + !$acc loop gang reduction(+:sum) ! { dg-bogus "gang reduction on an orphan loop" } + do i = 1, n + sum = sum + 1 + end do + !$acc end serial +end subroutine serial + +subroutine kernels_combined + implicit none + + integer, parameter :: n = 100 + integer :: i, sum + sum = 0 + + !$acc kernels loop gang reduction(+:sum) ! { dg-bogus "gang reduction on an orphan loop" } + do i = 1, n + sum = sum + 1 + end do +end subroutine kernels_combined + +subroutine parallel_combined + implicit none + + integer, parameter :: n = 100 + integer :: i, sum + sum = 0 + + !$acc parallel loop gang reduction(+:sum) ! { dg-bogus "gang reduction on an orphan loop" } + do i = 1, n + sum = sum + 1 + end do +end subroutine parallel_combined + +subroutine serial_combined + implicit none + + integer, parameter :: n = 100 + integer :: i, sum + sum = 0 + + !$acc serial loop gang reduction(+:sum) ! { dg-bogus "gang reduction on an orphan loop" } + ! { dg-warning "region contains gang partitioned code but is not gang partitioned" "" { target *-*-* } .-1 } + do i = 1, n + sum = sum + 1 + end do +end subroutine serial_combined -- 2.17.1 --=-=-=--