Re: [openacc] tile, independent, default, private and firstprivate support in c/++

[Thomas Schwinge <thomas at codesourcery dot com>]

Hi!

Nathan, question here about clause splitting for combined OpenACC
"parallel loop" and "kernels loop" constructs, in particular:

    #pragma acc parallel loop reduction([...])

On Thu, 5 Nov 2015 06:47:58 -0800, Cesar Philippidis <cesar@codesourcery.com> wrote:
> On 11/05/2015 04:14 AM, Thomas Schwinge wrote:
> 
> > On Tue, 3 Nov 2015 14:16:59 -0800, Cesar Philippidis <cesar@codesourcery.com> wrote:
> >> This patch does the following to the c and c++ front ends:
> > 
> >>  * updates c_oacc_split_loop_clauses to filter out the loop clauses
> >>    from combined parallel/kernels loops
> > 
> >> 	gcc/c-family/
> >> 	* c-omp.c (c_oacc_split_loop_clauses): Make TILE, GANG, WORKER, VECTOR,
> >> 	AUTO, SEQ and independent as loop clauses.  Associate REDUCTION
> >> 	clauses with parallel and kernels.
> > 
> >> --- a/gcc/c-family/c-omp.c
> >> +++ b/gcc/c-family/c-omp.c
> >> @@ -709,12 +709,21 @@ c_oacc_split_loop_clauses (tree clauses, tree *not_loop_clauses)
> >>  
> >>        switch (OMP_CLAUSE_CODE (clauses))
> >>          {
> >> +	  /* Loop clauses.  */
> >>  	case OMP_CLAUSE_COLLAPSE:
> >> -	case OMP_CLAUSE_REDUCTION:
> >> +	case OMP_CLAUSE_TILE:
> >> +	case OMP_CLAUSE_GANG:
> >> +	case OMP_CLAUSE_WORKER:
> >> +	case OMP_CLAUSE_VECTOR:
> >> +	case OMP_CLAUSE_AUTO:
> >> +	case OMP_CLAUSE_SEQ:
> >> +	case OMP_CLAUSE_INDEPENDENT:
> >>  	  OMP_CLAUSE_CHAIN (clauses) = loop_clauses;
> >>  	  loop_clauses = clauses;
> >>  	  break;
> >>  
> >> +	  /* Parallel/kernels clauses.  */
> >> +
> >>  	default:
> >>  	  OMP_CLAUSE_CHAIN (clauses) = *not_loop_clauses;
> >>  	  *not_loop_clauses = clauses;
> > 
> > Contrary to your ChangeLog entry, this is not duplicating but is moving
> > OMP_CLAUSE_REDUCTION handling.  Is that intentional?  (And, doesn't
> > anything break in the testsuite?)
> 
> Sorry, I must have mis-phrased it. The spec is unclear here. There are
> three possible ways to interpret 'acc parallel loop reduction':
> 
>   1. acc parallel reduction
>      acc loop

This is what you propose in your patch, but I don't think that makes
sense, or does it?  I'm happy to learn otherwise, but in my current
understanding, a reduction clause needs to be attached (at least) to the
innermost construct where reductions are to be processed.  (Let's also
consider multi-level gang/worker/vector loops/reductions.)  So, either:

>   2. acc parallel
>      acc loop reduction

... this, or even this:

>   3. acc parallel reduction
>      acc loop reduction

..., which I'm not sure what the execution model implementation requires.
(Nathan?)

And while we're at it: the very same question also applies to the private
clause, which -- contrary to all other (as far as I remember) clauses --
also is applicable to both the parallel and loop constructs:

    #pragma acc parallel loop private([...])

... is to be decomposed into which of the following:

    #pragma acc parallel private([...])
    #pragma acc loop

    #pragma acc parallel
    #pragma acc loop private([...])

    #pragma acc parallel private([...])
    #pragma acc loop private([...])

(There is no private clause allowed to be specified with the kernels
construct for what it's worth, but that doesn't mean we couldn't use it
internally, of course, if so required.)


> You told me to make all of the front ends consistent, and since I
> started working on fortran first, I had c and c++ follow what it was doing.

Well, I had not asked you to "blindly" ;-) make the front ends
consistent, but rather to figure out why currently C/C++ vs. Fortran are
doing things differently, and then make that consistent, as needed.

> I haven't observed any regressions with this in in place. Then again,
> maybe we don't have sufficient test coverage. I'll do more testing.

Thanks!  (I do see some -- but not very many, curiously! -- regressions
when doing (only) this change on gomp-4_0-branch.)

As far as I'm concerned, this analysis/testing can also be done later on.
(Jakub?)  In that case, for now please just keep OMP_CLAUSE_REDUCTION as
it has been handled before.


GrÃÃe
 Thomas

Attachment: signature.asc
Description: PGP signature