On Fri, Nov 11, 2016 at 12:28:16PM +0300, Alexander Monakov wrote:
> On Fri, 11 Nov 2016, Jakub Jelinek wrote:
>
> > On Fri, Nov 11, 2016 at 11:52:58AM +0300, Alexander Monakov wrote:
> > > On Fri, 11 Nov 2016, Jakub Jelinek wrote:
> > > [...]
> > > > the intended outlining of SIMT regions for PTX offloading done (IMHO the
> > > > best place to do that is in omp expansion, not gimplification)
> > >
> > > Sorry, I couldn't find a good way to implement that during omp expansion.
> > > The
> > > reason I went for gimplification is automatic discovery of sharing
> > > clauses -
> > > I'm assuming in expansion it's very hard to try and fill omp_data_[sio]
> > > without
> > > gimplifier's help. Does this sound sensible?
> >
> > Sure, for discovery of needed sharing clauses the gimplifier has the right
> > infrastructure. But that doesn't mean you can't add those clauses at
> > gimplification time and do the outlining at omp expansion time.
> > That is what is done for omp parallel, task etc. as well. If the standard
> > OpenMP clauses can't serve that purpose, there is always the possibility of
> > adding further internal clauses, that would e.g. be only considered for the
> > SIMT stuff. For the outlining, our current infrastructure really wants to
> > have CFG etc., something you don't have at gimplification time.
>
> Yes, that is exactly what I'm doing. I'm first tweaking the gimplifier to
> inject
> a parallel region with an artificial _simtreg_ clause, transforming
>
> #pragma omp simd
> for (...)
>
> into
>
> #pragma omp parallel _simtreg_
> #pragma omp simd
> for (...)
>
> and then expansion of 'omp parallel' can check presence of _simtreg_ clause
> and
> emit a direct call rather than an invocation of GOMP_parallel.
Well, I meant keep #pragma omp simd as is, just add some data-sharing-like
clauses _simt_shared_(x) or whatever you need, then the omplower versioning
patch I've posted could e.g. drop those _simt_shared_ or whatever else you
need clauses for the omp simd without _simt_ clause, omp lowering then would
do whatever is needed for those _simt_shared_ clauses and finally omp
expansion would outline it. Adding omp parallel around the omp simd is just
weird, it has nothing to do with omp parallel.
Jakub
