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_devr...@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 <t...@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
signature.asc
Description: PGP signature
