Re: OpenACC Firstprivate

[Nathan Sidwell <nathan@acm.org>]

On 11/09/15 08:46, Jakub Jelinek wrote:
> On Sat, Nov 07, 2015 at 08:50:28AM -0500, Nathan Sidwell wrote:
>> Index: gcc/gimplify.c
>> ===================================================================

>
> If you want to switch to hexadecimal, you should change all values
> in the enum to hexadecimal for consistency.

ok.

>>
>>   /* 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.

I presumed it necessary, as it was on the branch.  will  remove.

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

Will investigate (later)

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

Ok.  It seemed (and it may become clearer with default handling added), that 
OpenACC  and OpenMP scanned scopes in opposite orders.  I remember trying to 
get the ACC code to scan in the same order, but came up blank.  Anyway, you're 
right, it should say what OpenACC is trying.


> 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),

I don't think that can be avoided.  The host doesn't have control over when the 
CTAs (a gang) start -- they may even be serialized onto the same physical HW. 
So each gang has to initialize its own instance.  Or did you mean something else?

> 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)
>    {
>      ...
>    }
> }

Hm, I suspect that is either ill formed or the std does not contemplate.

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

I would have to investigate ...

nathan