Re: OpenACC Firstprivate

[Jakub Jelinek <jakub@redhat.com>]

On Sat, Nov 07, 2015 at 08:50:28AM -0500, Nathan Sidwell wrote:
> Index: gcc/gimplify.c
> ===================================================================
> --- gcc/gimplify.c	(revision 229892)
> +++ gcc/gimplify.c	(working copy)
> @@ -108,9 +108,15 @@ enum omp_region_type
>    /* Data region with offloading.  */
>    ORT_TARGET = 32,
>    ORT_COMBINED_TARGET = 33,
> +
> +  ORT_ACC = 0x40,  /* An OpenACC region.  */
> +  ORT_ACC_DATA = ORT_ACC | ORT_TARGET_DATA, /* Data construct.  */
> +  ORT_ACC_PARALLEL = ORT_ACC | ORT_TARGET,  /* Parallel construct */
> +  ORT_ACC_KERNELS  = ORT_ACC | ORT_TARGET | 0x80,  /* Kernels construct.  */
> +
>    /* Dummy OpenMP region, used to disable expansion of
>       DECL_VALUE_EXPRs in taskloop pre body.  */
> -  ORT_NONE = 64
> +  ORT_NONE = 0x100
>  };

If you want to switch to hexadecimal, you should change all values
in the enum to hexadecimal for consistency.
>  
>  /* Gimplify hashtable helper.  */
> @@ -377,6 +383,12 @@ new_omp_context (enum omp_region_type re
>    else
>      c->default_kind = OMP_CLAUSE_DEFAULT_UNSPECIFIED;
>  
> +  c->combined_loop = false;
> +  c->distribute = false;
> +  c->target_map_scalars_firstprivate = false;
> +  c->target_map_pointers_as_0len_arrays = false;
> +  c->target_firstprivatize_array_bases = false;

Why this?  c is XCNEW allocated, so zero initialized.

> @@ -5667,11 +5682,13 @@ omp_add_variable (struct gimplify_omp_ct
>        /* We shouldn't be re-adding the decl with the same data
>  	 sharing class.  */
>        gcc_assert ((n->value & GOVD_DATA_SHARE_CLASS & flags) == 0);
> -      /* The only combination of data sharing classes we should see is
> -	 FIRSTPRIVATE and LASTPRIVATE.  */
>        nflags = n->value | flags;
> -      gcc_assert ((nflags & GOVD_DATA_SHARE_CLASS)
> -		  == (GOVD_FIRSTPRIVATE | GOVD_LASTPRIVATE)
> +      /* The only combination of data sharing classes we should see is
> +	 FIRSTPRIVATE and LASTPRIVATE.  However, OpenACC permits
> +	 reduction variables to be used in data sharing clauses.  */
> +      gcc_assert ((ctx->region_type & ORT_ACC) != 0
> +		  || ((nflags & GOVD_DATA_SHARE_CLASS)
> +		      == (GOVD_FIRSTPRIVATE | GOVD_LASTPRIVATE))
>  		  || (flags & GOVD_DATA_SHARE_CLASS) == 0);

Are you sure you want to give up on any kind of consistency checks for
OpenACC?  If only reduction is special on OpenACC, perhaps you could tweak
the assert for that instead?  Something that can be done incrementally of
course.

> +
> +	  /*  OpenMP doesn't look in outer contexts to find an
> +	      enclosing data clause.  */

I'm puzzled by the comment.  OpenMP does look in outer context for clauses
that need that (pretty much all closes but private), that is the do_outer:
recursion in omp_notice_variable.  Say for firstprivate in order to copy (or
copy construct) the private variable one needs the access to the outer
context's var etc.).
So perhaps it would help to document what you are doing here for OpenACC and
why.

> +	  struct gimplify_omp_ctx *octx = ctx->outer_context;
> +	  if ((ctx->region_type & ORT_ACC) && octx)
> +	    {
> +	      omp_notice_variable (octx, decl, in_code);
> +	      
> +	      for (; octx; octx = octx->outer_context)
> +		{
> +		  if (!(octx->region_type & (ORT_TARGET_DATA | ORT_TARGET)))
> +		    break;
> +		  splay_tree_node n2
> +		    = splay_tree_lookup (octx->variables,
> +					 (splay_tree_key) decl);
> +		  if (n2)
> +		    {
> +		      nflags |= GOVD_MAP;
> +		      goto found_outer;
> +		    }
> +		}
>  	    }
> -	  else if (nflags == flags)
> -	    nflags |= GOVD_MAP;
> +

The main issue I have is with the omp-low.c changes.
I see:
"2.5.9
private clause
The private clause is allowed on the parallel construct; it declares that a copy of each
item on the list will be created for each parallel gang.

2.5.10
firstprivate clause
The firstprivate clause is allowed on the parallel construct; it declares that a copy
of each item on the list will be created for each parallel gang, and that the copy will be
initialized with the value of that item on the host when the parallel construct is
encountered."

but looking at what you actually emit looks like standard present_copyin
clause I think with a private variable defined in the region where the
value of the present_copyin mapped variable is assigned to the private one.
This I'm afraid performs often two copies rather than just one (one to copy
the host value to the present_copyin mapped value, another one in the
region), but more importantly, if the var is already mapped, you could
initialize the private var with old data.
Say
  int arr[64];
// initialize arr
#pragma acc data copyin (arr)
{
  // modify arr on the host
  # pragma acc parallel firstprivate (arr)
  {
    ...
  }
}
Is that really what you want?  If not, any reason not to implement
GOMP_MAP_FIRSTPRIVATE and GOMP_MAP_FIRSTPRIVATE_INT on the libgomp oacc-*
side and just use the OpenMP firstprivate handling in omp-low.c?

	Jakub