| Message ID | 563F6E2D.7070107@mentor.com |
|---|---|
| State | New |
| Headers | show |
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 >> <cesar@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
On 19/11/15 01:03, Cesar Philippidis wrote: > 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 >>> <cesar@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. > Hi Cesar, that patch was submitted for trunk as part of the kernels support: https://gcc.gnu.org/ml/gcc-patches/2015-11/msg00994.html Thanks, - Tom
Ignore reduction clause on kernels directive 2015-11-08 Tom de Vries <tom@codesourcery.com> * c-omp.c (c_oacc_split_loop_clauses): Don't copy OMP_CLAUSE_REDUCTION, classify as loop clause. --- gcc/c-family/c-omp.c | 9 ++------- 1 file changed, 2 insertions(+), 7 deletions(-) diff --git a/gcc/c-family/c-omp.c b/gcc/c-family/c-omp.c index 8b30844..907d329 100644 --- a/gcc/c-family/c-omp.c +++ b/gcc/c-family/c-omp.c @@ -867,7 +867,7 @@ c_omp_check_loop_iv_exprs (location_t stmt_loc, tree declv, tree decl, tree c_oacc_split_loop_clauses (tree clauses, tree *not_loop_clauses) { - tree next, loop_clauses, t; + tree next, loop_clauses; loop_clauses = *not_loop_clauses = NULL_TREE; for (; clauses ; clauses = next) @@ -886,16 +886,11 @@ c_oacc_split_loop_clauses (tree clauses, tree *not_loop_clauses) case OMP_CLAUSE_SEQ: case OMP_CLAUSE_INDEPENDENT: case OMP_CLAUSE_PRIVATE: + case OMP_CLAUSE_REDUCTION: OMP_CLAUSE_CHAIN (clauses) = loop_clauses; loop_clauses = clauses; break; - /* Reductions belong in both constructs. */ - case OMP_CLAUSE_REDUCTION: - t = copy_node (clauses); - OMP_CLAUSE_CHAIN (t) = loop_clauses; - loop_clauses = t; - /* FIXME: device_type */ /* Parallel/kernels clauses. */ -- 1.9.1