Re: OpenACC Firstprivate

[Nathan Sidwell <nathan@acm.org>]

On 11/09/15 09:10, Jakub Jelinek wrote:
> On Mon, Nov 09, 2015 at 08:59:15AM -0500, Nathan Sidwell wrote:
>>> 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?
>
> So, what is the scope of the private and firstprivate vars in OpenACC?
> In OpenMP if a variable is private or firstprivate on the target construct,
> unless further privatized in inner constructs it is really shared among all
> the threads in all the teams (ro one var per all CTAs/workers in PTX terms).
> Is that the case for OpenACC too, or are the vars e.g. private to each CTA
> already or to each thread in each CTA, something different?
> If they are shared by all CTAs, then you should hopefully be able to use the
> GOMP_MAP_FIRSTPRIVATE{,_INT}, if not, then I'd say you should at least use
> those to provide you the initializer data to initialize your private vars
> from as a cheaper alternative to mapping.

I'm going to try and get clarification, but I think the intent is to initialize 
with the value seen on the device.  Consider:


int foo = 0;
#pragma acc data copyin(foo)
{
   #pragma acc parallel present(foo)
   {
     foo = 2;
   }

   if (expr){
     #pragma update host (foo)
   }

   #pragma acc parallel firstprivate (foo)
   {
   // which initialization value?
   }
}

Here we copy data to the device, then set it a distinct value there.  We 
conditionally update the host's instance from the device.

My thinking is that the intent of the firstprivate is to initialize with the 
value known on the device (and behave as-if copyin, if it's not there).  Not the 
value most recently seen on the host -- the update clause could change that, and 
may well be being used as a debugging aide, so it seems bizarre that it can 
change program semantics in such a way.