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

[Cesar Philippidis <cesar@codesourcery.com>]

On 10/26/2015 11:34 AM, Jakub Jelinek wrote:
> 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.

Generally variables cannot appear in multiple clauses. I'll add more
testing for this.

>> 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?

OK.

>> +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.

Is this a new OpenMP 4.5 feature? I'll take a closer look and see if
they are similar enough. I also noticed that OpenMP 4.5 has something
similar to OpenACC's enter/exit data construct now.

>> +  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.)?

I'm not sure. This is one of those things where multiple developers
worked on it, and the history got lost. I'll investigate it.

>> +    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.

OK.

>> 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?

Probably undefined with the way that OpenACC is defined.

> 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

OK.

Thanks for you feedback.

Cesar