On 11/08/2015 07:45 AM, Tom de Vries wrote: > On 07/11/15 12:45, Thomas Schwinge wrote: >> Hi! >> >> On Fri, 6 Nov 2015 15:31:23 -0800, Cesar Philippidis >> <ce...@codesourcery.com> wrote: >>> I've applied this patch to gomp-4_0-branch which backports most of my >>> front end changes from trunk. Note that I found a regression while >>> testing, which is also present in trunk. It looks like >>> kernels-acc-loop-reduction.c is failing because I'm incorrectly >>> propagating the reduction variable to both to the kernels and loop >>> constructs for combined 'acc kernels loop'. The problem here is that >>> kernels don't support the reduction clause. I'll fix that next week. >> >> Always need to consider both what the specification allows -- and thus >> what the front ends accept/refuse -- as well as what we might do >> differently, internally in later processing stages. I have not analyzed >> whether it makes sense to have the OMP_CLAUSE_REDUCTION of a combined >> "kernels loop reduction([...])" construct be attached to the outer >> OACC_KERNELS or inner OACC_LOOP, or duplicated for both. >> >> Tom, if you need a solution for that right now/want to restore the >> previous behavior (attached to innter OACC_LOOP only), here's what you >> should try: in gcc/c-family/c-omp.c:c_oacc_split_loop_clauses remove the >> special handling for OMP_CLAUSE_REDUCTION, and move it to "Loop clauses" >> section, > > Committed to gomp-4_0-branch, as attached.
Can you port this patch to trunk? Originally we were attaching the reduction clause to both the acc loop and parallel construct so that the reduction variable would get a copy clause implicitly. However, Nathan later interpreted #pragma acc parallel reduction(+:var) as #pragma acc parallel reduction(+:var) private(var) Therefore, the burden is on the user to ensure that 'var' is transferred to the parallel region in an appropriate data clause. As a result, we only need to associate reductions with loops now. So your patch is good for trunk. Cesar
