From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: from esa2.mentor.iphmx.com (esa2.mentor.iphmx.com [68.232.141.98]) by sourceware.org (Postfix) with ESMTPS id 3A8E23857C53; Tue, 7 Jul 2020 10:43:05 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.3.2 sourceware.org 3A8E23857C53 Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=Thomas_Schwinge@mentor.com IronPort-SDR: LFcEOgqJU3b7QSKMrA47fTW5Qp+3MCOrrtmQTGji9eJqhH0qKHvdU/l45a/nwm9j2wnMSSZekN ER4xt1Ofsq5gFh8xTSS8akI5GxmwpQ0Yeh9U04/QECmN9Xn0Rtp+vSmS5/nAw2mWgASzD0/sld 2U6EkNAtooy9qdY8jK47BnaEWH8BgX3YEYBpaNBMIzh5tsw6j/qbmGple2bnqzcVIpHB665kiP SbEfqB2lSthXQz5x1cpy8NOcgvupyEt4qPLYZkoTt0vxlWh1CHAL8sC5LB3RBDP0chZvKbQbs9 Llw= X-IronPort-AV: E=Sophos;i="5.75,323,1589270400"; d="scan'208";a="50664970" Received: from orw-gwy-02-in.mentorg.com ([192.94.38.167]) by esa2.mentor.iphmx.com with ESMTP; 07 Jul 2020 02:43:03 -0800 IronPort-SDR: bqc/YSPXnx+uYw2aDXlQqvm17KcoVcZ6aHroZg9TYku+/0bYSPd6HWVAjl06gPZlbS2CyZLa9Y ARu73QZJDLALZZdedUq4MNeJdLFkNmDu3b3LOiIyOQBvp8BgNETAzIYzvI6Exuyjzj63f7D2bg 39wGJ16YjhvicBvnWZUI0+O2jjmCILDkvHRuUJds8ISThquypa/SZcSVBA6en5RMjO74/2t9Dn ET69DTQfh+UNxYJyhK9cbP8Vi1pg2SqhvKoPyAjOgQHtibRjwc9hkMJ7bk3hhLXyqxE3yzF98U owA= From: Thomas Schwinge To: Frederik Harwath , 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> User-Agent: Notmuch/0.29.1+93~g67ed7df (https://notmuchmail.org) Emacs/25.2.2 (x86_64-pc-linux-gnu) Date: Tue, 7 Jul 2020 12:42:53 +0200 Message-ID: MIME-Version: 1.0 Content-Type: text/plain; charset="utf-8" Content-Transfer-Encoding: quoted-printable X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: SVR-IES-MBX-04.mgc.mentorg.com (139.181.222.4) To SVR-IES-MBX-03.mgc.mentorg.com (139.181.222.3) X-Spam-Status: No, score=-13.6 required=5.0 tests=BAYES_00, GIT_PATCH_0, HEADER_FROM_DIFFERENT_DOMAINS, KAM_DMARC_STATUS, KAM_LOTSOFHASH, 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: gcc-patches@gcc.gnu.org X-Mailman-Version: 2.1.29 Precedence: list List-Id: Gcc-patches mailing list List-Unsubscribe: , List-Archive: List-Post: List-Help: List-Subscribe: , X-List-Received-Date: Tue, 07 Jul 2020 10:43:07 -0000 Hi Frederik! (CC added, for everything touching gfortran.) On 2020-07-07T10:52:08+0200, Frederik Harwath w= rote: > This patch fixes the check for reductions on orphaned gang loops This is the "Make OpenACC orphan gang reductions errors" functionality originally added in gomp-4_0-branch r247461. > the Fortran frontend which (in contrast to the C, C++ frontends) > erroneously rejects reductions on gang loops that are contained in > "kernels" constructs and which hence are not orphaned. > > 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"). Remember that 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. Or the other way round: a 'loop' construct is orphaned if it appears inside a 'routine' region, right? > The patch has been tested by running the GCC and libgomp testsuites. > The latter tests ran with offloading to nvptx although that should not > be important here unless there was some very subtle reason for > forbidding the gang reductions on kernels loops. As expect, there seems > to be no such reason, i.e. I observed no regressions with the patch. Note that the aforementioned gomp-4_0-branch r247461, openacc-gcc-7-branch commit 0554f9f79325960c72166327d442a553cd35bad9, and openacc-gcc-8-branch commit 65dd9cf3b3c45d64d72967df1e4a54778cb4e35f still do contain the appropriate 'kernels' handling. Just in openacc-gcc-9-branch commit 533beb2ec19f8486e4b1b645a153746f96b41f04 this got (a) mixed together with a bunch of other, unrelated changes ("Various OpenACC reduction enhancements"), and (b) the 'kernels' handling got removed. Julian (Git author), or Kwok (Git committer), do you remember any rationale for that? Later, this then got picked into devel/omp/gcc-9 commit 3fa4bb72dcb3b9171952a0eca5310bb8811d5ffd, and devel/omp/gcc-10 commit 6b3e1f7f05cd360bbd356b3f78511aa2ec3f40c3. > Can I include the patch in OG10? Unless Julian/Kwok speak up soon: OK, thanks. Reviewed-by: Thomas Schwinge 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.) Gr=C3=BC=C3=9Fe Thomas > From 7320635211fff3a773beb0de1914dbfcc317ab37 Mon Sep 17 00:00:00 2001 > From: Frederik Harwath > Date: Tue, 7 Jul 2020 10:41:21 +0200 > Subject: [PATCH] libgomp, Fortran: Fix OpenACC "gang reduction on an orph= an > 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-07 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 error message is not emitted for > non-orphaned loops. > > * c-c++-common/goacc/orphan-reductions-2.c: Likewise for C and C++. > --- > gcc/fortran/openmp.c | 13 +++- > .../c-c++-common/goacc/orphan-reductions-2.c | 69 +++++++++++++++++++ > .../gfortran.dg/goacc/orphan-reductions-2.f90 | 58 ++++++++++++++++ > 3 files changed, 137 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.f= 90 > > diff --git a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c > index 28408c4c99a..83c498112a8 100644 > --- a/gcc/fortran/openmp.c > +++ b/gcc/fortran/openmp.c > @@ -5926,9 +5926,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 =3D=3D EXEC_OACC_KERNELS || code->op =3D=3D 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 > @@ -6222,7 +6229,7 @@ resolve_oacc_loop_blocks (gfc_code *code) > for (c =3D omp_current_ctx; c; c =3D c->previous) > if (!oacc_is_loop (c->code)) > break; > - if (c =3D=3D NULL || !oacc_is_parallel_or_serial (c->code)) > + if (c =3D=3D NULL || !oacc_is_compute_construct (c->code)) > gfc_error ("gang reduction on an orphan loop at %L", &code->loc); > } > > 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..2b651fd2b9f > --- /dev/null > +++ b/gcc/testsuite/c-c++-common/goacc/orphan-reductions-2.c > @@ -0,0 +1,69 @@ > +/* 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 =3D 0, s2 =3D 0; > +#pragma acc kernels > + { > +#pragma acc loop gang reduction(+:s1) /* { dg-bogus "gang reduction on a= n orphan loop" } */ > + for (i =3D 0; i < n; i++) > + s1 =3D s1 + 2; > + > +#pragma acc loop gang reduction(+:s2) /* { dg-bogus "gang reduction on a= n orphan loop" } */ > + for (i =3D 0; i < n; i++) > + s2 =3D s2 + 2; > + } > + return s1 + s2; > +} > + > +int > +parallel (int n) > +{ > + int i, s1 =3D 0, s2 =3D 0; > +#pragma acc parallel > + { > +#pragma acc loop gang reduction(+:s1) /* { dg-bogus "gang reduction on a= n orphan loop" } */ > + for (i =3D 0; i < n; i++) > + s1 =3D s1 + 2; > + > +#pragma acc loop gang reduction(+:s2) /* { dg-bogus "gang reduction on a= n orphan loop" } */ > + for (i =3D 0; i < n; i++) > + s2 =3D s2 + 2; > + } > + return s1 + s2; > +} > + > +int > +parallel_combined (int n) > +{ > + int i, s1 =3D 0, s2 =3D 0; > +#pragma acc parallel loop gang reduction(+:s1) /* { dg-bogus "gang reduc= tion on an orphan loop" } */ > + for (i =3D 0; i < n; i++) > + s1 =3D s1 + 2; > + > +#pragma acc parallel loop gang reduction(+:s2) /* { dg-bogus "gang reduc= tion on an orphan loop" } */ > + for (i =3D 0; i < n; i++) > + s2 =3D s2 + 2; > + > + return s1 + s2; > +} > + > +int > +kernels_combined (int n) > +{ > + int i, s1 =3D 0, s2 =3D 0; > +#pragma acc kernels loop gang reduction(+:s1) /* { dg-bogus "gang reduct= ion on an orphan loop" } */ > + for (i =3D 0; i < n; i++) > + s1 =3D s1 + 2; > + > +#pragma acc kernels loop gang reduction(+:s2) /* { dg-bogus "gang reduct= ion on an orphan loop" } */ > + for (i =3D 0; i < n; i++) > + s2 =3D s2 + 2; > + > + return s1 + s2; > +} > + > diff --git a/gcc/testsuite/gfortran.dg/goacc/orphan-reductions-2.f90 b/gc= c/testsuite/gfortran.dg/goacc/orphan-reductions-2.f90 > new file mode 100644 > index 00000000000..13887a059fe > --- /dev/null > +++ b/gcc/testsuite/gfortran.dg/goacc/orphan-reductions-2.f90 > @@ -0,0 +1,58 @@ > +! 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 =3D 100 > + integer :: i, sum > + sum =3D 0 > + > + !$acc kernels > + !$acc loop gang reduction(+:sum) ! { dg-bogus "gang reduction on an or= phan loop" } > + do i =3D 1, n > + sum =3D sum + 1 > + end do > + !$acc end kernels > +end subroutine kernels > + > +subroutine parallel > + implicit none > + > + integer, parameter :: n =3D 100 > + integer :: i, sum > + sum =3D 0 > + > + !$acc parallel > + !$acc loop gang reduction(+:sum) ! { dg-bogus "gang reduction on an or= phan loop" } > + do i =3D 1, n > + sum =3D sum + 1 > + end do > + !$acc end parallel > +end subroutine parallel > + > +subroutine kernels_combined > + implicit none > + > + integer, parameter :: n =3D 100 > + integer :: i, sum > + sum =3D 0 > + > + !$acc kernels loop gang reduction(+:sum) ! { dg-bogus "gang reduction = on an orphan loop" } > + do i =3D 1, n > + sum =3D sum + 1 > + end do > +end subroutine kernels_combined > + > +subroutine parallel_combined > + implicit none > + > + integer, parameter :: n =3D 100 > + integer :: i, sum > + sum =3D 0 > + > + !$acc parallel loop gang reduction(+:sum) ! { dg-bogus "gang reduction= on an orphan loop" } > + do i =3D 1, n > + sum =3D sum + 1 > + end do > +end subroutine parallel_combined ----------------- Mentor Graphics (Deutschland) GmbH, Arnulfstra=C3=9Fe 201, 80634 M=C3=BCnch= en / Germany Registergericht M=C3=BCnchen HRB 106955, Gesch=C3=A4ftsf=C3=BChrer: Thomas = Heurung, Alexander Walter