[4/16] Implement -foffload-alias

On 09/11/15 16:35, Tom de Vries wrote:
> Hi,
>
> this patch series for stage1 trunk adds support to:
> - parallelize oacc kernels regions using parloops, and
> - map the loops onto the oacc gang dimension.
>
> The patch series contains these patches:
>
>       1    Insert new exit block only when needed in
>          transform_to_exit_first_loop_alt
>       2    Make create_parallel_loop return void
>       3    Ignore reduction clause on kernels directive
>       4    Implement -foffload-alias
>       5    Add in_oacc_kernels_region in struct loop
>       6    Add pass_oacc_kernels
>       7    Add pass_dominator_oacc_kernels
>       8    Add pass_ch_oacc_kernels
>       9    Add pass_parallelize_loops_oacc_kernels
>      10    Add pass_oacc_kernels pass group in passes.def
>      11    Update testcases after adding kernels pass group
>      12    Handle acc loop directive
>      13    Add c-c++-common/goacc/kernels-*.c
>      14    Add gfortran.dg/goacc/kernels-*.f95
>      15    Add libgomp.oacc-c-c++-common/kernels-*.c
>      16    Add libgomp.oacc-fortran/kernels-*.f95
>
> The first 9 patches are more or less independent, but patches 10-16 are
> intended to be committed at the same time.
>
> Bootstrapped and reg-tested on x86_64.
>
> Build and reg-tested with nvidia accelerator, in combination with a
> patch that enables accelerator testing (which is submitted at
> https://gcc.gnu.org/ml/gcc-patches/2015-10/msg01771.html ).
>
> I'll post the individual patches in reply to this message.

this patch addresses the problem that once the offloading region has 
been split off from the original function, alias analysis can no longer 
use information available in the original function that would allow it 
to do a more precise analysis for the offloading function. [ At some 
point we could use fipa-pta for that, as discussed in PR46032, but 
that's not feasible now. ]

The basic idea behind the patch is that for typical usage, the base 
pointers used in an offloaded region are non-aliasing. The patch works 
by adding restrict to the types of the fields used to pass data to an 
offloading region.


The patch implements a new option
-foffload-alias=<none|pointer|all>.

The option -foffload-alias=none instructs the compiler to assume that
object references and pointer dereferences in an offload region do not
alias.

The option -foffload-alias=pointer instructs the compiler to assume that 
objects references in an offload region do not alias.

The option -foffload-alias=all instructs the compiler to make no
assumptions about aliasing in offload regions.

The default value is -foffload-alias=none.

Thanks,
- Tom

On Mon, 9 Nov 2015, Tom de Vries wrote:

> On 09/11/15 16:35, Tom de Vries wrote:
> > Hi,
> > 
> > this patch series for stage1 trunk adds support to:
> > - parallelize oacc kernels regions using parloops, and
> > - map the loops onto the oacc gang dimension.
> > 
> > The patch series contains these patches:
> > 
> >       1    Insert new exit block only when needed in
> >          transform_to_exit_first_loop_alt
> >       2    Make create_parallel_loop return void
> >       3    Ignore reduction clause on kernels directive
> >       4    Implement -foffload-alias
> >       5    Add in_oacc_kernels_region in struct loop
> >       6    Add pass_oacc_kernels
> >       7    Add pass_dominator_oacc_kernels
> >       8    Add pass_ch_oacc_kernels
> >       9    Add pass_parallelize_loops_oacc_kernels
> >      10    Add pass_oacc_kernels pass group in passes.def
> >      11    Update testcases after adding kernels pass group
> >      12    Handle acc loop directive
> >      13    Add c-c++-common/goacc/kernels-*.c
> >      14    Add gfortran.dg/goacc/kernels-*.f95
> >      15    Add libgomp.oacc-c-c++-common/kernels-*.c
> >      16    Add libgomp.oacc-fortran/kernels-*.f95
> > 
> > The first 9 patches are more or less independent, but patches 10-16 are
> > intended to be committed at the same time.
> > 
> > Bootstrapped and reg-tested on x86_64.
> > 
> > Build and reg-tested with nvidia accelerator, in combination with a
> > patch that enables accelerator testing (which is submitted at
> > https://gcc.gnu.org/ml/gcc-patches/2015-10/msg01771.html ).
> > 
> > I'll post the individual patches in reply to this message.
> 
> this patch addresses the problem that once the offloading region has been
> split off from the original function, alias analysis can no longer use
> information available in the original function that would allow it to do a
> more precise analysis for the offloading function. [ At some point we could
> use fipa-pta for that, as discussed in PR46032, but that's not feasible now. ]
> 
> The basic idea behind the patch is that for typical usage, the base pointers
> used in an offloaded region are non-aliasing. The patch works by adding
> restrict to the types of the fields used to pass data to an offloading region.
> 
> 
> The patch implements a new option
> -foffload-alias=<none|pointer|all>.
> 
> The option -foffload-alias=none instructs the compiler to assume that
> object references and pointer dereferences in an offload region do not
> alias.
> 
> The option -foffload-alias=pointer instructs the compiler to assume that
> objects references in an offload region do not alias.
> 
> The option -foffload-alias=all instructs the compiler to make no
> assumptions about aliasing in offload regions.
> 
> The default value is -foffload-alias=none.

I think global options for this is nonsense.  Please follow what
we do for #pragma GCC ivdep for example, thus allow the alias
behavior to be specified per "region" (whatever makes sense here
in the context of offloading).

Thanks,
Richard.

> Thanks,
> - Tom
> 
>

On Wed, Nov 11, 2015 at 11:51:02AM +0100, Richard Biener wrote:
> > The option -foffload-alias=pointer instructs the compiler to assume that
> > objects references in an offload region do not alias.
> > 
> > The option -foffload-alias=all instructs the compiler to make no
> > assumptions about aliasing in offload regions.
> > 
> > The default value is -foffload-alias=none.
> 
> I think global options for this is nonsense.  Please follow what
> we do for #pragma GCC ivdep for example, thus allow the alias
> behavior to be specified per "region" (whatever makes sense here
> in the context of offloading).

Yeah, completely agreed.  I don't see why the offloaded region would be in
any way special, they are C/C++/Fortran code as any other.
What we can and should improve is teach IPA aliasing/points to analysis
about the way we lower the host vs. offloading region boundary, so that
if alias analysis on the caller of GOMP_target_ext/GOACC_parallel_keyed
determines something it can be used on the offloaded function side and vice
versa, but a switch like the above is just wrong.

	Jakub

On 11/11/15 12:00, Jakub Jelinek wrote:
> On Wed, Nov 11, 2015 at 11:51:02AM +0100, Richard Biener wrote:
>>> The option -foffload-alias=pointer instructs the compiler to assume that
>>> objects references in an offload region do not alias.
>>>
>>> The option -foffload-alias=all instructs the compiler to make no
>>> assumptions about aliasing in offload regions.
>>>
>>> The default value is -foffload-alias=none.
>>
>> I think global options for this is nonsense.  Please follow what
>> we do for #pragma GCC ivdep for example, thus allow the alias
>> behavior to be specified per "region" (whatever makes sense here
>> in the context of offloading).

So, IIUC, instead of a global option foffload-alias, you're saying 
something like the following would be acceptable:
...
#pragma GCC offload-alias=<none|pointer|all>
#pragma acc kernels copyin (a[0:N], b[0:N]) copyout (c[0:N])
   {
     #pragma acc loop
     for (COUNTERTYPE ii = 0; ii < N; ii++)
       c[ii] = a[ii] + b[ii];
   }
...
?

I suppose that would work (though a global option would allow us to 
easily switch between none/pointer/all values in a large number of 
files, something that might be useful when f.i. running an openacc  test 
suite).

> Yeah, completely agreed.  I don't see why the offloaded region would be in
> any way special, they are C/C++/Fortran code as any other.
> What we can and should improve is teach IPA aliasing/points to analysis
> about the way we lower the host vs. offloading region boundary, so that
> if alias analysis on the caller of GOMP_target_ext/GOACC_parallel_keyed
> determines something it can be used on the offloaded function side and vice
> versa,

I agree this would be a nice way to solve the aliasing info problem, but 
considering the remark of Richard at 
https://gcc.gnu.org/bugzilla/show_bug.cgi?id=46032#c19 :
...
Not that I think IPA PTA is anywhere near production ready
...
I haven't considered proceeding in that direction.

Thanks,
- Tom

> but a switch like the above is just wrong.

On Thu, 12 Nov 2015, Tom de Vries wrote:

> On 11/11/15 12:00, Jakub Jelinek wrote:
> > On Wed, Nov 11, 2015 at 11:51:02AM +0100, Richard Biener wrote:
> > > > The option -foffload-alias=pointer instructs the compiler to assume that
> > > > objects references in an offload region do not alias.
> > > > 
> > > > The option -foffload-alias=all instructs the compiler to make no
> > > > assumptions about aliasing in offload regions.
> > > > 
> > > > The default value is -foffload-alias=none.
> > > 
> > > I think global options for this is nonsense.  Please follow what
> > > we do for #pragma GCC ivdep for example, thus allow the alias
> > > behavior to be specified per "region" (whatever makes sense here
> > > in the context of offloading).
> 
> So, IIUC, instead of a global option foffload-alias, you're saying something
> like the following would be acceptable:
> ...
> #pragma GCC offload-alias=<none|pointer|all>
> #pragma acc kernels copyin (a[0:N], b[0:N]) copyout (c[0:N])
>   {
>     #pragma acc loop
>     for (COUNTERTYPE ii = 0; ii < N; ii++)
>       c[ii] = a[ii] + b[ii];
>   }
> ...
> ?
> 
> I suppose that would work (though a global option would allow us to easily
> switch between none/pointer/all values in a large number of files, something
> that might be useful when f.i. running an openacc  test suite).
> 
> > Yeah, completely agreed.  I don't see why the offloaded region would be in
> > any way special, they are C/C++/Fortran code as any other.
> > What we can and should improve is teach IPA aliasing/points to analysis
> > about the way we lower the host vs. offloading region boundary, so that
> > if alias analysis on the caller of GOMP_target_ext/GOACC_parallel_keyed
> > determines something it can be used on the offloaded function side and vice
> > versa,
> 
> I agree this would be a nice way to solve the aliasing info problem, but
> considering the remark of Richard at
> https://gcc.gnu.org/bugzilla/show_bug.cgi?id=46032#c19 :
> ...
> Not that I think IPA PTA is anywhere near production ready

Just to clarify on that sentence:
 1) we lack good testing coverage for IPA PTA so wrong-code bugs might 
still exist
 2) IPA PTA can use a _lot_ of memory and compile-time
 3) for existing wrong-code issues I have merely dumbed down the
use of the analysis result resulting in weaker alias analysis compared to
the local PTA (for some cases)

Because of 2) and no good way to avoid this I decided to not make
fixing 3) a priority (and 1) still holds).

