Message ID | 56F41D05.7040700@mentor.com |
---|---|
State | New |
Headers | show |
On 24/03/16 17:59, Tom de Vries wrote: > Hi, > > This patch fixes an incorrect warning for the oacc copy clause. > > Consider this test-case: > ... > void > foo (void) > { > int i; > > #pragma acc kernels > { > 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 kernels region. > > The warning occurs because there's an implicit copy(i) clause on the > kernels region, and that copy generates a read of i before the region, > and a write to i in region. > > The patch silences the warning by marking the variable in the copy > 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 > 0001-Remove-incorrect-warning-for-kernels-copy-clause.patch > > > Remove incorrect warning for kernels copy clause > > 2016-03-24 Tom de Vries <tom@codesourcery.com> > > * omp-low.c (lower_omp_target): Set TREE_NO_WARNING for oacc copy > clause. > > * c-c++-common/goacc/uninit-copy-clause.c: New test. > * gfortran.dg/goacc/uninit-copy-clause.f95: New test. > > --- > gcc/omp-low.c | 6 +++- > .../c-c++-common/goacc/uninit-copy-clause.c | 38 ++++++++++++++++++++++ > .../gfortran.dg/goacc/uninit-copy-clause.f95 | 29 +++++++++++++++++ > 3 files changed, 72 insertions(+), 1 deletion(-) > > diff --git a/gcc/omp-low.c b/gcc/omp-low.c > index 3fd6eb3..d107961 100644 > --- a/gcc/omp-low.c > +++ b/gcc/omp-low.c > @@ -16083,7 +16083,11 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) > || map_kind == GOMP_MAP_POINTER > || map_kind == GOMP_MAP_TO_PSET > || map_kind == GOMP_MAP_FORCE_DEVICEPTR) > - gimplify_assign (avar, var, &ilist); > + { > + if (is_gimple_omp_oacc (ctx->stmt)) > + TREE_NO_WARNING (var) = 1; > + gimplify_assign (avar, var, &ilist); > + } > avar = build_fold_addr_expr (avar); > gimplify_assign (x, avar, &ilist); > if ((GOMP_MAP_COPY_FROM_P (map_kind) > diff --git a/gcc/testsuite/c-c++-common/goacc/uninit-copy-clause.c b/gcc/testsuite/c-c++-common/goacc/uninit-copy-clause.c > new file mode 100644 > index 0000000..b3cc445 > --- /dev/null > +++ b/gcc/testsuite/c-c++-common/goacc/uninit-copy-clause.c > @@ -0,0 +1,38 @@ > +/* { dg-do compile } */ > +/* { dg-additional-options "-Wuninitialized" } */ > + > +void > +foo (void) > +{ > + int i; > + > +#pragma acc kernels > + { > + i = 1; > + } > + > +} > + > +void > +foo2 (void) > +{ > + int i; > + > +#pragma acc kernels copy (i) > + { > + i = 1; > + } > + > +} > + > +void > +foo3 (void) > +{ > + int i; > + > +#pragma acc kernels copyin(i) > + { > + i = 1; > + } > + > +} > diff --git a/gcc/testsuite/gfortran.dg/goacc/uninit-copy-clause.f95 b/gcc/testsuite/gfortran.dg/goacc/uninit-copy-clause.f95 > new file mode 100644 > index 0000000..b2aae1d > --- /dev/null > +++ b/gcc/testsuite/gfortran.dg/goacc/uninit-copy-clause.f95 > @@ -0,0 +1,29 @@ > +! { dg-do compile } > +! { dg-additional-options "-Wuninitialized" } > + > +subroutine foo > + integer :: i > + > + !$acc kernels > + i = 1 > + !$acc end kernels > + > +end subroutine foo > + > +subroutine foo2 > + integer :: i > + > + !$acc kernels copy (i) > + i = 1 > + !$acc end kernels > + > +end subroutine foo2 > + > +subroutine foo3 > + integer :: i > + > + !$acc kernels copyin (i) > + i = 1 > + !$acc end kernels > + > +end subroutine foo3 >
Remove incorrect warning for kernels copy clause 2016-03-24 Tom de Vries <tom@codesourcery.com> * omp-low.c (lower_omp_target): Set TREE_NO_WARNING for oacc copy clause. * c-c++-common/goacc/uninit-copy-clause.c: New test. * gfortran.dg/goacc/uninit-copy-clause.f95: New test. --- gcc/omp-low.c | 6 +++- .../c-c++-common/goacc/uninit-copy-clause.c | 38 ++++++++++++++++++++++ .../gfortran.dg/goacc/uninit-copy-clause.f95 | 29 +++++++++++++++++ 3 files changed, 72 insertions(+), 1 deletion(-) diff --git a/gcc/omp-low.c b/gcc/omp-low.c index 3fd6eb3..d107961 100644 --- a/gcc/omp-low.c +++ b/gcc/omp-low.c @@ -16083,7 +16083,11 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) || map_kind == GOMP_MAP_POINTER || map_kind == GOMP_MAP_TO_PSET || map_kind == GOMP_MAP_FORCE_DEVICEPTR) - gimplify_assign (avar, var, &ilist); + { + if (is_gimple_omp_oacc (ctx->stmt)) + TREE_NO_WARNING (var) = 1; + gimplify_assign (avar, var, &ilist); + } avar = build_fold_addr_expr (avar); gimplify_assign (x, avar, &ilist); if ((GOMP_MAP_COPY_FROM_P (map_kind) diff --git a/gcc/testsuite/c-c++-common/goacc/uninit-copy-clause.c b/gcc/testsuite/c-c++-common/goacc/uninit-copy-clause.c new file mode 100644 index 0000000..b3cc445 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/uninit-copy-clause.c @@ -0,0 +1,38 @@ +/* { dg-do compile } */ +/* { dg-additional-options "-Wuninitialized" } */ + +void +foo (void) +{ + int i; + +#pragma acc kernels + { + i = 1; + } + +} + +void +foo2 (void) +{ + int i; + +#pragma acc kernels copy (i) + { + i = 1; + } + +} + +void +foo3 (void) +{ + int i; + +#pragma acc kernels copyin(i) + { + i = 1; + } + +} diff --git a/gcc/testsuite/gfortran.dg/goacc/uninit-copy-clause.f95 b/gcc/testsuite/gfortran.dg/goacc/uninit-copy-clause.f95 new file mode 100644 index 0000000..b2aae1d --- /dev/null +++ b/gcc/testsuite/gfortran.dg/goacc/uninit-copy-clause.f95 @@ -0,0 +1,29 @@ +! { dg-do compile } +! { dg-additional-options "-Wuninitialized" } + +subroutine foo + integer :: i + + !$acc kernels + i = 1 + !$acc end kernels + +end subroutine foo + +subroutine foo2 + integer :: i + + !$acc kernels copy (i) + i = 1 + !$acc end kernels + +end subroutine foo2 + +subroutine foo3 + integer :: i + + !$acc kernels copyin (i) + i = 1 + !$acc end kernels + +end subroutine foo3