Message ID | 56583E91.4090603@mentor.com |
---|---|
State | New |
Headers | show |
On Fri, Nov 27, 2015 at 12:29:21PM +0100, Tom de Vries wrote: > Fix oacc kernels default mapping for scalars > > 2015-11-27 Tom de Vries <tom@codesourcery.com> > > * gimplify.c (enum gimplify_omp_var_data): Add enum value > GOVD_MAP_FORCE. > (oacc_default_clause): Fix default for scalars in oacc kernels. > (gimplify_adjust_omp_clauses_1): Handle GOVD_MAP_FORCE. > > * c-c++-common/goacc/kernels-default-2.c: New test. > * c-c++-common/goacc/kernels-default.c: New test. > > --- > gcc/gimplify.c | 19 ++++++++++++++----- > gcc/testsuite/c-c++-common/goacc/kernels-default-2.c | 17 +++++++++++++++++ > gcc/testsuite/c-c++-common/goacc/kernels-default.c | 14 ++++++++++++++ > 3 files changed, 45 insertions(+), 5 deletions(-) > > diff --git a/gcc/gimplify.c b/gcc/gimplify.c > index fcac745..68d90bf 100644 > --- a/gcc/gimplify.c > +++ b/gcc/gimplify.c > @@ -87,6 +87,9 @@ enum gimplify_omp_var_data > /* Flag for GOVD_MAP, if it is always, to or always, tofrom mapping. */ > GOVD_MAP_ALWAYS_TO = 65536, > > + /* Flag for GOVD_MAP, if it is a forced mapping. */ > + GOVD_MAP_FORCE = 131072, The patch has been outdated already when posted, there is GOVD_WRITTEN at this spot. Once you fix this, it is ok for trunk. Jakub
Hi! Copying Nathan for your information, in case you had not yet seen that gcc/gimplify.c:oacc_default_clause change: On Fri, 27 Nov 2015 12:29:21 +0100, Tom de Vries <Tom_deVries@mentor.com> wrote: > The OpenACC 2.0a standard says this about the default mapping for > variables used in a kernels region: > ... > An array or variable of aggregate data type 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 > present_or_copy clause for the 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. > ... > > But atm, all variables including the scalar ones have 'present_or_copy' > defaults. > > This patch makes sure scalar variables get the 'copy' default. > > Bootstrapped and reg-tested on x86_64. OK for stage3 trunk? I see: PASS: gfortran.dg/goacc/reduction-2.f95 -O scan-tree-dump-times gimple "acc loop private.k. reduction..:a." 1 PASS: gfortran.dg/goacc/reduction-2.f95 -O scan-tree-dump-times gimple "acc loop private.p. reduction..:a." 1 [-PASS:-]{+FAIL:+} gfortran.dg/goacc/reduction-2.f95 -O scan-tree-dump-times gimple "target oacc_kernels map.tofrom:a .len: 4.." 1 PASS: gfortran.dg/goacc/reduction-2.f95 -O scan-tree-dump-times gimple "target oacc_parallel firstprivate.a." 1 PASS: gfortran.dg/goacc/reduction-2.f95 -O (test for excess errors) I guess that one needs to be updated? > Fix oacc kernels default mapping for scalars > > 2015-11-27 Tom de Vries <tom@codesourcery.com> > > * gimplify.c (enum gimplify_omp_var_data): Add enum value > GOVD_MAP_FORCE. > (oacc_default_clause): Fix default for scalars in oacc kernels. > (gimplify_adjust_omp_clauses_1): Handle GOVD_MAP_FORCE. > > * c-c++-common/goacc/kernels-default-2.c: New test. > * c-c++-common/goacc/kernels-default.c: New test. > > --- > gcc/gimplify.c | 19 ++++++++++++++----- > gcc/testsuite/c-c++-common/goacc/kernels-default-2.c | 17 +++++++++++++++++ > gcc/testsuite/c-c++-common/goacc/kernels-default.c | 14 ++++++++++++++ > 3 files changed, 45 insertions(+), 5 deletions(-) > > diff --git a/gcc/gimplify.c b/gcc/gimplify.c > index fcac745..68d90bf 100644 > --- a/gcc/gimplify.c > +++ b/gcc/gimplify.c > @@ -87,6 +87,9 @@ enum gimplify_omp_var_data > /* Flag for GOVD_MAP, if it is always, to or always, tofrom mapping. */ > GOVD_MAP_ALWAYS_TO = 65536, > > + /* Flag for GOVD_MAP, if it is a forced mapping. */ > + GOVD_MAP_FORCE = 131072, > + > GOVD_DATA_SHARE_CLASS = (GOVD_SHARED | GOVD_PRIVATE | GOVD_FIRSTPRIVATE > | GOVD_LASTPRIVATE | GOVD_REDUCTION | GOVD_LINEAR > | GOVD_LOCAL) > @@ -5976,8 +5979,12 @@ oacc_default_clause (struct gimplify_omp_ctx *ctx, tree decl, unsigned flags) > gcc_unreachable (); > > case ORT_ACC_KERNELS: > - /* Everything under kernels are default 'present_or_copy'. */ > + /* Scalars are default 'copy' under kernels, non-scalars are default > + 'present_or_copy'. */ > flags |= GOVD_MAP; > + if (!AGGREGATE_TYPE_P (TREE_TYPE (decl))) > + flags |= GOVD_MAP_FORCE; > + > rkind = "kernels"; > break; > > @@ -7489,10 +7496,12 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n, void *data) > } > else if (code == OMP_CLAUSE_MAP) > { > - OMP_CLAUSE_SET_MAP_KIND (clause, > - flags & GOVD_MAP_TO_ONLY > - ? GOMP_MAP_TO > - : GOMP_MAP_TOFROM); > + int kind = (flags & GOVD_MAP_TO_ONLY > + ? GOMP_MAP_TO > + : GOMP_MAP_TOFROM); > + if (flags & GOVD_MAP_FORCE) > + kind |= GOMP_MAP_FLAG_FORCE; > + OMP_CLAUSE_SET_MAP_KIND (clause, kind); > if (DECL_SIZE (decl) > && TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST) > { > diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-default-2.c b/gcc/testsuite/c-c++-common/goacc/kernels-default-2.c > new file mode 100644 > index 0000000..232b123 > --- /dev/null > +++ b/gcc/testsuite/c-c++-common/goacc/kernels-default-2.c > @@ -0,0 +1,17 @@ > +/* { dg-additional-options "-O2" } */ > +/* { dg-additional-options "-fdump-tree-gimple" } */ > + > +#define N 2 > + > +void > +foo (void) > +{ > + unsigned int a[N]; > + > +#pragma acc kernels > + { > + a[0]++; > + } > +} > + > +/* { dg-final { scan-tree-dump-times "map\\(tofrom" 1 "gimple" } } */ > diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-default.c b/gcc/testsuite/c-c++-common/goacc/kernels-default.c > new file mode 100644 > index 0000000..58cd5e1 > --- /dev/null > +++ b/gcc/testsuite/c-c++-common/goacc/kernels-default.c > @@ -0,0 +1,14 @@ > +/* { dg-additional-options "-O2" } */ > +/* { dg-additional-options "-fdump-tree-gimple" } */ > + > +void > +foo (void) > +{ > + unsigned int i; > +#pragma acc kernels > + { > + i++; > + } > +} > + > +/* { dg-final { scan-tree-dump-times "map\\(force_tofrom" 1 "gimple" } } */ Grüße Thomas
On 02/12/15 19:03, Thomas Schwinge wrote: > Hi! > > Copying Nathan for your information, in case you had not yet seen that > gcc/gimplify.c:oacc_default_clause change: > > On Fri, 27 Nov 2015 12:29:21 +0100, Tom de Vries<Tom_deVries@mentor.com> wrote: >> >The OpenACC 2.0a standard says this about the default mapping for >> >variables used in a kernels region: >> >... >> >An array or variable of aggregate data type 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 >> >present_or_copy clause for the 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. >> >... >> > >> >But atm, all variables including the scalar ones have 'present_or_copy' >> >defaults. >> > >> >This patch makes sure scalar variables get the 'copy' default. >> > >> >Bootstrapped and reg-tested on x86_64. OK for stage3 trunk? > I see: > > PASS: gfortran.dg/goacc/reduction-2.f95 -O scan-tree-dump-times gimple "acc loop private.k. reduction..:a." 1 > PASS: gfortran.dg/goacc/reduction-2.f95 -O scan-tree-dump-times gimple "acc loop private.p. reduction..:a." 1 > [-PASS:-]{+FAIL:+} gfortran.dg/goacc/reduction-2.f95 -O scan-tree-dump-times gimple "target oacc_kernels map.tofrom:a .len: 4.." 1 > PASS: gfortran.dg/goacc/reduction-2.f95 -O scan-tree-dump-times gimple "target oacc_parallel firstprivate.a." 1 > PASS: gfortran.dg/goacc/reduction-2.f95 -O (test for excess errors) > > I guess that one needs to be updated? > Seems to have been fixed by Cesar in r231204. Thanks, - Tom
Fix oacc kernels default mapping for scalars 2015-11-27 Tom de Vries <tom@codesourcery.com> * gimplify.c (enum gimplify_omp_var_data): Add enum value GOVD_MAP_FORCE. (oacc_default_clause): Fix default for scalars in oacc kernels. (gimplify_adjust_omp_clauses_1): Handle GOVD_MAP_FORCE. * c-c++-common/goacc/kernels-default-2.c: New test. * c-c++-common/goacc/kernels-default.c: New test. --- gcc/gimplify.c | 19 ++++++++++++++----- gcc/testsuite/c-c++-common/goacc/kernels-default-2.c | 17 +++++++++++++++++ gcc/testsuite/c-c++-common/goacc/kernels-default.c | 14 ++++++++++++++ 3 files changed, 45 insertions(+), 5 deletions(-) diff --git a/gcc/gimplify.c b/gcc/gimplify.c index fcac745..68d90bf 100644 --- a/gcc/gimplify.c +++ b/gcc/gimplify.c @@ -87,6 +87,9 @@ enum gimplify_omp_var_data /* Flag for GOVD_MAP, if it is always, to or always, tofrom mapping. */ GOVD_MAP_ALWAYS_TO = 65536, + /* Flag for GOVD_MAP, if it is a forced mapping. */ + GOVD_MAP_FORCE = 131072, + GOVD_DATA_SHARE_CLASS = (GOVD_SHARED | GOVD_PRIVATE | GOVD_FIRSTPRIVATE | GOVD_LASTPRIVATE | GOVD_REDUCTION | GOVD_LINEAR | GOVD_LOCAL) @@ -5976,8 +5979,12 @@ oacc_default_clause (struct gimplify_omp_ctx *ctx, tree decl, unsigned flags) gcc_unreachable (); case ORT_ACC_KERNELS: - /* Everything under kernels are default 'present_or_copy'. */ + /* Scalars are default 'copy' under kernels, non-scalars are default + 'present_or_copy'. */ flags |= GOVD_MAP; + if (!AGGREGATE_TYPE_P (TREE_TYPE (decl))) + flags |= GOVD_MAP_FORCE; + rkind = "kernels"; break; @@ -7489,10 +7496,12 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n, void *data) } else if (code == OMP_CLAUSE_MAP) { - OMP_CLAUSE_SET_MAP_KIND (clause, - flags & GOVD_MAP_TO_ONLY - ? GOMP_MAP_TO - : GOMP_MAP_TOFROM); + int kind = (flags & GOVD_MAP_TO_ONLY + ? GOMP_MAP_TO + : GOMP_MAP_TOFROM); + if (flags & GOVD_MAP_FORCE) + kind |= GOMP_MAP_FLAG_FORCE; + OMP_CLAUSE_SET_MAP_KIND (clause, kind); if (DECL_SIZE (decl) && TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST) { diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-default-2.c b/gcc/testsuite/c-c++-common/goacc/kernels-default-2.c new file mode 100644 index 0000000..232b123 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/kernels-default-2.c @@ -0,0 +1,17 @@ +/* { dg-additional-options "-O2" } */ +/* { dg-additional-options "-fdump-tree-gimple" } */ + +#define N 2 + +void +foo (void) +{ + unsigned int a[N]; + +#pragma acc kernels + { + a[0]++; + } +} + +/* { dg-final { scan-tree-dump-times "map\\(tofrom" 1 "gimple" } } */ diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-default.c b/gcc/testsuite/c-c++-common/goacc/kernels-default.c new file mode 100644 index 0000000..58cd5e1 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/kernels-default.c @@ -0,0 +1,14 @@ +/* { dg-additional-options "-O2" } */ +/* { dg-additional-options "-fdump-tree-gimple" } */ + +void +foo (void) +{ + unsigned int i; +#pragma acc kernels + { + i++; + } +} + +/* { dg-final { scan-tree-dump-times "map\\(force_tofrom" 1 "gimple" } } */