[PATCH] Handle BUILT_IN_GOACC_PARALLEL in ipa-pta

[Tom de Vries]

On 01/12/15 15:44, Jakub Jelinek wrote:
> On Tue, Dec 01, 2015 at 03:25:42PM +0100, Tom de Vries wrote:
>> Handle BUILT_IN_GOACC_PARALLEL in ipa-pta
>>
>> 2015-12-01  Tom de Vries  <tom@codesourcery.com>
>>
>> 	* tree-ssa-structalias.c (find_func_aliases_for_builtin_call)
>> 	(find_func_clobbers, ipa_pta_execute): Handle BUILT_IN_GOACC_PARALLEL.
>
> Isn't this cheating though?  The kernel will be called with those addresses
> only if doing host fallback

Let's take a look at goacc/kernels-alias-ipa-pta.c:
...
unsigned int a[N];
unsigned int b[N];
unsigned int c[N];

#pragma acc kernels pcopyout (a, b, c)
{
   a[0] = 0;
   b[0] = 1;
   c[0] = a[0];
}
...

If we execute on the host, the a, b and c used in the kernels region 
will be the a, b and c declared outside the region.

If we execute on a non-shared mem accelerator, the a, b and c used in 
the kernels region will be copies of a, b and c in the accelerator 
memory: a.1, b.1 and c.1.

This patch tells ipa-pta (which has no notion of a.1, b.1 and c.1) that 
we're using declared a, b, and c, in the kernels region, while on the 
accelerator we're really using a.1, b.1 and c.1, so in that sense it's 
cheating.

However, given that declared a, b and c are disjunct, we know that their 
copies will be disjunct, so by pretending that declared a, b and c are 
used in the kernels region, we get conclusions which are also valid when 
we use a.1, b.1 and c.1 instead in the kernels region.

So, for this patch to be incorrect we have to find an example where 
ipa-pta finds that two memory references are not aliasing, while on the 
accelerator those memory references are really aliasing. AFAICT there 
are no such examples.

> (and for GOMP_target_ext even not for that
> always - firstprivate vars will have the addresses replaced by addresses of
> alloca-ed copies of those objects).

I don't think firstprivate vars is a problem, I think the opposite would 
a problem: merging vars on the accelerator which are disjunct on the host.

> I haven't studied in detail what exactly IPA-PTA does, so maybe it is good
> enough to pretend that.

AFAIU, it's good enough, because the points-to information is only used 
to prove non-aliases.

Does this explanation address your concern?

Thanks,
- Tom