This is the mail archive of the gcc-patches@gcc.gnu.org mailing list for the GCC project.


Index Nav: [Date Index] [Subject Index] [Author Index] [Thread Index]
Message Nav: [Date Prev] [Date Next] [Thread Prev] [Thread Next]
Other format: [Raw text]

[PING][PATCH] Remove incorrect warning for kernels copy clause


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



Index Nav: [Date Index] [Subject Index] [Author Index] [Thread Index]
Message Nav: [Date Prev] [Date Next] [Thread Prev] [Thread Next]