From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: (qmail 52539 invoked by alias); 29 Apr 2015 11:52:46 -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 52525 invoked by uid 89); 29 Apr 2015 11:52:45 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-0.9 required=5.0 tests=AWL,BAYES_50,RCVD_IN_DNSWL_LOW,SPF_PASS 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, 29 Apr 2015 11:52:44 +0000 Received: from nat-ies.mentorg.com ([192.94.31.2] helo=SVR-IES-FEM-01.mgc.mentorg.com) by relay1.mentorg.com with esmtp id 1YnQXg-0007Ou-2U from Thomas_Schwinge@mentor.com ; Wed, 29 Apr 2015 04:52:40 -0700 Received: from feldtkeller.schwinge.homeip.net (137.202.0.76) by SVR-IES-FEM-01.mgc.mentorg.com (137.202.0.104) with Microsoft SMTP Server id 14.3.224.2; Wed, 29 Apr 2015 12:52:38 +0100 From: Thomas Schwinge To: Jakub Jelinek CC: , Kirill Yukhin , Richard Henderson Subject: Re: [gomp4.1] Initial support for some OpenMP 4.1 construct parsing In-Reply-To: <20150429111406.GE1751@tucnak.redhat.com> References: <20150429111406.GE1751@tucnak.redhat.com> User-Agent: Notmuch/0.9-101-g81dad07 (http://notmuchmail.org) Emacs/24.3.1 (x86_64-pc-linux-gnu) Date: Wed, 29 Apr 2015 11:55:00 -0000 Message-ID: <874mnzrw1z.fsf@schwinge.name> MIME-Version: 1.0 Content-Type: multipart/signed; boundary="=-=-="; micalg=pgp-sha1; protocol="application/pgp-signature" X-SW-Source: 2015-04/txt/msg01863.txt.bz2 --=-=-= Content-Type: text/plain; charset=utf-8 Content-Transfer-Encoding: quoted-printable Content-length: 4576 Hi Jakub! (I have not yet read any version of the OpenMP 4.1 standard.) On Wed, 29 Apr 2015 13:14:06 +0200, Jakub Jelinek wrote: > --- gcc/tree.def.jj 2015-04-29 10:58:01.663745452 +0200 > +++ gcc/tree.def 2015-04-29 11:34:18.293302684 +0200 > +/* OpenMP - #pragma omp target enter data [clause1 ... clauseN] > + Operand 0: OMP_TARGET_ENTER_DATA_CLAUSES: List of clauses. */ > +DEFTREECODE (OMP_TARGET_ENTER_DATA, "omp_target_enter_data", tcc_stateme= nt, 1) > + > +/* OpenMP - #pragma omp target exit data [clause1 ... clauseN] > + Operand 0: OMP_TARGET_EXIT_DATA_CLAUSES: List of clauses. */ > +DEFTREECODE (OMP_TARGET_EXIT_DATA, "omp_target_exit_data", tcc_statement= , 1) ;-) Heh, OpenMP catching up with OpenACC? > --- gcc/c/c-parser.c.jj 2015-04-29 10:59:20.760504260 +0200 > +++ gcc/c/c-parser.c 2015-04-29 11:03:04.641991095 +0200 > map ( variable-list ) >=20=20 > map-kind: > - alloc | to | from | tofrom */ > + alloc | to | from | tofrom > + > + OpenMP 4.1: > + map-kind: > + alloc | to | from | tofrom | delete > + > + map ( always [,] map-kind: variable-list ) */ "Funnily", OpenACC 2.5 is said to be dropping the "always" semantics that OpenMP 4.1 is now adding... That is, OpenACC 2.5's "copy" will then be the same as "present_or_copy", and so on. > static tree > c_parser_omp_target_data (location_t loc, c_parser *parser) > { > - tree stmt =3D make_node (OMP_TARGET_DATA); > - TREE_TYPE (stmt) =3D void_type_node; > - > - OMP_TARGET_DATA_CLAUSES (stmt) > + tree clauses > =3D c_parser_omp_all_clauses (parser, OMP_TARGET_DATA_CLAUSE_MASK, > "#pragma omp target data"); > + if (find_omp_clause (clauses, OMP_CLAUSE_MAP) =3D=3D NULL_TREE) > + { > + error_at (loc, > + "%<#pragma omp target data%> must contain at least one " > + "% clause"); > + return NULL_TREE; > + } > + > + tree stmt =3D make_node (OMP_TARGET_DATA); > + TREE_TYPE (stmt) =3D void_type_node; > + OMP_TARGET_DATA_CLAUSES (stmt) =3D clauses; Even if it doesn't make a lot of sense, my reading of OpenMP 4.0 has been that a target data construct without any map clauses is still valid. (But I have not verified that now.) > --- gcc/omp-low.c.jj 2015-04-29 10:58:01.489748182 +0200 > +++ gcc/omp-low.c 2015-04-29 11:03:04.646991017 +0200 > @@ -8994,6 +9017,11 @@ expand_omp_target (struct omp_region *re > case GF_OMP_TARGET_KIND_UPDATE: > start_ix =3D BUILT_IN_GOMP_TARGET_UPDATE; > break; > + case GF_OMP_TARGET_KIND_ENTER_DATA: > + case GF_OMP_TARGET_KIND_EXIT_DATA: > + /* FIXME */ > + start_ix =3D BUILT_IN_GOMP_TARGET_UPDATE; > + break; Actually, I've also wondered (but never verified) whether all of the OpenACC enter, exit, and update constructs can be implemented via the same libgomp API. > @@ -11488,6 +11520,9 @@ lower_omp_target (gimple_stmt_iterator * > default: > gcc_unreachable (); > } > + /* FIXME: Temporary hack. */ > + if (talign_shift =3D=3D 3) > + tkind &=3D ~GOMP_MAP_FLAG_FORCE; > gcc_checking_assert (tkind > < (HOST_WIDE_INT_C (1U) << talign_shift)); > talign =3D ceil_log2 (talign); Assuming you'll be changing the relevant functions' APIs, and assigning new ABI versions, don't forget to also drop the unused offload_table formal parameter (assuming that it remains unused). > --- gcc/tree-pretty-print.c.jj 2015-04-29 10:58:01.663745452 +0200 > +++ gcc/tree-pretty-print.c 2015-04-29 11:03:04.648990986 +0200 > @@ -550,22 +560,22 @@ dump_omp_clause (pretty_printer *pp, tre > pp_string (pp, "tofrom"); > break; > case GOMP_MAP_FORCE_ALLOC: > - pp_string (pp, "force_alloc"); > + pp_string (pp, "always,alloc"); > break; > case GOMP_MAP_FORCE_TO: > - pp_string (pp, "force_to"); > + pp_string (pp, "always,to"); > break; > case GOMP_MAP_FORCE_FROM: > - pp_string (pp, "force_from"); > + pp_string (pp, "always,from"); > break; > case GOMP_MAP_FORCE_TOFROM: > - pp_string (pp, "force_tofrom"); > + pp_string (pp, "always,tofrom"); > break; > case GOMP_MAP_FORCE_PRESENT: > pp_string (pp, "force_present"); > break; > case GOMP_MAP_FORCE_DEALLOC: > - pp_string (pp, "force_dealloc"); > + pp_string (pp, "always,delete"); > break; > case GOMP_MAP_FORCE_DEVICEPTR: > pp_string (pp, "force_deviceptr"); Makes some sense to me to use the same "always,*" syntax also for the other "force_*" ones, given that these are artificial ("non-OpenACC") descriptions anyway. Gr=C3=BC=C3=9Fe, Thomas --=-=-= Content-Type: application/pgp-signature Content-length: 472 -----BEGIN PGP SIGNATURE----- Version: GnuPG v1 iQEcBAEBAgAGBQJVQMX5AAoJEPoxNhtoi6COrkYH/1YEWs1IztkuK8mDPd8XwUFY 6kTp/jZ6+UdLMT3ehlNyqQGE09oI4C6gXa2nY93/AeKQgqBx2nwjvt48UVfUw9bG LaD16LIHED/aLlX7oV7C7Z/Vx+uOnT/8AA4EhewrcDUz/1wXIQC9yHrCjMbvFecJ pkHja92ARSqbfd9uJQpAuY+gQAvvvKtHyJfh7tamCJ3idDcIPDXgZEkvGf8+8OBh 0SU4nkDpOxuKsSYHD12ZGU37h+hOuBIK6WtK/aXbnmvs6ZvOKdGcy4rifcv52+9l gkOOOMrzaOBOUIaUSDjx5FlaAmhHebAuVJ/aIuxSZMutTiBi75Q1qXCvbGy63LM= =KKH1 -----END PGP SIGNATURE----- --=-=-=--