Re: [PATCH, 4/16] Implement -foffload-alias

[Richard Biener <rguenther at suse dot de>]

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
> 
> 
> 
> 

-- 
Richard Biener <rguenther@suse.de>
SUSE LINUX GmbH, GF: Felix Imendoerffer, Jane Smithard, Graham Norton, HRB 21284 (AG Nuernberg)