public inbox for gcc-patches@gcc.gnu.org
 help / color / mirror / Atom feed
* [PATCH] Remove incorrect warning for kernels copy clause
@ 2016-03-24 17:03 Tom de Vries
  2016-04-05 10:16 ` [PING][PATCH] " Tom de Vries
  0 siblings, 1 reply; 3+ messages in thread
From: Tom de Vries @ 2016-03-24 17:03 UTC (permalink / raw)
  To: Jakub Jelinek; +Cc: GCC Patches

[-- Attachment #1: Type: text/plain, Size: 771 bytes --]

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

[-- Attachment #2: 0001-Remove-incorrect-warning-for-kernels-copy-clause.patch --]
[-- Type: text/x-patch, Size: 2462 bytes --]

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

^ permalink raw reply	[flat|nested] 3+ messages in thread

* [PING][PATCH] Remove incorrect warning for kernels copy clause
  2016-03-24 17:03 [PATCH] Remove incorrect warning for kernels copy clause Tom de Vries
@ 2016-04-05 10:16 ` Tom de Vries
  2016-04-08  8:54   ` Tom de Vries
  0 siblings, 1 reply; 3+ messages in thread
From: Tom de Vries @ 2016-04-05 10:16 UTC (permalink / raw)
  To: Jakub Jelinek; +Cc: GCC Patches

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
>

^ permalink raw reply	[flat|nested] 3+ messages in thread

* Re: [PING][PATCH] Remove incorrect warning for kernels copy clause
  2016-04-05 10:16 ` [PING][PATCH] " Tom de Vries
@ 2016-04-08  8:54   ` Tom de Vries
  0 siblings, 0 replies; 3+ messages in thread
From: Tom de Vries @ 2016-04-08  8:54 UTC (permalink / raw)
  To: Jakub Jelinek; +Cc: GCC Patches

[-- Attachment #1: Type: text/plain, Size: 1068 bytes --]

On 05/04/16 12:16, Tom de Vries wrote:
> 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.
>

Committed as attached, with only the testcases. The code change has been 
committed as part of the fix for PR70550.

Thanks,
- TOm

[-- Attachment #2: 0001-Add-goacc-uninit-copy-clause.-c-f95-testcases.patch --]
[-- Type: text/x-patch, Size: 1675 bytes --]

Add goacc/uninit-copy-clause.{c,f95} testcases

2016-03-24  Tom de Vries  <tom@codesourcery.com>

	* c-c++-common/goacc/uninit-copy-clause.c: New test.
	* gfortran.dg/goacc/uninit-copy-clause.f95: New test.

---
 .../c-c++-common/goacc/uninit-copy-clause.c        | 38 ++++++++++++++++++++++
 .../gfortran.dg/goacc/uninit-copy-clause.f95       | 29 +++++++++++++++++
 2 files changed, 67 insertions(+)

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

^ permalink raw reply	[flat|nested] 3+ messages in thread

end of thread, other threads:[~2016-04-08  8:54 UTC | newest]

Thread overview: 3+ messages (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2016-03-24 17:03 [PATCH] Remove incorrect warning for kernels copy clause Tom de Vries
2016-04-05 10:16 ` [PING][PATCH] " Tom de Vries
2016-04-08  8:54   ` Tom de Vries

This is a public inbox, see mirroring instructions
for how to clone and mirror all data and code used for this inbox;
as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).