From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: (qmail 3404 invoked by alias); 24 Jan 2014 19:33:54 -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 3369 invoked by uid 89); 24 Jan 2014 19:33:53 -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-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; Fri, 24 Jan 2014 19:33:51 +0000 Received: from svr-orw-fem-01.mgc.mentorg.com ([147.34.98.93]) by relay1.mentorg.com with esmtp id 1W6mVc-0006aJ-FZ from Thomas_Schwinge@mentor.com ; Fri, 24 Jan 2014 11:33:44 -0800 Received: from SVR-IES-FEM-02.mgc.mentorg.com ([137.202.0.106]) by svr-orw-fem-01.mgc.mentorg.com over TLS secured channel with Microsoft SMTPSVC(6.0.3790.4675); Fri, 24 Jan 2014 11:33:44 -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.2.247.3; Fri, 24 Jan 2014 19:33:41 +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: <52E159BD.9040407@samsung.com> References: <52E158EF.9050009@samsung.com> <52E1595D.9000007@samsung.com> <52E1597E.8070809@samsung.com> <52E15999.6010505@samsung.com> <52E159BD.9040407@samsung.com> User-Agent: Notmuch/0.9-101-g81dad07 (http://notmuchmail.org) Emacs/23.4.1 (x86_64-pc-linux-gnu) Date: Fri, 24 Jan 2014 19:33:00 -0000 Message-ID: <874n4tqink.fsf@schwinge.name> MIME-Version: 1.0 Content-Type: multipart/signed; boundary="=-=-="; micalg=pgp-sha1; protocol="application/pgp-signature" X-SW-Source: 2014-01/txt/msg01613.txt.bz2 --=-=-= Content-Type: text/plain; charset=utf-8 Content-Transfer-Encoding: quoted-printable Content-length: 9209 Hi! On Thu, 23 Jan 2014 22:04:45 +0400, Ilmir Usmanov w= rote: > Subject: [PATCH 4/6] OpenACC GENERIC nodes > --- a/gcc/gimplify.c > +++ b/gcc/gimplify.c > @@ -7846,6 +7907,18 @@ 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: > + ret =3D GS_ALL_DONE; > + break; Would it make sense, until it is really implemented, to add something like: sorry ("directive not yet implemented"); That way it'd be more obvious that this still missing. > --- a/gcc/omp-low.c > +++ b/gcc/omp-low.c > @@ -1543,10 +1543,12 @@ scan_sharing_clauses (tree clauses, omp_context *= ctx) > break; > /* FALLTHRU */ >=20=20 > - case OMP_CLAUSE_FIRSTPRIVATE: > - case OMP_CLAUSE_REDUCTION: > case OMP_CLAUSE_LINEAR: > gcc_assert (gimple_code (ctx->stmt) !=3D GIMPLE_OACC_PARALLEL); > + /* FALLTHRU. */ > + > + case OMP_CLAUSE_FIRSTPRIVATE: /* =3D=3D OACC_CLAUSE_FIRSTPRIVATE. */ > + case OMP_CLAUSE_REDUCTION: /* =3D=3D OACC_CLAUSE_REDUCTION. */ > decl =3D OMP_CLAUSE_DECL (c); > do_private: > if (is_variable_sized (decl)) By the way, what I have been using these gcc_asserts for, is to mark all code paths that I had not yet verified to be suitable for OpenACC. As we're now close to adding additional GIMPLE_OACC_* codes, maybe we should have some gimple_code_is_oacc function instead of just watching out for GIMPLE_OACC_PARALLEL? That is, if we generally do agree that we want to continue putting such gcc_asserts into the code? > @@ -1713,6 +1719,35 @@ scan_sharing_clauses (tree clauses, omp_context *c= tx) > install_var_local (decl, ctx); > break; >=20=20 > + case OACC_CLAUSE_COPY: > + case OACC_[...] > + /* Not implemented yet. */ > + break; > @@ -1822,6 +1856,38 @@ scan_sharing_clauses (tree clauses, omp_context *c= tx) > gcc_assert (gimple_code (ctx->stmt) !=3D GIMPLE_OACC_PARALLEL); > break; >=20=20 > + case OACC_CLAUSE_COPY: > + case OACC_[...] > + /* Not implemented yet. */ > + break; Likewise for these, I wonder, add a sorry messages? > --- a/gcc/tree-core.h > +++ b/gcc/tree-core.h > @@ -216,12 +216,18 @@ enum omp_clause_code { > /* OpenMP clause: private (variable_list). */ > OMP_CLAUSE_PRIVATE, >=20=20 > + /* OpenACC clause: private (variable_list). */ > + OACC_CLAUSE_PRIVATE =3D OMP_CLAUSE_PRIVATE, I prefer to avoid such duplicate names; here and in other instances in the following. I think they mostly just add redundancy. Instead, I prefer to map (in my mind) the existing OMP_CLAUSE_* to OACC_or_OMP_CLAUSE_*. Is that acceptable for you? > /* OpenMP clause: copyin (variable_list). */ > OMP_CLAUSE_COPYIN, >=20=20 > + /* OpenACC clause: copyin (variable_list). */ > + OACC_CLAUSE_COPYIN =3D OMP_CLAUSE_COPYIN, > + > /* OpenMP clause: copyprivate (variable_list). */ > OMP_CLAUSE_COPYPRIVATE, >=20=20 > @@ -261,12 +273,79 @@ enum omp_clause_code { > /* OpenMP clause: map ({alloc:,to:,from:,tofrom:,}variable-list). */ > OMP_CLAUSE_MAP, >=20=20 > + /* OpenACC clause: copy (variable_list). */ > + OACC_CLAUSE_COPY, > + > + /* OpenACC clause: [data clauses...] Now, the data clauses. In my patch series posted at , I have chosen the approach to map the OpenACC data clauses all into the existing OMP_CLAUSE_MAP's map_kind, and extend that as required. That seemed the more appropriate approach to me, as it's the very same functionality. Also, OMP_CLAUSE_COPYIN semantically is something different from OACC_CLAUSE_COPYIN. > + /* OpenACC clause: host (variable_list). */ > + OACC_CLAUSE_HOST, If we're adding new names for implementing OpenACC things, maybe we should also name these OMP_*, to keep things simple to read in the code that uses them. While they clearly do stem from OpenMP, with a future generalization task in mind, I understand the OMP_CLAUSE_* (and generally OMP_*) names to no longer mean something specific to OpenMP (or OpenACC, for that matter), but instead to define a generic interface to the respective functionality, that just happens to have OMP_* names due to its OpenMP legacy. For myself, I decided to apply this to clauses only, so -- at least for now -- I have not spent any time on generalizing the OpenACC and OpenMP directives themselves. (For example, very roughly, would an OpenACC parallel directive be something like a (suitably extended) OpenMP parallel directive inside a OpenMP target region?) So, I think it's find to add new tree codes for the OpenACC directives (as you've done), and any generalization work can be done later, once we have a good understanding of what the required basic primitives actually are. Back to the OpenACC host clause. However, to not give the impression that there is a OpenMP host clause, maybe this one should be named OMP_CLAUSE_OACC_HOST? (And, by the way, I do agree that naming this *_HOST instead of *_SELF is better, for *_HOST is the more descriptive name of the two synonyms.) But then, I'm not sure we need this clause at all: can't we just again use the existing OMP_CLAUSE_MAP with an appropriate map_kind, attached to a *_UPDATE directive, to describe what the OpenACC host clause (and others) semantically mean? (I have not yet looked at that in detail. If you can't comment on this right now, I'm fine with keeping that as it is now, and we can still change that later.) > + /* OpenACC clause: wait [(integer-expression-list)]. */ > + OACC_CLAUSE_WAIT, > + > + /* Internal structure to hold OpenACC wait directive's variable_list. > + #pragma acc wait [(integer-expression-list)]. */ > + OACC_NO_CLAUSE_WAIT, > + > + /* Internal structure to hold OpenACC cache directive's variable-list. > + #pragma acc cache (variable-_ist). */ > + OACC_NO_CLAUSE_CACHE, Hmm, yeah, while *_NO_CLAUSE_* perhaps isn't the most beautiful approach, I think it's fine at least for now. > @@ -660,7 +763,6 @@ enum annot_expr_kind { > annot_expr_ivdep_kind > }; >=20=20 > - > /*----------------------------------------------------------------------= ----- > Type definitions > ------------------------------------------------------------------------= ---*/ Even though myself, I don't care too much about this topic, it is generally considered good proactice in GCC, to try to avoid such "spurious" changes, to reduce the patches' volume. If changing such things as formatting in regions not directly related to the code you're editiing, as well as fixing typos is comments, and so on, that should be separate patches, which most often can be committed directly on trunk, as "obvious" changes. > --- a/gcc/tree-pretty-print.c > +++ b/gcc/tree-pretty-print.c > @@ -54,6 +54,248 @@ static void do_niy (pretty_printer *, const_tree); > static pretty_printer buffer; > static int initialized =3D 0; >=20=20 > +unsigned char > +dump_oacc_body (int flags, tree node, int spc, > + unsigned char is_expr, pretty_printer* buffer) That's the same as dump_omp_body, right? Again, putting this into a function instead of making it a goto label does make some sense, but it is an independent change. > +void > +dump_oacc_clause_remap (const char* name, > +void > +dump_oacc_clause (pretty_printer *buffer, tree clause, int spc, int flag= s) Same here (dump_omp_clause:print_remap, and dump_omp_clause). You are right that if we use dump_omp_clause for the OpenACC clauses, we won't be (easily) able to exactly reproduce the input directives, especially given there are OMP_CLAUSE_* shared between OpenACC and OpenMP. So, with my patch series for initial support for OpenACC data clauses, I now get things like: =C2=BB#pragma acc parallel map(from:d) map(force_tofrom:b)=C2= =AB meaning =C2=BB#pragma acc parallel present_or_copyin(d) copy(b)=C2=AB, but = I don't think this is a problem, as that is for debugging output only. Would that work for you? > --- a/gcc/tree.h > +++ b/gcc/tree.h > @@ -1390,6 +1385,91 @@ 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_BODY(NODE) \ > + TREE_OPERAND (NODE, 0) > + > +#define OACC_[...] Again, I'd just re-use (and extend as required) the existing OMP_* ones. > @@ -3585,6 +3665,11 @@ extern tree build_translation_unit_decl (tree); > extern tree build_block (tree, tree, tree, tree); > extern tree build_empty_stmt (location_t); > extern tree build_omp_clause (location_t, enum omp_clause_code); > +static inline tree > +build_oacc_clause (location_t loc, enum omp_clause_code code) > +{ > + return build_omp_clause (loc, code); > +} ..., and here, just use build_omp_clause at the call sites, too. Gr=C3=BC=C3=9Fe, Thomas --=-=-= Content-Type: application/pgp-signature Content-length: 489 -----BEGIN PGP SIGNATURE----- Version: GnuPG v1.4.12 (GNU/Linux) iQEcBAEBAgAGBQJS4sAPAAoJENuKOtuXzphJwcsH/jtBq1wRvivUZBcEVzgS/7ez firwL65rg0yPAaNCA1Mv2fEDNg+QOGxwBQ4GMSzSTa3yNEOsXF8njynADEXUEOTX 9E0LOUD7OUsVX98PCxc9hWjqlBW8ACKlAllCoows1rJQA4O0JejnlJduqdSV1RnH yh++/E1uG5xTxFkTjqSRFW841wX6wSeXeNCkstIJwmKIv/c1mnqfYFm4p4pPk+FH F0iRni+Dl64iqW8FRfePx9aao2Q6MGXUStxrl6Sr4zJFrzE6vI2Y0+8AxpJTwML8 7/wc1gljkXXQJQNOE7Es3enY5KbOe4iwpsMjlKW0SGXRPVwsbSpD4XTObj1BW6I= =cDOl -----END PGP SIGNATURE----- --=-=-=--