Richard.

> ...
> I haven't considered proceeding in that direction.
> 
> Thanks,
> - Tom
> 
> > but a switch like the above is just wrong.
> 
>

On 13/11/15 09:46, Richard Biener wrote:
> On Thu, 12 Nov 2015, Tom de Vries wrote:
>
>> On 11/11/15 12:00, Jakub Jelinek wrote:
>>> On Wed, Nov 11, 2015 at 11:51:02AM +0100, Richard Biener wrote:
>>>>> The option -foffload-alias=pointer instructs the compiler to assume that
>>>>> objects references in an offload region do not alias.
>>>>>
>>>>> The option -foffload-alias=all instructs the compiler to make no
>>>>> assumptions about aliasing in offload regions.
>>>>>
>>>>> The default value is -foffload-alias=none.
>>>>
>>>> I think global options for this is nonsense.  Please follow what
>>>> we do for #pragma GCC ivdep for example, thus allow the alias
>>>> behavior to be specified per "region" (whatever makes sense here
>>>> in the context of offloading).
>>
>> So, IIUC, instead of a global option foffload-alias, you're saying something
>> like the following would be acceptable:
>> ...
>> #pragma GCC offload-alias=<none|pointer|all>
>> #pragma acc kernels copyin (a[0:N], b[0:N]) copyout (c[0:N])
>>    {
>>      #pragma acc loop
>>      for (COUNTERTYPE ii = 0; ii < N; ii++)
>>        c[ii] = a[ii] + b[ii];
>>    }
>> ...
>> ?
>>
>> I suppose that would work (though a global option would allow us to easily
>> switch between none/pointer/all values in a large number of files, something
>> that might be useful when f.i. running an openacc  test suite).
>>
>>> Yeah, completely agreed.  I don't see why the offloaded region would be in
>>> any way special, they are C/C++/Fortran code as any other.
>>> What we can and should improve is teach IPA aliasing/points to analysis
>>> about the way we lower the host vs. offloading region boundary, so that
>>> if alias analysis on the caller of GOMP_target_ext/GOACC_parallel_keyed
>>> determines something it can be used on the offloaded function side and vice
>>> versa,
>>
>> I agree this would be a nice way to solve the aliasing info problem, but
>> considering the remark of Richard at
>> https://gcc.gnu.org/bugzilla/show_bug.cgi?id=46032#c19 :
>> ...
>> Not that I think IPA PTA is anywhere near production ready
>
> Just to clarify on that sentence:
>   1) we lack good testing coverage for IPA PTA so wrong-code bugs might
> still exist
>   2) IPA PTA can use a _lot_ of memory and compile-time
>   3) for existing wrong-code issues I have merely dumbed down the
> use of the analysis result resulting in weaker alias analysis compared to
> the local PTA (for some cases)
>
> Because of 2) and no good way to avoid this I decided to not make
> fixing 3) a priority (and 1) still holds).
>

