On Thu, Oct 22, 2015 at 09:08:30AM -0400, Nathan Sidwell wrote: > On 10/22/15 07:10, Julian Brown wrote: > >On Thu, 22 Oct 2015 10:05:30 +0200 > >Richard Biener <richard.guent...@gmail.com> wrote: > > >>So you'd need to be more precise as to what properties you are trying > >>to preserve by placing a single stmt somewhere. > > > >FWIW an earlier, abandoned attempt at solving the same problem was > >discussed in the following thread, continuing through June: > > > > https://gcc.gnu.org/ml/gcc-patches/2015-05/msg02647.html > > > >Though the details of lowering of OpenACC constructs have changed with > >Nathan's current patches, the underlying problem remains the same. PTX > >requires certain operations (bar.sync) to be executed uniformly by all > >threads in a CTA. IIUC this affects "JOIN" points across all > >workers/vectors in a gang, in particular (though this is generic code, > >other -- particularly GPU -- targets may have similar restrictions). > > > Richard, does this answer your question?
I agree with Richard that it would be better to write more about what kind of IL changes are acceptable with IFN_UNIQUE in the IL and what are not. E.g. is inlining ok (I'd hope yes)? Is function splitting ok (bet as long as all IFN_UNIQUE calls stay in one or the other part, but not both)? Various loop optimization, ... Jakub