This is the mail archive of the gcc-patches@gcc.gnu.org mailing list for the GCC project.


Index Nav: [Date Index] [Subject Index] [Author Index] [Thread Index]
Message Nav: [Date Prev] [Date Next] [Thread Prev] [Thread Next]
Other format: [Raw text]

Re: gomp-nvptx branch - middle-end changes


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.

(a few days ago I've sent you privately a patch implementing the above)

Thanks.
Alexander


Index Nav: [Date Index] [Subject Index] [Author Index] [Thread Index]
Message Nav: [Date Prev] [Date Next] [Thread Prev] [Thread Next]