Re: [OpenACC] declare directive

[Jakub Jelinek <jakub at redhat dot com>]

On Fri, Nov 06, 2015 at 10:08:26AM -0600, James Norris wrote:
> 2015-10-27  James Norris  <jnorris@codesourcery.com>
> 	    Joseph Myers  <joseph@codesourcery.com>
> 
> 	gcc/
> 	* c-family/c-pragma.c (oacc_pragmas): Add entry for declare directive. 
> 	* c-family/c-pragma.h (enum pragma_kind): Add PRAGMA_OACC_DECLARE.

c-family/, c/, cp/ have all their own ChangeLog files, and the entries
that go in there should not have the path prefixes.

> +  for (t = clauses; t; t = OMP_CLAUSE_CHAIN (t))
> +    {
> +      location_t loc = OMP_CLAUSE_LOCATION (t);
> +      tree decl = OMP_CLAUSE_DECL (t);
> +      tree devres = NULL_TREE;
> +      if (!DECL_P (decl))
> +	{
> +	  error_at (loc, "subarray in %<#pragma acc declare%>");

Is that the OpenACC term that you use everywhere?  In OpenMP
those are array sections.

> +	case GOMP_MAP_LINK:
> +	  if (!global_bindings_p () && !DECL_EXTERNAL (decl))
> +	    {
> +	      error_at (loc,
> +			"%qD must be a global variable in"
> +			"%<#pragma acc declare link%>",
> +			decl);
> +	      error = true;
> +	      continue;

What about TREE_STATIC?

> diff --git a/gcc/gimplify.c b/gcc/gimplify.c
> index fa34858..fbd4a69 100644
> --- a/gcc/gimplify.c
> +++ b/gcc/gimplify.c
> @@ -157,6 +157,7 @@ struct gimplify_omp_ctx
>    bool target_map_scalars_firstprivate;
>    bool target_map_pointers_as_0len_arrays;
>    bool target_firstprivatize_array_bases;
> +  gomp_target *declare_returns;
>  };

This declare_returns stuff looks broken.
What exactly do you want to achieve?  Why do you record it
in gimplify_omp_ctx, but then only look at it in gimplify_body?
The way you abuse gimplify_omp_ctx for that is just too ugly.
All the gimplification code expects that when you enter some OpenMP/OpenACC
construct, you create new context if it needs one, then process the body
of the construct, then pop it up.  The declare_returns stuff
violates all of this.  Do you want to enqueue all the statements
at the end of the body?  Then just stick it into some gimple_seq
outside of the gimplify_omp_ctx, and process in there.  Or if you
want to process it when leaving some construct, arrange for that.

> @@ -5841,6 +5863,8 @@ omp_default_clause (struct gimplify_omp_ctx *ctx, tree decl,
>        flags |= GOVD_FIRSTPRIVATE;
>        break;
>      case OMP_CLAUSE_DEFAULT_UNSPECIFIED:
> +      if (is_global_var (decl) && device_resident_p (decl))
> +	flags |= GOVD_MAP_TO_ONLY | GOVD_MAP;

I don't think you want to do this except for (selected or all?)
OpenACC contexts.  Say, I don't see why one couldn't e.g. try to mix
OpenMP host parallelization or tasking with OpenACC offloading,
and that affecting in weird way OpenMP semantics.

> +      /* Any clauses that affect the state of a variable prior
> +         to return are saved and dealt with after the body has
> +         been gimplified.  */
> +
> +      gimplify_scan_omp_clauses (&clauses, pre_p, ORT_TARGET_DATA,
> +				 OACC_DECLARE);
> +
> +      c = gimplify_omp_ctxp;
> +      gimplify_omp_ctxp = c->outer_context;
> +      delete_omp_context (c);

Why don't you call gimplify_adjust_omp_clauses instead?

> +      stmt = gimple_build_omp_target (NULL, GF_OMP_TARGET_KIND_OACC_DECLARE,
> +				      clauses);
> +      gimplify_omp_ctxp->declare_returns = stmt;

But the above is the main thing I have trouble with.
What happens if you have multiple #pragma acc declare directives?
Is the stmt from the other ones lost?
I'd expect something like gimple_seq in there instead and pushing
stmts into the sequence.
I don't know what is the semantics of the construct, if it is something
close to say target data in OpenMP or acc data in OpenACC, the addition
of code to unmap the variables is performed using a cleanup, otherwise
how do you handle exceptions, or goto and all other kinds of abnormal
returns from the function.

> @@ -160,6 +151,25 @@ varpool_node::get_create (tree decl)
>        node->force_output = 1;
>  #endif
>      }
> +}
> +
> +/* Return varpool node assigned to DECL.  Create new one when needed.  */
> +varpool_node *
> +varpool_node::get_create (tree decl)
> +{
> +  varpool_node *node = varpool_node::get (decl);
> +  gcc_checking_assert (TREE_CODE (decl) == VAR_DECL);
> +  if (node)
> +    {
> +      if (!node->offloadable)
> +	make_offloadable (node, decl);

I don't like this change at all, that is significant extra work
every time somebody calls this function.
Just do what we do on the OpenMP side, if you add some "omp declare target"
or similar attribute and the varpool node for it has been already created,
set offloadable in the FE.

	Jakub