Hi,

thanks for the explanation. Filed as PR68331 - '[meta-bug] fipa-pta issues'.

Any feedback on the '#pragma GCC offload-alias=<none|pointer|all>' bit 
above? Is that sort of what you had in mind?

Thanks,
- Tom

On Fri, 13 Nov 2015, Tom de Vries wrote:

> On 13/11/15 09:46, Richard Biener wrote:
> > On Thu, 12 Nov 2015, Tom de Vries wrote:
> > 
> > > On 11/11/15 12:00, Jakub Jelinek wrote:
> > > > On Wed, Nov 11, 2015 at 11:51:02AM +0100, Richard Biener wrote:
> > > > > > The option -foffload-alias=pointer instructs the compiler to assume
> > > > > > that
> > > > > > objects references in an offload region do not alias.
> > > > > > 
> > > > > > The option -foffload-alias=all instructs the compiler to make no
> > > > > > assumptions about aliasing in offload regions.
> > > > > > 
> > > > > > The default value is -foffload-alias=none.
> > > > > 
> > > > > I think global options for this is nonsense.  Please follow what
> > > > > we do for #pragma GCC ivdep for example, thus allow the alias
> > > > > behavior to be specified per "region" (whatever makes sense here
> > > > > in the context of offloading).
> > > 
> > > So, IIUC, instead of a global option foffload-alias, you're saying
> > > something
> > > like the following would be acceptable:
> > > ...
> > > #pragma GCC offload-alias=<none|pointer|all>
> > > #pragma acc kernels copyin (a[0:N], b[0:N]) copyout (c[0:N])
> > >    {
> > >      #pragma acc loop
> > >      for (COUNTERTYPE ii = 0; ii < N; ii++)
> > >        c[ii] = a[ii] + b[ii];
> > >    }
> > > ...
> > > ?
> > > 
> > > I suppose that would work (though a global option would allow us to easily
> > > switch between none/pointer/all values in a large number of files,
> > > something
> > > that might be useful when f.i. running an openacc  test suite).
> > > 
> > > > Yeah, completely agreed.  I don't see why the offloaded region would be
> > > > in
> > > > any way special, they are C/C++/Fortran code as any other.
> > > > What we can and should improve is teach IPA aliasing/points to analysis
> > > > about the way we lower the host vs. offloading region boundary, so that
> > > > if alias analysis on the caller of GOMP_target_ext/GOACC_parallel_keyed
> > > > determines something it can be used on the offloaded function side and
> > > > vice
> > > > versa,
> > > 
> > > I agree this would be a nice way to solve the aliasing info problem, but
> > > considering the remark of Richard at
> > > https://gcc.gnu.org/bugzilla/show_bug.cgi?id=46032#c19 :
> > > ...
> > > Not that I think IPA PTA is anywhere near production ready
> > 
> > Just to clarify on that sentence:
> >   1) we lack good testing coverage for IPA PTA so wrong-code bugs might
> > still exist
> >   2) IPA PTA can use a _lot_ of memory and compile-time
> >   3) for existing wrong-code issues I have merely dumbed down the
> > use of the analysis result resulting in weaker alias analysis compared to
> > the local PTA (for some cases)
> > 
> > Because of 2) and no good way to avoid this I decided to not make
> > fixing 3) a priority (and 1) still holds).
> > 
> 
> Hi,
> 
> thanks for the explanation. Filed as PR68331 - '[meta-bug] fipa-pta issues'.
> 
> Any feedback on the '#pragma GCC offload-alias=<none|pointer|all>' bit above?
> Is that sort of what you had in mind?

