Fix oacc kernels default mapping for scalars

Hi,

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?

Thanks,
- Tom

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