This is the mail archive of the gcc-patches@gcc.gnu.org mailing list for the GCC project.


Index Nav: [Date Index] [Subject Index] [Author Index] [Thread Index]
Message Nav: [Date Prev] [Date Next] [Thread Prev] [Thread Next]
Other format: [Raw text]

Re: [patch, libgomp, OpenACC] Additional enter/exit data map handling


Ping.

On 2016/11/3 10:22 PM, Chung-Lin Tang wrote:
> 
> Ping this patch again.
> 
> On 2016/9/21 12:43 AM, Cesar Philippidis wrote:
>>> +/* Returns the number of mappings associated with the pointer or pset. PSET
>>>> +   have three mappings, whereas pointer have two.  */
>>>> +
>>>>  static int
>>>> -find_pset (int pos, size_t mapnum, unsigned short *kinds)
>>>> +find_pointer (int pos, size_t mapnum, unsigned short *kinds)
>>>>  {
>>>>    if (pos + 1 >= mapnum)
>>>>      return 0;
>>>>  
>>>>    unsigned char kind = kinds[pos+1] & 0xff;
>>>>  
>>>> -  return kind == GOMP_MAP_TO_PSET;
>>>> +  if (kind == GOMP_MAP_TO_PSET)
>>>> +    return 3;
>>>> +  else if (kind == GOMP_MAP_POINTER)
>>>> +    return 2;
>>>> +
>>>> +  return 0;
>>>>  }
>> Is this still necessary with the firstprivatization of subarrays
>> pointers? Well, it might be for fortran. Conceptually, the gimplifier
>> should prune out those unnecessary firstprivate pointer clauses for
>> executable constructs such as enter/exit data and update.
> 
> It appears that GOMP_MAP_POINTER/GOMP_MAP_TO_PSET maps are currently
> created only from the Fortran FE, so I think your description is accurate.
> 
>> Actually, this is one area in the spec where the intent of enter/exit
>> data conflicts with what it describes. If you look at the runtime
>> documentation for, say, acc_create, it states that
>>
>>   acc_create (pvar, n*sizeof(var))
>>
>> is equivalent to
>>
>>   acc enter data create (pvar[n])
>>
>> And to free acc_create, you use acc_delete. So in theory, you should be
>> able to
>>
>>   #pragma acc enter data create (pvar[n])
>>   acc_free (pvar)
>>
>> but this may result in a memory leak if the pointer mapping isn't freed.
> 
> Upon re-reading the OpenACC spec, it appears that acc_malloc/acc_free are supposed
> to be "dumb" allocation/deallocation interfaces, i.e. the implementation is likely
> to be something that directly wires to the alloc_func/free_func plugin hooks.
> I don't think it's supposed to be something that works with the enter/exit data directives,
> or anything that works on the maps managed by libgomp.
> 
> Chung-Lin
> 
> 
> 


Index Nav: [Date Index] [Subject Index] [Author Index] [Thread Index]
Message Nav: [Date Prev] [Date Next] [Thread Prev] [Thread Next]