Yes.  Whether that makes sense is another question of course.  You can
annotate memory references with MR_DEPENDENCE_BASE/CLIQUE yourself
as well if you know dependences without the users intervention.

Richard.

> Thanks,
> - Tom
> 
> 
> 
>

On Fri, Nov 13, 2015 at 12:29:51PM +0100, Richard Biener wrote:
> > thanks for the explanation. Filed as PR68331 - '[meta-bug] fipa-pta issues'.
> > 
> > Any feedback on the '#pragma GCC offload-alias=<none|pointer|all>' bit above?
> > Is that sort of what you had in mind?
> 
> Yes.  Whether that makes sense is another question of course.  You can
> annotate memory references with MR_DEPENDENCE_BASE/CLIQUE yourself
> as well if you know dependences without the users intervention.

I really don't like even the GCC offload-alias, I just don't see anything
special on the offload code.  Not to mention that the same issue is already
with other outlined functions, like OpenMP tasks or parallel regions, those
aren't offloaded, yet they can suffer from worse alias/points-to analysis
too.

We simply have some compiler internal interface between the caller and
callee of the outlined regions, each interface in between those has
its own structure type used to communicate the info;
we can attach attributes on the fields, or some flags to indicate some
properties interesting from aliasing POV.  We don't really need to perform
full IPA-PTA, perhaps it would be enough to a) record somewhere in cgraph
the relationship in between such callers and callees (for offloading regions
we already have "omp target entrypoint" attribute on the callee and a
singler caller), tell LTO if possible not to split those into different
partitions if easily possible, and then just for these pairs perform
aliasing/points-to analysis in the caller and the result record using
cliques/special attributes/whatever to the callee side, so that the callee
(outlined OpenMP/OpenACC/Cilk+ region) can then improve its alias analysis.

	Jakub

On 11/11/15 12:00, Jakub Jelinek wrote:
> On Wed, Nov 11, 2015 at 11:51:02AM +0100, Richard Biener wrote:
>>> The option -foffload-alias=pointer instructs the compiler to assume that
>>> objects references in an offload region do not alias.
>>>
>>> The option -foffload-alias=all instructs the compiler to make no
>>> assumptions about aliasing in offload regions.
>>>
>>> The default value is -foffload-alias=none.
>>
>> I think global options for this is nonsense.  Please follow what
>> we do for #pragma GCC ivdep for example, thus allow the alias
>> behavior to be specified per "region" (whatever makes sense here
>> in the context of offloading).
>
> Yeah, completely agreed.  I don't see why the offloaded region would be in
> any way special, they are C/C++/Fortran code as any other.
> What we can and should improve is teach IPA aliasing/points to analysis
> about the way we lower the host vs. offloading region boundary, so that
> if alias analysis on the caller of GOMP_target_ext/GOACC_parallel_keyed
> determines something it can be used on the offloaded function side and vice
> versa, but a switch like the above is just wrong.

Filed the GOMP_target_ext bit as PR 68675 - Handle GOMP_target_ext 
optimally in ipa-pta.

Thanks,
- Tom