From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: (qmail 20193 invoked by alias); 10 Dec 2014 10:10:32 -0000 Mailing-List: contact gcc-patches-help@gcc.gnu.org; run by ezmlm Precedence: bulk List-Id: List-Archive: List-Post: List-Help: Sender: gcc-patches-owner@gcc.gnu.org Received: (qmail 20181 invoked by uid 89); 10 Dec 2014 10:10:31 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-1.9 required=5.0 tests=AWL,BAYES_00,RCVD_IN_DNSWL_NONE autolearn=ham version=3.3.2 X-HELO: relay1.mentorg.com Received: from relay1.mentorg.com (HELO relay1.mentorg.com) (192.94.38.131) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with ESMTP; Wed, 10 Dec 2014 10:10:29 +0000 Received: from nat-ies.mentorg.com ([192.94.31.2] helo=SVR-IES-FEM-02.mgc.mentorg.com) by relay1.mentorg.com with esmtp id 1XyeDx-0006gP-L3 from Thomas_Schwinge@mentor.com ; Wed, 10 Dec 2014 02:10:26 -0800 Received: from feldtkeller.schwinge.homeip.net (137.202.0.76) by SVR-IES-FEM-02.mgc.mentorg.com (137.202.0.106) with Microsoft SMTP Server id 14.3.181.6; Wed, 10 Dec 2014 10:10:24 +0000 From: Thomas Schwinge To: Jakub Jelinek , Subject: Re: Nested OpenACC/OpenMP constructs (was: OpenACC GIMPLE_OACC_* -- or not?) In-Reply-To: <87sign7rgv.fsf@kepler.schwinge.homeip.net> References: <1383766943-8863-2-git-send-email-thomas@codesourcery.com> <1383766943-8863-3-git-send-email-thomas@codesourcery.com> <1383766943-8863-4-git-send-email-thomas@codesourcery.com> <1383766943-8863-5-git-send-email-thomas@codesourcery.com> <1383766943-8863-6-git-send-email-thomas@codesourcery.com> <1383766943-8863-7-git-send-email-thomas@codesourcery.com> <1383766943-8863-8-git-send-email-thomas@codesourcery.com> <1383766943-8863-9-git-send-email-thomas@codesourcery.com> <8761s5joir.fsf@schwinge.name> <87sihoczm0.fsf@kepler.schwinge.homeip.net> <20141112134502.GC5026@tucnak.redhat.com> <87sign7rgv.fsf@kepler.schwinge.homeip.net> User-Agent: Notmuch/0.9-101-g81dad07 (http://notmuchmail.org) Emacs/24.4.1 (i586-pc-linux-gnu) Date: Wed, 10 Dec 2014 10:10:00 -0000 Message-ID: <87iohj7r3o.fsf@kepler.schwinge.homeip.net> MIME-Version: 1.0 Content-Type: multipart/signed; boundary="=-=-="; micalg=pgp-sha1; protocol="application/pgp-signature" X-SW-Source: 2014-12/txt/msg00892.txt.bz2 --=-=-= Content-Type: text/plain; charset=utf-8 Content-Transfer-Encoding: quoted-printable Content-length: 8463 Hi Jakub! On Wed, 10 Dec 2014 11:02:24 +0100, I wrote: > OpenACC: Rework nested constructs checking. >=20=20=20=20=20 > gcc/ > * omp-low.c (scan_omp_target): Remove taskreg_nesting_level and > target_nesting_level assertions. > (check_omp_nesting_restrictions): Rework OpenACC constructs > handling. Update and extend the relevant test cases. Regarding the check_omp_nesting_restrictions rework: > --- gcc/omp-low.c > +++ gcc/omp-low.c > @@ -2706,46 +2700,26 @@ scan_omp_teams (gimple stmt, omp_context *outer_c= tx) > scan_omp (gimple_omp_body_ptr (stmt), ctx); > } >=20=20 > -/* Check OpenMP nesting restrictions. */ > +/* Check nesting restrictions. */ > static bool > check_omp_nesting_restrictions (gimple stmt, omp_context *ctx) > { > - /* TODO: While the OpenACC specification does allow for certain kinds = of > - nesting, we don't support many of these yet. */ > - if (is_gimple_omp (stmt) > - && is_gimple_omp_oacc_specifically (stmt)) > + /* TODO: Some OpenACC/OpenMP nesting should be allowed. */ > + > + /* No nesting of non-OpenACC STMT (that is, an OpenMP one, or a GOMP b= uiltin) > + inside an OpenACC CTX. */ > + if (!(is_gimple_omp (stmt) > + && is_gimple_omp_oacc_specifically (stmt))) > { > - /* Regular handling of OpenACC loop constructs. */ > - if (gimple_code (stmt) =3D=3D GIMPLE_OMP_FOR > - && gimple_omp_for_kind (stmt) =3D=3D GF_OMP_FOR_KIND_OACC_LOOP) > - goto cont; > - /* No nesting of OpenACC STMT inside any OpenACC or OpenMP CTX dif= ferent > - from an OpenACC data construct. */ > - for (omp_context *ctx_ =3D ctx; ctx_ !=3D NULL; ctx_ =3D ctx_->out= er) > - if (is_gimple_omp (ctx_->stmt) > - && !(gimple_code (ctx_->stmt) =3D=3D GIMPLE_OMP_TARGET > - && (gimple_omp_target_kind (ctx_->stmt) > - =3D=3D GF_OMP_TARGET_KIND_OACC_DATA))) > - { > - error_at (gimple_location (stmt), > - "may not be nested"); > - return false; > - } > - } > - else > - { > - /* No nesting of non-OpenACC STMT (that is, an OpenMP one, or a GO= MP > - builtin) inside any OpenACC CTX. */ > for (omp_context *ctx_ =3D ctx; ctx_ !=3D NULL; ctx_ =3D ctx_->out= er) > if (is_gimple_omp (ctx_->stmt) > && is_gimple_omp_oacc_specifically (ctx_->stmt)) > { > error_at (gimple_location (stmt), > - "may not be nested"); > + "non-OpenACC construct inside of OpenACC region"); > return false; > } > } > - cont: >=20=20 > if (ctx !=3D NULL) > { > @@ -3003,20 +2977,74 @@ check_omp_nesting_restrictions (gimple stmt, omp_= context *ctx) > break; > case GIMPLE_OMP_TARGET: > for (; ctx !=3D NULL; ctx =3D ctx->outer) > - if (gimple_code (ctx->stmt) =3D=3D GIMPLE_OMP_TARGET > - && gimple_omp_target_kind (ctx->stmt) =3D=3D GF_OMP_TARGET_KIND_REG= ION) > - { > - const char *name; > - switch (gimple_omp_target_kind (stmt)) > - { > - case GF_OMP_TARGET_KIND_REGION: name =3D "target"; break; > - case GF_OMP_TARGET_KIND_DATA: name =3D "target data"; break; > - case GF_OMP_TARGET_KIND_UPDATE: name =3D "target update"; break; > - default: gcc_unreachable (); > - } > - warning_at (gimple_location (stmt), 0, > - "%s construct inside of target region", name); > - } > + { > + if (gimple_code (ctx->stmt) !=3D GIMPLE_OMP_TARGET) > + { > + if (is_gimple_omp (stmt) > + && is_gimple_omp_oacc_specifically (stmt) > + && is_gimple_omp (ctx->stmt)) > + { > + error_at (gimple_location (stmt), > + "OpenACC construct inside of non-OpenACC region"); > + return false; > + } > + continue; > + } > + > + const char *stmt_name, *ctx_stmt_name; > + switch (gimple_omp_target_kind (stmt)) > + { > + case GF_OMP_TARGET_KIND_REGION: stmt_name =3D "target"; break; > + case GF_OMP_TARGET_KIND_DATA: stmt_name =3D "target data"; break; > + case GF_OMP_TARGET_KIND_UPDATE: stmt_name =3D "target update"; brea= k; > + case GF_OMP_TARGET_KIND_OACC_PARALLEL: stmt_name =3D "parallel"; br= eak; > + case GF_OMP_TARGET_KIND_OACC_KERNELS: stmt_name =3D "kernels"; brea= k; > + case GF_OMP_TARGET_KIND_OACC_DATA: stmt_name =3D "data"; break; > + case GF_OMP_TARGET_KIND_OACC_UPDATE: stmt_name =3D "update"; break; > + case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA: stmt_name =3D "enter/= exit data"; break; > + default: gcc_unreachable (); > + } > + switch (gimple_omp_target_kind (ctx->stmt)) > + { > + case GF_OMP_TARGET_KIND_REGION: ctx_stmt_name =3D "target"; break; > + case GF_OMP_TARGET_KIND_DATA: ctx_stmt_name =3D "target data"; brea= k; > + case GF_OMP_TARGET_KIND_OACC_PARALLEL: ctx_stmt_name =3D "parallel"= ; break; > + case GF_OMP_TARGET_KIND_OACC_KERNELS: ctx_stmt_name =3D "kernels"; = break; > + case GF_OMP_TARGET_KIND_OACC_DATA: ctx_stmt_name =3D "data"; break; > + default: gcc_unreachable (); > + } > + > + /* OpenACC/OpenMP mismatch? */ > + if (is_gimple_omp_oacc_specifically (stmt) > + !=3D is_gimple_omp_oacc_specifically (ctx->stmt)) > + { > + error_at (gimple_location (stmt), > + "%s %s construct inside of %s %s region", > + (is_gimple_omp_oacc_specifically (stmt) > + ? "OpenACC" : "OpenMP"), stmt_name, > + (is_gimple_omp_oacc_specifically (ctx->stmt) > + ? "OpenACC" : "OpenMP"), ctx_stmt_name); > + return false; > + } > + if (is_gimple_omp_offloaded (ctx->stmt)) > + { > + /* No GIMPLE_OMP_TARGET inside offloaded OpenACC CTX. */ > + if (is_gimple_omp_oacc_specifically (ctx->stmt)) > + { > + error_at (gimple_location (stmt), > + "%s construct inside of %s region", > + stmt_name, ctx_stmt_name); > + return false; > + } > + else > + { > + gcc_assert (!is_gimple_omp_oacc_specifically (stmt)); > + warning_at (gimple_location (stmt), 0, > + "%s construct inside of %s region", > + stmt_name, ctx_stmt_name); > + } > + } > + } > break; > default: > break; For guarding against OpenMP target regressions, I used the following tests -- OK to commit to trunk? gcc/testsuite/c-c++-common/gomp/nesting-1.c | 77 ++++++++++++++++++++= ++++ gcc/testsuite/c-c++-common/gomp/nesting-warn-1.c | 23 +++++++ 2 files changed, 100 insertions(+) diff --git gcc/testsuite/c-c++-common/gomp/nesting-1.c gcc/testsuite/c-c++-= common/gomp/nesting-1.c new file mode 100644 index 0000000..70b3e8d --- /dev/null +++ gcc/testsuite/c-c++-common/gomp/nesting-1.c @@ -0,0 +1,77 @@ +void +f_omp_parallel (void) +{ +#pragma omp parallel + { + int i; + +#pragma omp parallel + ; + +#pragma omp target + ; + +#pragma omp target data + ; + +#pragma omp target update to(i) + +#pragma omp target data + { +#pragma omp parallel + ; + +#pragma omp target + ; + +#pragma omp target data + ; + +#pragma omp target update to(i) + } + } +} + +void +f_omp_target (void) +{ +#pragma omp target + { +#pragma omp parallel + ; + } +} + +void +f_omp_target_data (void) +{ +#pragma omp target data + { + int i; + +#pragma omp parallel + ; + +#pragma omp target + ; + +#pragma omp target data + ; + +#pragma omp target update to(i) + +#pragma omp target data + { +#pragma omp parallel + ; + +#pragma omp target + ; + +#pragma omp target data + ; + +#pragma omp target update to(i) + } + } +} diff --git gcc/testsuite/c-c++-common/gomp/nesting-warn-1.c gcc/testsuite/c= -c++-common/gomp/nesting-warn-1.c new file mode 100644 index 0000000..6a65699 --- /dev/null +++ gcc/testsuite/c-c++-common/gomp/nesting-warn-1.c @@ -0,0 +1,23 @@ +void +f_omp_target (void) +{ +#pragma omp target + { + int i; + +#pragma omp target /* { dg-warning "target construct inside of target regi= on" } */ + ; +#pragma omp target data /* { dg-warning "target data construct inside of t= arget region" } */ + ; +#pragma omp target update to(i) /* { dg-warning "target update construct i= nside of target region" } */ + +#pragma omp parallel + { +#pragma omp target /* { dg-warning "target construct inside of target regi= on" } */ + ; +#pragma omp target data /* { dg-warning "target data construct inside of t= arget region" } */ + ; +#pragma omp target update to(i) /* { dg-warning "target update construct i= nside of target region" } */ + } + } +} Gr=C3=BC=C3=9Fe, Thomas --=-=-= Content-Type: application/pgp-signature; name="signature.asc" Content-length: 472 -----BEGIN PGP SIGNATURE----- Version: GnuPG v1 iQEcBAEBAgAGBQJUiBwLAAoJEK3/DN1sMFFtu20IAJezq6L2vS0EomQY4ABLgTtC Thfk7zTkTYe5zJe2zFlFmunCB7WQASqC0vNBclb4PQbmG28d9rYPlL3wnAKmp6Cn Y+EmK7DEUjI9gtLmu42nEo4+gQ6Sh3V6264HlahDytutGxiPLcaezMle8jO1mNRJ b1HwLPfLl90a2J3F5C2NfdDKoqiM2G00Fjz5sV1DYyLJYnOdqOsosyJ0ZMIe2f1v n2L1SHktqaW6qEjHtZ+DW6RWNTOXK12zensv5VLQrstHGIfRpZM9LS3S2GD3cRQn z/FjJBICWaJ2KE5tsFKKNN5m5iWXqHtaO/2eA3uZvcB6AoZqJYvRW68edAUN/mc= =yYRY -----END PGP SIGNATURE----- --=-=-=--