From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: (qmail 16168 invoked by alias); 11 Feb 2014 16:51:35 -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 16147 invoked by uid 89); 11 Feb 2014 16:51:35 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-2.1 required=5.0 tests=AWL,BAYES_00 autolearn=ham version=3.3.2 X-Spam-User: qpsmtpd, 2 recipients 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; Tue, 11 Feb 2014 16:51:33 +0000 Received: from svr-orw-exc-10.mgc.mentorg.com ([147.34.98.58]) by relay1.mentorg.com with esmtp id 1WDGYS-0005Ln-1w from Thomas_Schwinge@mentor.com ; Tue, 11 Feb 2014 08:51:28 -0800 Received: from SVR-IES-FEM-03.mgc.mentorg.com ([137.202.0.108]) by SVR-ORW-EXC-10.mgc.mentorg.com with Microsoft SMTPSVC(6.0.3790.4675); Tue, 11 Feb 2014 08:51:28 -0800 Received: from feldtkeller.schwinge.homeip.net (137.202.0.76) by SVR-IES-FEM-03.mgc.mentorg.com (137.202.0.108) with Microsoft SMTP Server id 14.2.247.3; Tue, 11 Feb 2014 16:51:25 +0000 From: Thomas Schwinge To: Ilmir Usmanov CC: Evgeny Gavrin , GarbuzovViacheslav , Dmitri Botcharnikov , , , Subject: Re: [PATCH 4/6] [GOMP4] OpenACC 1.0+ support in fortran front-end In-Reply-To: <52EB85F7.5020807@samsung.com> References: <52E158EF.9050009@samsung.com> <877g9pqmt2.fsf@schwinge.name> <52E65B24.9070403@samsung.com> <87iot5pgqb.fsf@schwinge.name> <52EB8437.3060602@samsung.com> <52EB84C3.4010407@samsung.com> <52EB8537.4090708@samsung.com> <52EB8596.2010808@samsung.com> <52EB85F7.5020807@samsung.com> User-Agent: Notmuch/0.9-101-g81dad07 (http://notmuchmail.org) Emacs/23.4.1 (x86_64-pc-linux-gnu) Date: Tue, 11 Feb 2014 16:51:00 -0000 Message-ID: <8761olk2yk.fsf@schwinge.name> MIME-Version: 1.0 Content-Type: multipart/signed; boundary="=-=-="; micalg=pgp-sha1; protocol="application/pgp-signature" X-SW-Source: 2014-02/txt/msg00685.txt.bz2 --=-=-= Content-Type: text/plain; charset=utf-8 Content-Transfer-Encoding: quoted-printable Content-length: 13873 Hi! On Fri, 31 Jan 2014 15:16:07 +0400, Ilmir Usmanov w= rote: > OpenACC 1.0 support -- GENERIC nodes and gimplify stubs. Thanks! This one is nearly ready for commit, and can then go in already, while the Fortran front end patches are still being dicussed. :-) Please merge =C2=BBOpenACC 1.0 support -- documentation=C2=AB into this pat= ch. In gcc/doc/generic.texi, please also add an entry for OACC_DECLARE, which is currently missing. Ideally, gcc/doc/generic.texi should also be updated for the OpenACC clauses (as well as recently added OpenMP ones), but as that one's outdated already, this is not a prerequisite. For ChangeLog files updates (on gomp-4_0-branch, use the respective ChangeLog.gomp files, by the way), should just you be listed as the author, or also your colleagues? > --- a/gcc/gimplify.c > +++ b/gcc/gimplify.c > @@ -6157,6 +6166,20 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_se= q *pre_p, > remove =3D true; > break; >=20=20 > + case OMP_CLAUSE_HOST: > + case [...] > + case OMP_CLAUSE_VECTOR_LENGTH: > case OMP_CLAUSE_NOWAIT: > case OMP_CLAUSE_ORDERED: > case OMP_CLAUSE_UNTIED: Indentation (one tab instead of two spaces). > @@ -6476,6 +6499,23 @@ gimplify_adjust_omp_clauses (tree *list_p) > } > break; >=20=20 > + case OMP_CLAUSE_HOST: > + case [...] > + case OMP_CLAUSE_VECTOR_LENGTH: > + sorry ("Clause not supported yet"); > + break; > + > case OMP_CLAUSE_REDUCTION: > case OMP_CLAUSE_COPYIN: > case OMP_CLAUSE_COPYPRIVATE: Indentation. Also, shouldn't the sorry be moved to gimplify_scan_omp_clauses, and a =C2=BBremove =3D true=C2=AB be added in there, and gimplify_adjust_omp_clau= ses then just contain a gcc_unreachable (or, move the new =C2=BBcase OMP_CLAUSE_*=C2= =AB next to the existing default: that already contains gcc_unreachable)? > @@ -7988,6 +8028,19 @@ gimplify_expr (tree *expr_p, gimple_seq *pre_p, gi= mple_seq *post_p, > ret =3D GS_ALL_DONE; > break; >=20=20 > + case OACC_KERNELS: > + case OACC_DATA: > + case OACC_CACHE: > + case OACC_WAIT: > + case OACC_HOST_DATA: > + case OACC_DECLARE: > + case OACC_UPDATE: > + case OACC_ENTER_DATA: > + case OACC_EXIT_DATA: > + sorry ("directive not yet implemented"); > + ret =3D GS_ALL_DONE; > + break; > + > case OMP_PARALLEL: > gimplify_omp_parallel (expr_p, pre_p); > ret =3D GS_ALL_DONE; Indentation. Further down in gimplify_expr, shouldn't these new OACC_* codes also be added to the =C2=BBThese expressions should already be in gimple IR form=C2= =AB assert? > --- a/gcc/omp-low.c > +++ b/gcc/omp-low.c > @@ -1491,6 +1491,18 @@ fixup_child_record_type (omp_context *ctx) > TREE_TYPE (ctx->receiver_decl) =3D build_pointer_type (type); > } >=20=20 > +static bool > +gimple_code_is_oacc (const_gimple g) > +{ > + switch (gimple_code (g)) > + { > + case GIMPLE_OACC_PARALLEL: > + return true; > + default: > + return false; > + } > +} > + Eventually, this will probably end up next to CASE_GIMPLE_OP/is_gimple_op in gimple.h (or the latter be reworked to be able to ask for is_omp vs. is_oacc vs. is_omp_or_oacc), but it's fine to do that once we actually need it in files other than just omp-low.c, and once we support more GIMPLE_OACC_* codes. > @@ -1710,17 +1732,34 @@ scan_sharing_clauses (tree clauses, omp_context *= ctx) > + case OMP_CLAUSE_HOST: > + case OMP_CLAUSE_OACC_DEVICE: > + case OMP_CLAUSE_DEVICE_RESIDENT: > + case OMP_CLAUSE_USE_DEVICE: > + case OMP_CLAUSE_ASYNC: > + case OMP_CLAUSE_GANG: > + case OMP_CLAUSE_WAIT: > + case OMP_NO_CLAUSE_CACHE: > + case OMP_CLAUSE_INDEPENDENT: > + case OMP_CLAUSE_WORKER: > + case OMP_CLAUSE_VECTOR: > + case OMP_CLAUSE_NUM_GANGS: > + case OMP_CLAUSE_NUM_WORKERS: > + case OMP_CLAUSE_VECTOR_LENGTH: > + sorry ("Clause not supported yet"); > + break; > + > default: > gcc_unreachable (); > } Indentation. > @@ -1827,9 +1876,26 @@ scan_sharing_clauses (tree clauses, omp_context *c= tx) > case OMP_CLAUSE__LOOPTEMP_: > case OMP_CLAUSE_TO: > case OMP_CLAUSE_FROM: > - gcc_assert (gimple_code (ctx->stmt) !=3D GIMPLE_OACC_PARALLEL); > + gcc_assert (!gimple_code_is_oacc (ctx->stmt)); > break; >=20=20 > + case OMP_CLAUSE_HOST: > + case OMP_CLAUSE_OACC_DEVICE: > + case OMP_CLAUSE_DEVICE_RESIDENT: > + case OMP_CLAUSE_USE_DEVICE: > + case OMP_CLAUSE_ASYNC: > + case OMP_CLAUSE_GANG: > + case OMP_CLAUSE_WAIT: > + case OMP_NO_CLAUSE_CACHE: > + case OMP_CLAUSE_INDEPENDENT: > + case OMP_CLAUSE_WORKER: > + case OMP_CLAUSE_VECTOR: > + case OMP_CLAUSE_NUM_GANGS: > + case OMP_CLAUSE_NUM_WORKERS: > + case OMP_CLAUSE_VECTOR_LENGTH: > + sorry ("Clause not supported yet"); > + break; > + > default: > gcc_unreachable (); > } Indentation. > --- a/gcc/tree-core.h > +++ b/gcc/tree-core.h > @@ -213,19 +213,19 @@ enum omp_clause_code { > (c_parser_omp_variable_list). */ > OMP_CLAUSE_ERROR =3D 0, >=20=20 > - /* OpenMP clause: private (variable_list). */ > + /* OpenMP/OpenACC clause: private (variable_list). */ > OMP_CLAUSE_PRIVATE, >=20=20 > /* OpenMP clause: shared (variable_list). */ > OMP_CLAUSE_SHARED, >=20=20 > - /* OpenMP clause: firstprivate (variable_list). */ > + /* OpenMP/OpenACC clause: firstprivate (variable_list). */ > OMP_CLAUSE_FIRSTPRIVATE, >=20=20 > /* OpenMP clause: lastprivate (variable_list). */ > OMP_CLAUSE_LASTPRIVATE, >=20=20 > - /* OpenMP clause: reduction (operator:variable_list). > + /* OpenMP/OpenACC clause: reduction (operator:variable_list). > OMP_CLAUSE_REDUCTION_CODE: The tree_code of the operator. > Operand 1: OMP_CLAUSE_REDUCTION_INIT: Stmt-list to initialize the v= ar. > Operand 2: OMP_CLAUSE_REDUCTION_MERGE: Stmt-list to merge private v= ar > @@ -265,10 +265,40 @@ enum omp_clause_code { > OpenMP clause: map ({alloc:,to:,from:,tofrom:,}variable-list). */ > OMP_CLAUSE_MAP, >=20=20 > + /* OpenACC clause: host (variable_list). */ > + OMP_CLAUSE_HOST, > + > + /* OpenACC clause: device (variable_list). */ > + OMP_CLAUSE_OACC_DEVICE, > + > + /* OpenACC clause: device_resident (variable_list). */ > + OMP_CLAUSE_DEVICE_RESIDENT, > + > + /* OpenACC clause: use_device (variable_list). */ > + OMP_CLAUSE_USE_DEVICE, > + > + /* OpenACC clause: async [(integer-expression)]. */ > + OMP_CLAUSE_ASYNC, > + > + /* OpenACC clause: gang [(gang-argument-list)].=20 > + Where=20 > + gang-argument-list: [gang-argument-list, ] gang-argument=20 > + gang-argument: [num:] integer-expression > + | static: size-expression > + size-expression: * | integer-expression. */ > + OMP_CLAUSE_GANG, > + > + /* OpenACC clause/directive: wait [(integer-expression-list)]. */ > + OMP_CLAUSE_WAIT, > + > + /* Internal structure to hold OpenACC cache directive's variable-list. > + #pragma acc cache (variable-list). */ > + OMP_NO_CLAUSE_CACHE, > + > /* Internal clause: temporary for combined loops expansion. */ > OMP_CLAUSE__LOOPTEMP_, >=20=20 > - /* OpenMP clause: if (scalar-expression). */ > + /* OpenMP/OpenACC clause: if (scalar-expression). */ > OMP_CLAUSE_IF, >=20=20 > /* OpenMP clause: num_threads (integer-expression). */ > @@ -281,12 +311,13 @@ enum omp_clause_code { > OMP_CLAUSE_NOWAIT, >=20=20 > /* OpenMP clause: ordered. */ > + /* OpenACC clause: seq. */ > OMP_CLAUSE_ORDERED, >=20=20 > /* OpenMP clause: default. */ > OMP_CLAUSE_DEFAULT, >=20=20 > - /* OpenMP clause: collapse (constant-integer-expression). */ > + /* OpenMP/OpenACC clause: collapse (constant-integer-expression). */ > OMP_CLAUSE_COLLAPSE, >=20=20 > /* OpenMP clause: untied. */ > @@ -338,7 +369,25 @@ enum omp_clause_code { > OMP_CLAUSE_TASKGROUP, >=20=20 > /* Internally used only clause, holding SIMD uid. */ > - OMP_CLAUSE__SIMDUID_ > + OMP_CLAUSE__SIMDUID_, > + > + /* OpenACC clause: independent. */ > + OMP_CLAUSE_INDEPENDENT, > + > + /* OpenACC clause: worker [( [num:] integer-expression)]. */ > + OMP_CLAUSE_WORKER, > + > + /* OpenACC clause: vector [( [length:] integer-expression)]. */ > + OMP_CLAUSE_VECTOR, > + > + /* OpenACC clause: num_gangs (integer-expression). */ > + OMP_CLAUSE_NUM_GANGS, > + > + /* OpenACC clause: num_workers (integer-expression). */ > + OMP_CLAUSE_NUM_WORKERS, > + > + /* OpenACC clause: vector_length (integer-expression). */ > + OMP_CLAUSE_VECTOR_LENGTH > }; Unless I'm confused, some of the new clauses (those that deal with more than just a single operand) will need to get special handling added =C3=A0 = la the other omp_clause_subcode variants in tree-core.h:tree_omp_clause? (It is fine to rework that in a later patch, once we actually implement these clauses.) > --- a/gcc/tree-pretty-print.c > +++ b/gcc/tree-pretty-print.c > @@ -326,6 +326,21 @@ dump_omp_clause (pretty_printer *buffer, tree clause= , int spc, int flags) OMP_CLAUSE_WAIT is missing. > @@ -2384,6 +2452,41 @@ dump_generic_node (pretty_printer *buffer, tree no= de, int spc, int flags, OACC_ENTER_DATA and OACC_EXIT_DATA are missing. > --- a/gcc/tree.c > +++ b/gcc/tree.c > @@ -327,7 +350,13 @@ const char * const omp_clause_code_name[] =3D > + "indepentend", Typo: independent. > @@ -11034,6 +11063,18 @@ walk_tree_1 (tree *tp, walk_tree_fn func, void *= data, OMP_CLAUSE_GANG missing. > case OMP_CLAUSE: > switch (OMP_CLAUSE_CODE (*tp)) > { > + case OMP_CLAUSE_HOST: > + case OMP_CLAUSE_OACC_DEVICE: > + case OMP_CLAUSE_DEVICE_RESIDENT: > + case OMP_CLAUSE_USE_DEVICE: > + case OMP_NO_CLAUSE_CACHE: > + case OMP_CLAUSE_ASYNC: > + case OMP_CLAUSE_WORKER: > + case OMP_CLAUSE_VECTOR: > + case OMP_CLAUSE_NUM_GANGS: > + case OMP_CLAUSE_NUM_WORKERS: > + case OMP_CLAUSE_VECTOR_LENGTH: > + case OMP_CLAUSE_WAIT: > case OMP_CLAUSE_PRIVATE: > case OMP_CLAUSE_SHARED: > case OMP_CLAUSE_FIRSTPRIVATE: Indentation. > @@ -11056,6 +11097,7 @@ walk_tree_1 (tree *tp, walk_tree_fn func, void *d= ata, > WALK_SUBTREE (OMP_CLAUSE_OPERAND (*tp, 0)); > /* FALLTHRU */ >=20=20 > + case OMP_CLAUSE_INDEPENDENT: > case OMP_CLAUSE_NOWAIT: > case OMP_CLAUSE_ORDERED: > case OMP_CLAUSE_DEFAULT: Indentation. > --- a/gcc/tree.def > +++ b/gcc/tree.def > @@ -1017,6 +1017,56 @@ DEFTREECODE (MEM_REF, "mem_ref", tcc_reference, 2) >=20=20 > DEFTREECODE (OACC_PARALLEL, "oacc_parallel", tcc_statement, 2) >=20=20 > +/* #pragma acc kernels */ > +/* Operand 0: BODY > + Operand 1: CLAUSES=20=20=20 > +*/ > +DEFTREECODE (OACC_KERNELS, "oacc_kernels", tcc_statement, 2) > + > +/* #pragma acc data */ > +/* Operand 0: BODY > + Operand 1: CLAUSES=20=20=20 > +*/ > +DEFTREECODE (OACC_DATA, "oacc_data", tcc_statement, 2) > + > +/* #pragma acc host_data */ > +/* Operand 0: BODY > + Operand 1: CLAUSES=20=20=20 > +*/ > +DEFTREECODE (OACC_HOST_DATA, "oacc_host_data", tcc_statement, 2) > + > +/* #pragma acc declare */ > +/* Operand 0: CLAUSES=20=20=20 > +*/ > +DEFTREECODE (OACC_DECLARE, "oacc_declare", tcc_statement, 1) > +[...] | [...] | DEFTREECODE (OMP_PARALLEL, "omp_parallel", tcc_statement, 2) | [...] | DEFTREECODE (OMP_SINGLE, "omp_single", tcc_statement, 2) | [...] | DEFTREECODE (OMP_CRITICAL, "omp_critical", tcc_statement, 2) | [...] If you look at tree.h:OMP_BODY and OMP_CLAUSES: #define OMP_BODY(NODE) \ TREE_OPERAND (TREE_RANGE_CHECK (NODE, OACC_PARALLEL, OMP_CRITICAL), 0) #define OMP_CLAUSES(NODE) \ TREE_OPERAND (TREE_RANGE_CHECK (NODE, OACC_PARALLEL, OMP_SINGLE), 1) ..., you'll see that any codes from OACC_PARALLEL to OMP_CRITICAL are expected to have a body as operand 0, and all from OACC_PARALLEL to OMP_SINGLE are expected to have a list of clauses as operand 1. This doesn't matter as long as all code paths are using the more specialized OMP_*_BODY/OACC_*_BODY and OMP_*_CLAUSES/OACC_*_CLAUSES accessor macros instead of the generic OMP_BODY and OMP_CLAUSES ones when working with such OMP_* nodes, but to reduce confusion for readers of this list/code, we should anyway pay attention, that is, distribute the new OACC_* tree codes in tree.def to the appropriate places, to fit into the existing scheme (as good as possible). Please then also try to maintain the same order of OACC_* codes in tree.h, tree.c, and so on: > --- a/gcc/tree.h > +++ b/gcc/tree.h > @@ -1367,6 +1367,59 @@ extern void protected_set_expr_location (tree, loc= ation_t); > #define OMP_CLAUSE_DEFAULT_KIND(NODE) \ > (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_DEFAULT)->omp_clause.subco= de.default_kind) >=20=20 > +/* OpenACC directives and clause accessors. */ > + > +#define OACC_KERNELS_BODY(NODE) TREE_OPERAND (OACC_KERNELS_CHECK(NO= DE), 0) > +#define OACC_KERNELS_CLAUSES(NODE) TREE_OPERAND (OACC_KERNELS_CHECK(NO= DE), 1) ..., that is, move these (and the following ones) after the existing OACC_PARALLEL_*, and so on, as appropriate. > +#define OACC_DATA_BODY(NODE) \ > + TREE_OPERAND (OACC_DATA_CHECK (NODE), 0) > + > +#define OACC_DATA_CLAUSES(NODE) \ > + TREE_OPERAND (OACC_DATA_CHECK (NODE), 1) No empty line between macros for the same tree code (OACC_DATA). > +/* OpenACC clauses */ > +#define OMP_CLAUSE_[...]_EXPR(NODE) \ Please also try to sort these into the appropriate order (that is, as defined in tree-core.h:omp_clause_code) -- if possible: I see that it is not really consistent right now. (This would also apply to the other files where any OMP_*/OACC_* are used, for example the pretty printers, but alas, if it's not in the right order now, then don't bother too much about it.) With these issues addressed, this patch is ready for commit to gomp-4_0-branch. Use your own judgement; if you feel confident, just commit it, or otherwise post it again for a final review -- as you prefer. Gr=C3=BC=C3=9Fe, Thomas --=-=-= Content-Type: application/pgp-signature Content-length: 489 -----BEGIN PGP SIGNATURE----- Version: GnuPG v1.4.12 (GNU/Linux) iQEcBAEBAgAGBQJS+lUDAAoJENuKOtuXzphJ4m0IAIPwHcQfSJ3ez4TsSGSxkALL Mq9GpOr6tB694P39GdcZWz55iTau7iQo0gtaiSjd98OcEXCH5K+IiChshzVg2hVi VAM2tS9jIxd/mLeJ72mlQBGo98jfV1i4h0ZW4UCYLDwrnkkrEnqnhKjIf4YqP20o GKGTJtAopG0T8J/wMgpZUPtNSEHaaruro1NVuC905zGYNaCAtBS5eQIv5usQKSUQ cW9RW2ifXKMIEQ+kj6Hos/wWeqWveeCJzBmVLBr6YbOQmarBzd2W1L210pcIqChH OkiZi5yOU4ksateg7JvCHRIanRdzOat3ThPDkDCgy3flY4KLqFTB3vHeKnBIWnE= =LlPU -----END PGP SIGNATURE----- --=-=-=--