[PR70895] Add copy mapping for reductions on OpenACC loop directives

On 2016/8/16 7:06 PM, Thomas Schwinge wrote:
> Hi!
> 
> On Mon, 15 Aug 2016 19:25:48 +0800, Chung-Lin Tang <cltang@codesourcery.com> wrote:
>> per the discussion on the bugzilla PR page, reductions on OpenACC loop
>> directives will automatically get a copy clause mapping on an enclosing
>> parallel construct (unless bounded by a local variable or an explicit
>> firstprivate clause).
>>
>> There is also a patch for libgomp testsuite cases. Asides from the
>> fortran case which now needs explicit firstprivate clauses to work,
>> other C/C++ cases have been adjusted to remove explicit copy clauses.
>> (I have not exhaustively searched everywhere to eliminate them though)
>>
>> This has been tested using gomp-4_0-branch, which is based on GCC 6,
>> which is what this PR was originally filed for.
>>
>> I will be committing this soon for gomp-4_0-branch,
>> is this okay for gcc-6-branch and trunk as well?
> 
> On Mon, 15 Aug 2016 15:23:14 +0200, Jakub Jelinek <jakub@redhat.com> wrote:
>> The gimplify.c change is ok for trunk and 6.3 (after 6.2 is released).
>> As for the testsuite, I'll leave it to Thomas/Nathan on what they prefer,
>> I'd think that having explicit clauses in e.g. half of the testcases and
>> implicit ones in the other half wouldn't hurt, so that both are tested
>> enough.
> 
> ACK, but from a quick scan it seems as if there's still sufficient
> coverage remaining with explicit usage.
> 
> What I'd like to see changed/added, though, is some libgomp.oacc-fortran
> test coverage of the new implicit copy clauses, and a handful of C/C++ as
> well as Fortran tree-scanning tests in gcc/testsuite/ -- basically, to
> document the expected behavior.

I've added the kernels case assertion you mentioned below, and committed the
gimplify patch to trunk, gcc-6-branch, and gomp-4_0-branch.
I've also added some of those testing items you mentioned, see the attached updated
testsuite patches.

>> +  /* For reductions clauses in OpenACC loop directives, by default create a
>> +     copy clause on the enclosing parallel construct for carrying back the
>> +     results.  */
>> +  if (ctx->region_type == ORT_ACC && (flags & GOVD_REDUCTION))
> 
> Should this be "ctx->region_type & ORT_ACC" instead of "=="?  I suppose
> the same thing also applies to OpenACC nested parallelism:
> 
>     #pragma acc parallel
>       {
>         [...]
>     #pragma acc parallel reduction([...])
>         {
>           [...]
> 
> ..., which we're not supporting right now, but that'll make it easier to
> spot once adding such support.

I'm ignoring this issue for now; this patch only deals with acc loop directives,
hence "== ORT_ACC". For Cesar's code path that deals with adding copy clauses for
parallel construct reductions, see the case in gimple_adjust_omp_clauses().
IMHO, to infer how we'll deal with nested parallelism right now is too long a shot.

> I suppose we can also run into this for ORT_ACC_KERNELS (if not, should
> mark/document that with gcc_unreachable or some such); per OpenACC 2.0a,
> 2.5.2 Kernels Construct, "a scalar variable referenced in the kernels
> construct that does not appear in a data clause for the construct or any
> enclosing data construct will be treated as if it appeared in a copy
> clause", so we should assert that this already is a GOVD_MAP.

Fair enough, I've added that in the gimplify patch.

> Can we also run into this for ORT_ACC_DATA and ORT_ACC_HOST_DATA, but
> nothing needs to be done for these?

Don't we have other mechanisms for enforcing proper nesting? data/host_data
constructs are supposed to be host code. They should not be directly enclosing loop directives,
as what we're concerned here.

The attached are the final patches I committed.

Thanks,
Chung-Lin

2016-08-18  Chung-Lin Tang  <cltang@codesourcery.com>

        PR middle-end/70895
        gcc/
        * gimplify.c (omp_add_variable): Adjust/add variable mapping on
        enclosing parallel construct for reduction variables on OpenACC loop
        directives.

        gcc/testsuite/
        * gfortran.dg/goacc/loop-tree-1.f90: Add gimple scan-tree-dump test.
        * c-c++-common/goacc/reduction-1.c: Likewise.
        * c-c++-common/goacc/reduction-2.c: Likewise.
        * c-c++-common/goacc/reduction-3.c: Likewise.
        * c-c++-common/goacc/reduction-4.c: Likewise.

        libgomp/
        * testsuite/libgomp.oacc-fortran/reduction-7.f90: Add explicit
        firstprivate clauses.
        * testsuite/libgomp.oacc-fortran/reduction-6.f90: Remove explicit
        copy clauses.
        * testsuite/libgomp.oacc-c-c++-common/reduction-7.c: Likewise.
        * testsuite/libgomp.oacc-c-c++-common/reduction-cplx-flt.c: Likewise.
        * testsuite/libgomp.oacc-c-c++-common/reduction-flt.c: Likewise.
        * testsuite/libgomp.oacc-c-c++-common/collapse-2.c: Likewise.
        * testsuite/libgomp.oacc-c-c++-common/loop-red-wv-1.c: Likewise.
        * testsuite/libgomp.oacc-c-c++-common/collapse-4.c: Likewise.
        * testsuite/libgomp.oacc-c-c++-common/loop-red-v-1.c: Likewise.
        * testsuite/libgomp.oacc-c-c++-common/reduction-cplx-dbl.c: Likewise.
        * testsuite/libgomp.oacc-c-c++-common/loop-red-g-1.c: Likewise.
        * testsuite/libgomp.oacc-c-c++-common/loop-red-gwv-1.c: Likewise.
        * testsuite/libgomp.oacc-c-c++-common/loop-red-w-1.c: Likewise.
        * testsuite/libgomp.oacc-c-c++-common/reduction-dbl.c: Likewise.

[PR70895] Add copy mapping for reductions on OpenACC loop directives

Commit Message

Patch