This is the mail archive of the
gcc-patches@gcc.gnu.org
mailing list for the GCC project.
Re: openacc reference reductions
- From: Jakub Jelinek <jakub at redhat dot com>
- To: Nathan Sidwell <nathan at acm dot org>
- Cc: Cesar Philippidis <cesar at codesourcery dot com>, "gcc-patches at gcc dot gnu dot org" <gcc-patches at gcc dot gnu dot org>
- Date: Fri, 8 Apr 2016 16:21:48 +0200
- Subject: Re: openacc reference reductions
- Authentication-results: sourceware.org; auth=none
- References: <56BA06C3 dot 90606 at acm dot org> <56BA10FC dot 90705 at codesourcery dot com> <56CB2A76 dot 3090809 at codesourcery dot com> <57046C2B dot 6080002 at codesourcery dot com> <20160406142340 dot GZ19207 at tucnak dot redhat dot com> <57056FCA dot 7040602 at codesourcery dot com> <20160407095657 dot GD19207 at tucnak dot redhat dot com> <570734E3 dot 7030601 at codesourcery dot com> <20160408074054 dot GN19207 at tucnak dot redhat dot com> <5707BCB5 dot 2050607 at acm dot org>
- Reply-to: Jakub Jelinek <jakub at redhat dot com>
On Fri, Apr 08, 2016 at 07:14:13AM -0700, Nathan Sidwell wrote:
> On 04/08/16 00:40, Jakub Jelinek wrote:
>
> >>+/* OpenACC parallel reductions need a present_or_copy clause to ensure
> >>+ that the original variable used in the reduction gets updated on
> >>+ the host. Scan the list of clauses for reduction so that any existing
> >>+ data clause can be adjusted if necessary. */
> >>+ if (region_type == ORT_ACC_PARALLEL)
> >>+ {
> >>+ for (c = *list_p; c; c = OMP_CLAUSE_CHAIN (c))
> >>+ {
> >>+ tree decl = NULL_TREE;
> >>+
> >>+ if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_REDUCTION)
> >>+ continue;
> >>+
> >>+ decl = OMP_CLAUSE_DECL (c);
> >>+ omp_add_variable (ctx, decl, GOVD_REDUCTION);
> >>+ }
> >>+ }
> >>+
> >
> >And this looks also wrong, why?
>
> And besides, won't it break
> #pragma acc parallel firstprivate (x) reduction(+:x)
> {...}
>
> which the openacc guys at GTC indicated they didn't want to disallow
> (without further thinking).
Well, given that this patchset is supposed to fix one of the last few P1
blockers, I think we shouldn't be changing the FEs now to allow the above,
just get working what is already allowed, and defer the rest for GCC 7.
Yes, acc reductions are just very weird and to me they look insufficiently
well defined. In OpenMP, reduction is a data sharing clause like
private/firstprivate etc., so one can use either private, or firstprivate,
or reduction, but not multiple of them on the same decl; the exception is
firstprivate and lastprivate on the same decl is allowed. Furthermore,
reduction is not allowed at all on target construct, it can be only on
teams, parallel, for, simd, sections constructs, so the data sharing or
mapping on target is orthogonal on if you reduce on the inner constructs or
not.
Jakub