Remove incorrect warning for parallel firstprivate clause

Hi,

This patch fixes an incorrect warning for the oacc firstprivate clause.

Consider this test-case:
...
void
foo (void)
{
   int i;

#pragma acc parallel
   {
     i = 1;
   }
}
...


When compiling with -fopenacc -Wuninitialized, we get an 'is used 
uninitialized' warning for variable 'i', which is confusing given that 
'i' is not used, but only set in the parallel region.

The warning occurs because there's an implicit firstprivate(i) clause on 
the parallel region, and that firstprivate clause generates a read of i 
before the region, and a write to i in the region.

The patch silences the warning by marking the variable in the 
firstprivate clause with TREE_NO_WARNING.

Build and reg-tested with goacc.exp, gomp.exp and target-libgomp.

OK for trunk if bootstrap and reg-test succeeds?

Thanks,
- Tom

On 24/03/16 18:02, Tom de Vries wrote:
> Hi,
>
> This patch fixes an incorrect warning for the oacc firstprivate clause.
>
> Consider this test-case:
> ...
> void
> foo (void)
> {
>    int i;
>
> #pragma acc parallel
>    {
>      i = 1;
>    }
> }
> ...
>
>
> When compiling with -fopenacc -Wuninitialized, we get an 'is used
> uninitialized' warning for variable 'i', which is confusing given that
> 'i' is not used, but only set in the parallel region.
>
> The warning occurs because there's an implicit firstprivate(i) clause on
> the parallel region, and that firstprivate clause generates a read of i
> before the region, and a write to i in the region.
>
> The patch silences the warning by marking the variable in the
> firstprivate clause with TREE_NO_WARNING.
>
> Build and reg-tested with goacc.exp, gomp.exp and target-libgomp.
>
> OK for trunk if bootstrap and reg-test succeeds?

Ping.

Thanks,
- Tom

> 0002-Remove-incorrect-warning-for-parallel-firstprivate-clause.patch
>
>
> Remove incorrect warning for parallel firstprivate clause
>
> 2016-03-24  Tom de Vries  <tom@codesourcery.com>
>
> 	* omp-low.c (lower_omp_target): Set TREE_NO_WARNING for oacc
> 	firstprivate clause.
>
> 	* c-c++-common/goacc/uninit-firstprivate-clause.c: New test.
> 	* gfortran.dg/goacc/uninit-firstprivate-clause.f95: New test.
>
> ---
>   gcc/omp-low.c                                      |  5 ++++-
>   .../goacc/uninit-firstprivate-clause.c             | 25 ++++++++++++++++++++++
>   .../goacc/uninit-firstprivate-clause.f95           | 18 ++++++++++++++++
>   3 files changed, 47 insertions(+), 1 deletion(-)
>
> diff --git a/gcc/omp-low.c b/gcc/omp-low.c
> index d107961..41eb3c8 100644
> --- a/gcc/omp-low.c
> +++ b/gcc/omp-low.c
> @@ -16068,7 +16068,10 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
>   		  {
>   		    gcc_assert (is_gimple_omp_oacc (ctx->stmt));
>   		    if (!is_reference (var))
> -		      var = build_fold_addr_expr (var);
> +		      {
> +			TREE_NO_WARNING (var) = 1;
> +			var = build_fold_addr_expr (var);
> +		      }
>   		    else
>   		      talign = TYPE_ALIGN_UNIT (TREE_TYPE (TREE_TYPE (ovar)));
>   		    gimplify_assign (x, var, &ilist);
> diff --git a/gcc/testsuite/c-c++-common/goacc/uninit-firstprivate-clause.c b/gcc/testsuite/c-c++-common/goacc/uninit-firstprivate-clause.c
> new file mode 100644
> index 0000000..3d3a03e
> --- /dev/null
> +++ b/gcc/testsuite/c-c++-common/goacc/uninit-firstprivate-clause.c
> @@ -0,0 +1,25 @@
> +/* { dg-do compile } */
> +/* { dg-additional-options "-Wuninitialized" } */
> +
> +void
> +foo (void)
> +{
> +  int i;
> +
> +#pragma acc parallel
> +  {
> +    i = 1;
> +  }
> +}
> +
> +
> +void
> +foo2 (void)
> +{
> +  int i;
> +
> +#pragma acc parallel firstprivate (i)
> +  {
> +    i = 1;
> +  }
> +}
> diff --git a/gcc/testsuite/gfortran.dg/goacc/uninit-firstprivate-clause.f95 b/gcc/testsuite/gfortran.dg/goacc/uninit-firstprivate-clause.f95
> new file mode 100644
> index 0000000..c18765b
> --- /dev/null
> +++ b/gcc/testsuite/gfortran.dg/goacc/uninit-firstprivate-clause.f95
> @@ -0,0 +1,18 @@
> +! { dg-do compile }
> +! { dg-additional-options "-Wuninitialized" }
> +
> +subroutine test
> +  INTEGER :: i
> +
> +  !$acc parallel
> +  i = 1
> +  !$acc end parallel
> +end subroutine test
> +
> +subroutine test2
> +  INTEGER :: i
> +
> +  !$acc parallel firstprivate (i)
> +  i = 1
> +  !$acc end parallel
> +end subroutine test2
>

On Tue, Apr 05, 2016 at 12:17:16PM +0200, Tom de Vries wrote:
> On 24/03/16 18:02, Tom de Vries wrote:
> >Remove incorrect warning for parallel firstprivate clause
> >
> >2016-03-24  Tom de Vries  <tom@codesourcery.com>
> >
> >	* omp-low.c (lower_omp_target): Set TREE_NO_WARNING for oacc
> >	firstprivate clause.
> >
> >	* c-c++-common/goacc/uninit-firstprivate-clause.c: New test.
> >	* gfortran.dg/goacc/uninit-firstprivate-clause.f95: New test.
> >
> >---
> >  gcc/omp-low.c                                      |  5 ++++-
> >  .../goacc/uninit-firstprivate-clause.c             | 25 ++++++++++++++++++++++
> >  .../goacc/uninit-firstprivate-clause.f95           | 18 ++++++++++++++++
> >  3 files changed, 47 insertions(+), 1 deletion(-)
> >
> >diff --git a/gcc/omp-low.c b/gcc/omp-low.c
> >index d107961..41eb3c8 100644
> >--- a/gcc/omp-low.c
> >+++ b/gcc/omp-low.c
> >@@ -16068,7 +16068,10 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
> >  		  {
> >  		    gcc_assert (is_gimple_omp_oacc (ctx->stmt));
> >  		    if (!is_reference (var))
> >-		      var = build_fold_addr_expr (var);
> >+		      {
> >+			TREE_NO_WARNING (var) = 1;
> >+			var = build_fold_addr_expr (var);
> >+		      }

IMHO it should be done only if var is is_gimple_reg (var), otherwise all
that happens on the caller side is that you take the address of the actual
variable.  Also, I think it would be better to do this only
for implicit firstprivate (and map) clauses, if somebody uses explicit
firstprivate on a var, I think it is better to warn if the var is
uninitialized, the user can then use private clause instead.

BTW, some undesirable warnings are also on OpenMP code (I'm adding
TREE_NO_WARNING already in case of shared clause), I've filed PR70550
to track this and will attach there a patch soon.

	Jakub