Hi! On Thu, 23 Jan 2014 22:04:45 +0400, Ilmir Usmanov <i.usma...@samsung.com> wrote: > 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,
> gimple_seq *post_p,
> ret = GS_ALL_DONE;
> break;
>
> + 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 = 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 */
>
> - case OMP_CLAUSE_FIRSTPRIVATE:
> - case OMP_CLAUSE_REDUCTION:
> case OMP_CLAUSE_LINEAR:
> gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL);
> + /* FALLTHRU. */
> +
> + case OMP_CLAUSE_FIRSTPRIVATE: /* == OACC_CLAUSE_FIRSTPRIVATE. */
> + case OMP_CLAUSE_REDUCTION: /* == OACC_CLAUSE_REDUCTION. */
> decl = 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 *ctx)
> install_var_local (decl, ctx);
> break;
>
> + case OACC_CLAUSE_COPY:
> + case OACC_[...]
> + /* Not implemented yet. */
> + break;
> @@ -1822,6 +1856,38 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
> gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL);
> break;
>
> + 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,
>
> + /* OpenACC clause: private (variable_list). */
> + OACC_CLAUSE_PRIVATE = 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,
>
> + /* OpenACC clause: copyin (variable_list). */
> + OACC_CLAUSE_COPYIN = OMP_CLAUSE_COPYIN,
> +
> /* OpenMP clause: copyprivate (variable_list). */
> OMP_CLAUSE_COPYPRIVATE,
>
> @@ -261,12 +273,79 @@ enum omp_clause_code {
> /* OpenMP clause: map ({alloc:,to:,from:,tofrom:,}variable-list). */
> OMP_CLAUSE_MAP,
>
> + /* OpenACC clause: copy (variable_list). */
> + OACC_CLAUSE_COPY,
> +
> + /* OpenACC clause: [data clauses...]
Now, the data clauses. In my patch series posted at
<http://news.gmane.org/find-root.php?message_id=%3C87ppnuvbv6.fsf%40schwinge.name%3E>,
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
> };
>
> -
> /*---------------------------------------------------------------------------
> 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 = 0;
>
> +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 flags)
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: »#pragma acc parallel map(from:d) map(force_tofrom:b)«
meaning »#pragma acc parallel present_or_copyin(d) copy(b)«, 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,
> location_t);
> #define OMP_CLAUSE_DEFAULT_KIND(NODE) \
> (OMP_CLAUSE_SUBCODE_CHECK (NODE,
> OMP_CLAUSE_DEFAULT)->omp_clause.subcode.default_kind)
>
> +/* 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üße,
Thomas
pgpI8m2YXqBPo.pgp
Description: PGP signature
