diff mbox

Remove incorrect warning for kernels copy clause

Message ID 56F41D05.7040700@mentor.com
State New
Headers show

Commit Message

Tom de Vries March 24, 2016, 4:59 p.m. UTC
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?

Thanks,
- Tom

Comments

Tom de Vries April 5, 2016, 10:16 a.m. UTC | #1
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
>
diff mbox

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