This is the mail archive of the
gcc-patches@gcc.gnu.org
mailing list for the GCC project.
Re: [patch, libgomp, OpenACC] Additional enter/exit data map handling
- From: Chung-Lin Tang <chunglin_tang at mentor dot com>
- To: Cesar Philippidis <cesar at codesourcery dot com>, Chung-Lin Tang <cltang at codesourcery dot com>, gcc-patches <gcc-patches at gcc dot gnu dot org>, Jakub Jelinek <jakub at redhat dot com>
- Cc: Thomas Schwinge <thomas_schwinge at mentor dot com>
- Date: Tue, 6 Dec 2016 23:49:19 +0800
- Subject: Re: [patch, libgomp, OpenACC] Additional enter/exit data map handling
- Authentication-results: sourceware.org; auth=none
- References: <66f8ce3d-7206-ee8a-abaa-4bb25423e4eb@codesourcery.com> <b0d4e474-8431-a166-1e0b-c9344fe53cc5@codesourcery.com> <c4559e7f-796c-2071-5f7e-be3480cab72c@mentor.com>
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
>
>
>