Re: [Bulk] [OpenACC 0/7] host_data construct

[Jakub Jelinek <jakub at redhat dot com>]

On Fri, Oct 23, 2015 at 10:51:42AM -0500, James Norris wrote:
> @@ -12942,6 +12961,7 @@ c_finish_omp_clauses (tree clauses, bool is_omp, bool declare_simd)
>  	case OMP_CLAUSE_GANG:
>  	case OMP_CLAUSE_WORKER:
>  	case OMP_CLAUSE_VECTOR:
> +	case OMP_CLAUSE_USE_DEVICE:
>  	  pc = &OMP_CLAUSE_CHAIN (c);
>  	  continue;
>  

Are there any restrictions on whether you can specify the same var multiple
times in use_device clause?
#pragma acc host_data use_device (x) use_device (x) use_device (y, y, y)
?
If not, have you verified that the gimplifier doesn't ICE on it?  Generally
it doesn't like the same var being mentioned multiple times.
If yes, you can use e.g. the generic_head bitmap for that and in any case,
cover that with sufficient testsuite coverage.

> diff --git a/gcc/gimplify.c b/gcc/gimplify.c
> index ab9e540..0c32219 100644
> --- a/gcc/gimplify.c
> +++ b/gcc/gimplify.c
> @@ -93,6 +93,8 @@ enum gimplify_omp_var_data
>  
>    GOVD_MAP_0LEN_ARRAY = 32768,
>  
> +  GOVD_USE_DEVICE = 65536,
> +
>    GOVD_DATA_SHARE_CLASS = (GOVD_SHARED | GOVD_PRIVATE | GOVD_FIRSTPRIVATE
>  			   | GOVD_LASTPRIVATE | GOVD_REDUCTION | GOVD_LINEAR
>  			   | GOVD_LOCAL)
> @@ -116,7 +118,9 @@ enum omp_region_type
>    ORT_COMBINED_TARGET = 33,
>    /* Dummy OpenMP region, used to disable expansion of
>       DECL_VALUE_EXPRs in taskloop pre body.  */
> -  ORT_NONE = 64
> +  ORT_NONE = 64,
> +  /* An OpenACC host-data region.  */
> +  ORT_HOST_DATA = 128

I'd prefer ORT_NONE to be the last one, can you just renumber it and put
ORT_HOST_DATA before it?

> +static tree
> +gimplify_oacc_host_data_1 (tree *tp, int *walk_subtrees,
> +			   void *data ATTRIBUTE_UNUSED)
> +{

Your use_device sounds very similar to use_device_ptr clause in OpenMP,
which is allowed on #pragma omp target data construct and is implemented
quite a bit differently from this; it is unclear if the OpenACC standard
requires this kind of implementation, or you just chose to implement it this
way.  In particular, the GOMP_target_data call puts the variables mentioned
in the use_device_ptr clauses into the mapping structures (similarly how
map clause appears) and the corresponding vars are privatized within the
target data region (which is a host region, basically a fancy { } braces),
where the private variables contain the offloading device's pointers.

> +  splay_tree_node n = NULL;
> +  location_t loc = EXPR_LOCATION (*tp);
> +
> +  switch (TREE_CODE (*tp))
> +    {
> +    case ADDR_EXPR:
> +      {
> +	tree decl = TREE_OPERAND (*tp, 0);
> +
> +	switch (TREE_CODE (decl))
> +	  {
> +	  case ARRAY_REF:
> +	  case ARRAY_RANGE_REF:
> +	  case COMPONENT_REF:
> +	  case VIEW_CONVERT_EXPR:
> +	  case REALPART_EXPR:
> +	  case IMAGPART_EXPR:
> +	    if (TREE_CODE (TREE_OPERAND (decl, 0)) == VAR_DECL)
> +	      n = splay_tree_lookup (gimplify_omp_ctxp->variables,
> +				     (splay_tree_key) TREE_OPERAND (decl, 0));
> +	    break;

I must say this looks really strange, you throw away all the offsets
embedded in the component codes (fixed or variable).
Where comes the above list?  What about other components (say bit field refs,
etc.)?

> +    case VAR_DECL:

What is so special about VAR_DECLs?  Shouldn't PARM_DECLs / RESULT_DECLs
be treated the same way?
> --- a/libgomp/libgomp.map
> +++ b/libgomp/libgomp.map
> @@ -378,6 +378,7 @@ GOACC_2.0 {
>  	GOACC_wait;
>  	GOACC_get_thread_num;
>  	GOACC_get_num_threads;
> +	GOACC_deviceptr;
>  };
>  
>  GOACC_2.0.1 {

You shouldn't be adding new symbols into a symbol version that appeared in a
compiler that shipped already (GCC 5 already had GOACC_2.0 symbols).
So it should go into GOACC_2.0.1.

> diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c
> index af067d6..497ab92 100644
> --- a/libgomp/oacc-mem.c
> +++ b/libgomp/oacc-mem.c
> @@ -204,6 +204,38 @@ acc_deviceptr (void *h)
>    return d;
>  }
>  
> +/* This function is used as a helper in generated code to implement pointer
> +   lookup in host_data regions.  Unlike acc_deviceptr, it returns its argument
> +   unchanged on a shared-memory system (e.g. the host).  */
> +
> +void *
> +GOACC_deviceptr (void *h)
> +{
> +  splay_tree_key n;
> +  void *d;
> +  void *offset;
> +
> +  goacc_lazy_initialize ();
> +
> +  struct goacc_thread *thr = goacc_thread ();
> +
> +  if ((thr->dev->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM) == 0)
> +    {
> +      n = lookup_host (thr->dev, h, 1);

What is supposed to be the behavior when the h pointer points at object
boundary, rather than into the middle of existing mapped object?

Say you have:
  char a[16], b[0], c[16]; // b is GCC extension
Now, char *p = &a[5]; is unambiguous, either a is mapped, or not.
But, if p = &a[16];, then it could be either the one-past-last byte in a,
or it could be the start of b (== one-past-last byte in b) or it could be
the pointer to start of c.

In OpenMP 4.5, I had endless discussions about this and the end result is
that one-past-last byte addresses are unspecified behavior

	Jakub