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