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 <rguent...@suse.de> SUSE LINUX GmbH, GF: Felix Imendoerffer, Jane Smithard, Graham Norton, HRB 21284 (AG Nuernberg)