diff mbox series

[OpenACC] Remove spurious OpenACC error on combining "auto" with gang/worker/vector

Message ID 996e8fc7-c4c3-6bd8-efdc-84a482975dc7@codesourcery.com
State New
Headers show
Series [OpenACC] Remove spurious OpenACC error on combining "auto" with gang/worker/vector | expand

Commit Message

Gergö Barany Jan. 25, 2019, 2:18 p.m. UTC
On OpenACC loop constructs, it is OK according to the OpenACC spec to 
have both the "auto" clause and one or more of the "gang", "worker", or 
"vector" clauses. GCC emits errors for this combination; this patch 
eliminates that error.

OK for openacc-gcc-8-branch?

Thanks,
Gergö


     gcc/
     * omp-low.c (check_oacc_kernel_gwv): Remove spurious error message.
     * omp-offload.c (oacc_loop_fixed_partitions): Likewise.

     gcc/testsuite/
     * c-c++-common/goacc/combined-directives-3.c: Adjust test.
     * c-c++-common/goacc/loop-2-kernels.c: Likewise.
     * c-c++-common/goacc/loop-2-parallel.c: Likewise.

Comments

Thomas Schwinge Jan. 28, 2019, 8:15 p.m. UTC | #1
Hi Gergő!

On Fri, 25 Jan 2019 15:18:56 +0100, Gergö Barany <gergo@codesourcery.com> wrote:
> On OpenACC loop constructs, it is OK according to the OpenACC spec to 
> have both the "auto" clause and one or more of the "gang", "worker", or 
> "vector" clauses. GCC emits errors for this combination; this patch 
> eliminates that error.

Right.  As I'd mentioned internally, there might be more such things to
be fixed, which we shall re-visit later.  That said, your change is valid
as is, so...

> OK for openacc-gcc-8-branch?

OK.


Grüße
 Thomas


> From d8e7f1826d423de05e11afcb6e422ccaced0f6ea Mon Sep 17 00:00:00 2001
> From: =?UTF-8?q?Gerg=C3=B6=20Barany?= <gergo@codesourcery.com>
> Date: Wed, 23 Jan 2019 03:10:07 -0800
> Subject: [PATCH] Remove spurious OpenACC error on combining "auto" with
>  gang/worker/vector
> 
>     gcc/
>     * omp-low.c (check_oacc_kernel_gwv): Remove spurious error message.
>     * omp-offload.c (oacc_loop_fixed_partitions): Likewise.
> 
>     gcc/testsuite/
>     * c-c++-common/goacc/combined-directives-3.c: Adjust test.
>     * c-c++-common/goacc/loop-2-kernels.c: Likewise.
>     * c-c++-common/goacc/loop-2-parallel.c: Likewise.
> ---
>  gcc/ChangeLog.openacc                                    |  5 +++++
>  gcc/omp-low.c                                            |  3 ---
>  gcc/omp-offload.c                                        |  7 ++-----
>  gcc/testsuite/ChangeLog.openacc                          |  6 ++++++
>  gcc/testsuite/c-c++-common/goacc/combined-directives-3.c |  4 ++--
>  gcc/testsuite/c-c++-common/goacc/loop-2-kernels.c        | 12 ++++++------
>  gcc/testsuite/c-c++-common/goacc/loop-2-parallel.c       | 12 ++++++------
>  7 files changed, 27 insertions(+), 22 deletions(-)
> 
> diff --git a/gcc/ChangeLog.openacc b/gcc/ChangeLog.openacc
> index 932fb37..f3c741c 100644
> --- a/gcc/ChangeLog.openacc
> +++ b/gcc/ChangeLog.openacc
> @@ -1,5 +1,10 @@
>  2019-01-24  Gergö Barany  <gergo@codesourcery.com>
>  
> +	* omp-low.c (check_oacc_kernel_gwv): Remove spurious error message.
> +	* omp-offload.c (oacc_loop_fixed_partitions): Likewise.
> +
> +2019-01-24  Gergö Barany  <gergo@codesourcery.com>
> +
>  	* gimplify.c (oacc_default_clause): Refactor and unify computation of
>  	default mapping clauses.
>  
> diff --git a/gcc/omp-low.c b/gcc/omp-low.c
> index 72b6548..f48002e 100644
> --- a/gcc/omp-low.c
> +++ b/gcc/omp-low.c
> @@ -2397,9 +2397,6 @@ check_oacc_kernel_gwv (gomp_for *stmt, omp_context *ctx)
>        if (has_seq && (this_mask || has_auto))
>  	error_at (gimple_location (stmt), "%<seq%> overrides other"
>  		  " OpenACC loop specifiers");
> -      else if (has_auto && this_mask)
> -	error_at (gimple_location (stmt), "%<auto%> conflicts with other"
> -		  " OpenACC loop specifiers");
>  
>        if (this_mask & outer_mask)
>  	error_at (gimple_location (stmt), "inner loop uses same"
> diff --git a/gcc/omp-offload.c b/gcc/omp-offload.c
> index d428c6f..57a7a06 100644
> --- a/gcc/omp-offload.c
> +++ b/gcc/omp-offload.c
> @@ -1211,14 +1211,11 @@ oacc_loop_fixed_partitions (oacc_loop *loop, unsigned outer_mask)
>        bool maybe_auto
>  	= !seq_par && this_mask == (tiling ? this_mask & -this_mask : 0);
>  
> -      if ((this_mask != 0) + auto_par + seq_par > 1)
> +      if (seq_par && (this_mask != 0 || auto_par))
>  	{
>  	  if (noisy)
>  	    error_at (loop->loc,
> -		      seq_par
> -		      ? G_("%<seq%> overrides other OpenACC loop specifiers")
> -		      : G_("%<auto%> conflicts with other OpenACC loop "
> -			   "specifiers"));
> +		      G_("%<seq%> overrides other OpenACC loop specifiers"));
>  	  maybe_auto = false;
>  	  loop->flags &= ~OLF_AUTO;
>  	  if (seq_par)
> diff --git a/gcc/testsuite/ChangeLog.openacc b/gcc/testsuite/ChangeLog.openacc
> index 3bdce2e..3850d97 100644
> --- a/gcc/testsuite/ChangeLog.openacc
> +++ b/gcc/testsuite/ChangeLog.openacc
> @@ -1,3 +1,9 @@
> +2019-01-24  Gergö Barany  <gergo@codesourcery.com>
> +
> +	* c-c++-common/goacc/combined-directives-3.c: Adjust test.
> +	* c-c++-common/goacc/loop-2-kernels.c: Likewise.
> +	* c-c++-common/goacc/loop-2-parallel.c: Likewise.
> +
>  2019-01-09  Julian Brown  <julian@codesourcery.com>
>  
>  	* c-c++-common/cpp/openacc-define-3.c: Update expected value for
> diff --git a/gcc/testsuite/c-c++-common/goacc/combined-directives-3.c b/gcc/testsuite/c-c++-common/goacc/combined-directives-3.c
> index 77d4182..5aa84dc 100644
> --- a/gcc/testsuite/c-c++-common/goacc/combined-directives-3.c
> +++ b/gcc/testsuite/c-c++-common/goacc/combined-directives-3.c
> @@ -12,9 +12,9 @@ main ()
>      for (y = 0; y < 10; y++)
>        ;
>  
> -#pragma acc parallel loop gang auto /* { dg-error "'auto' conflicts with other OpenACC loop specifiers" } */
> +#pragma acc parallel loop gang seq /* { dg-error "'seq' overrides other OpenACC loop specifiers" } */
>    for (x = 0; x < 10; x++)
> -#pragma acc loop worker auto /* { dg-error "'auto' conflicts with other OpenACC loop specifiers" } */
> +#pragma acc loop worker seq /* { dg-error "'seq' overrides other OpenACC loop specifiers" } */
>      for (y = 0; y < 10; y++)
>  #pragma acc loop vector
>        for (z = 0; z < 10; z++)
> diff --git a/gcc/testsuite/c-c++-common/goacc/loop-2-kernels.c b/gcc/testsuite/c-c++-common/goacc/loop-2-kernels.c
> index 3a11ef5f..2608c12 100644
> --- a/gcc/testsuite/c-c++-common/goacc/loop-2-kernels.c
> +++ b/gcc/testsuite/c-c++-common/goacc/loop-2-kernels.c
> @@ -106,13 +106,13 @@ void K(void)
>  #pragma acc loop seq auto // { dg-error "'seq' overrides" }
>      for (i = 0; i < 10; i++)
>        { }
> -#pragma acc loop gang auto // { dg-error "'auto' conflicts" }
> +#pragma acc loop gang auto
>      for (i = 0; i < 10; i++)
>        { }
> -#pragma acc loop worker auto // { dg-error "'auto' conflicts" }
> +#pragma acc loop worker auto
>      for (i = 0; i < 10; i++)
>        { }
> -#pragma acc loop vector auto // { dg-error "'auto' conflicts" }
> +#pragma acc loop vector auto
>      for (i = 0; i < 10; i++)
>        { }
>    }
> @@ -177,13 +177,13 @@ void K(void)
>  #pragma acc kernels loop seq auto // { dg-error "'seq' overrides" }
>    for (i = 0; i < 10; i++)
>      { }
> -#pragma acc kernels loop gang auto // { dg-error "'auto' conflicts" }
> +#pragma acc kernels loop gang auto
>    for (i = 0; i < 10; i++)
>      { }
> -#pragma acc kernels loop worker auto // { dg-error "'auto' conflicts" }
> +#pragma acc kernels loop worker auto
>    for (i = 0; i < 10; i++)
>      { }
> -#pragma acc kernels loop vector auto // { dg-error "'auto' conflicts" }
> +#pragma acc kernels loop vector auto
>    for (i = 0; i < 10; i++)
>      { }
>  }
> diff --git a/gcc/testsuite/c-c++-common/goacc/loop-2-parallel.c b/gcc/testsuite/c-c++-common/goacc/loop-2-parallel.c
> index 27f7bbd..457a894 100644
> --- a/gcc/testsuite/c-c++-common/goacc/loop-2-parallel.c
> +++ b/gcc/testsuite/c-c++-common/goacc/loop-2-parallel.c
> @@ -90,13 +90,13 @@ void P(void)
>  #pragma acc loop seq auto // { dg-error "'seq' overrides" }
>      for (i = 0; i < 10; i++)
>        { }
> -#pragma acc loop gang auto // { dg-error "'auto' conflicts" }
> +#pragma acc loop gang auto
>      for (i = 0; i < 10; i++)
>        { }
> -#pragma acc loop worker auto // { dg-error "'auto' conflicts" }
> +#pragma acc loop worker auto
>      for (i = 0; i < 10; i++)
>        { }
> -#pragma acc loop vector auto // { dg-error "'auto' conflicts" }
> +#pragma acc loop vector auto
>      for (i = 0; i < 10; i++)
>        { }
>  
> @@ -150,13 +150,13 @@ void P(void)
>  #pragma acc parallel loop seq auto // { dg-error "'seq' overrides" }
>    for (i = 0; i < 10; i++)
>      { }
> -#pragma acc parallel loop gang auto // { dg-error "'auto' conflicts" }
> +#pragma acc parallel loop gang auto
>    for (i = 0; i < 10; i++)
>      { }
> -#pragma acc parallel loop worker auto // { dg-error "'auto' conflicts" }
> +#pragma acc parallel loop worker auto
>    for (i = 0; i < 10; i++)
>      { }
> -#pragma acc parallel loop vector auto // { dg-error "'auto' conflicts" }
> +#pragma acc parallel loop vector auto
>    for (i = 0; i < 10; i++)
>      { }
>  }
diff mbox series

Patch

From d8e7f1826d423de05e11afcb6e422ccaced0f6ea Mon Sep 17 00:00:00 2001
From: =?UTF-8?q?Gerg=C3=B6=20Barany?= <gergo@codesourcery.com>
Date: Wed, 23 Jan 2019 03:10:07 -0800
Subject: [PATCH] Remove spurious OpenACC error on combining "auto" with
 gang/worker/vector

    gcc/
    * omp-low.c (check_oacc_kernel_gwv): Remove spurious error message.
    * omp-offload.c (oacc_loop_fixed_partitions): Likewise.

    gcc/testsuite/
    * c-c++-common/goacc/combined-directives-3.c: Adjust test.
    * c-c++-common/goacc/loop-2-kernels.c: Likewise.
    * c-c++-common/goacc/loop-2-parallel.c: Likewise.
---
 gcc/ChangeLog.openacc                                    |  5 +++++
 gcc/omp-low.c                                            |  3 ---
 gcc/omp-offload.c                                        |  7 ++-----
 gcc/testsuite/ChangeLog.openacc                          |  6 ++++++
 gcc/testsuite/c-c++-common/goacc/combined-directives-3.c |  4 ++--
 gcc/testsuite/c-c++-common/goacc/loop-2-kernels.c        | 12 ++++++------
 gcc/testsuite/c-c++-common/goacc/loop-2-parallel.c       | 12 ++++++------
 7 files changed, 27 insertions(+), 22 deletions(-)

diff --git a/gcc/ChangeLog.openacc b/gcc/ChangeLog.openacc
index 932fb37..f3c741c 100644
--- a/gcc/ChangeLog.openacc
+++ b/gcc/ChangeLog.openacc
@@ -1,5 +1,10 @@ 
 2019-01-24  Gergö Barany  <gergo@codesourcery.com>
 
+	* omp-low.c (check_oacc_kernel_gwv): Remove spurious error message.
+	* omp-offload.c (oacc_loop_fixed_partitions): Likewise.
+
+2019-01-24  Gergö Barany  <gergo@codesourcery.com>
+
 	* gimplify.c (oacc_default_clause): Refactor and unify computation of
 	default mapping clauses.
 
diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index 72b6548..f48002e 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -2397,9 +2397,6 @@  check_oacc_kernel_gwv (gomp_for *stmt, omp_context *ctx)
       if (has_seq && (this_mask || has_auto))
 	error_at (gimple_location (stmt), "%<seq%> overrides other"
 		  " OpenACC loop specifiers");
-      else if (has_auto && this_mask)
-	error_at (gimple_location (stmt), "%<auto%> conflicts with other"
-		  " OpenACC loop specifiers");
 
       if (this_mask & outer_mask)
 	error_at (gimple_location (stmt), "inner loop uses same"
diff --git a/gcc/omp-offload.c b/gcc/omp-offload.c
index d428c6f..57a7a06 100644
--- a/gcc/omp-offload.c
+++ b/gcc/omp-offload.c
@@ -1211,14 +1211,11 @@  oacc_loop_fixed_partitions (oacc_loop *loop, unsigned outer_mask)
       bool maybe_auto
 	= !seq_par && this_mask == (tiling ? this_mask & -this_mask : 0);
 
-      if ((this_mask != 0) + auto_par + seq_par > 1)
+      if (seq_par && (this_mask != 0 || auto_par))
 	{
 	  if (noisy)
 	    error_at (loop->loc,
-		      seq_par
-		      ? G_("%<seq%> overrides other OpenACC loop specifiers")
-		      : G_("%<auto%> conflicts with other OpenACC loop "
-			   "specifiers"));
+		      G_("%<seq%> overrides other OpenACC loop specifiers"));
 	  maybe_auto = false;
 	  loop->flags &= ~OLF_AUTO;
 	  if (seq_par)
diff --git a/gcc/testsuite/ChangeLog.openacc b/gcc/testsuite/ChangeLog.openacc
index 3bdce2e..3850d97 100644
--- a/gcc/testsuite/ChangeLog.openacc
+++ b/gcc/testsuite/ChangeLog.openacc
@@ -1,3 +1,9 @@ 
+2019-01-24  Gergö Barany  <gergo@codesourcery.com>
+
+	* c-c++-common/goacc/combined-directives-3.c: Adjust test.
+	* c-c++-common/goacc/loop-2-kernels.c: Likewise.
+	* c-c++-common/goacc/loop-2-parallel.c: Likewise.
+
 2019-01-09  Julian Brown  <julian@codesourcery.com>
 
 	* c-c++-common/cpp/openacc-define-3.c: Update expected value for
diff --git a/gcc/testsuite/c-c++-common/goacc/combined-directives-3.c b/gcc/testsuite/c-c++-common/goacc/combined-directives-3.c
index 77d4182..5aa84dc 100644
--- a/gcc/testsuite/c-c++-common/goacc/combined-directives-3.c
+++ b/gcc/testsuite/c-c++-common/goacc/combined-directives-3.c
@@ -12,9 +12,9 @@  main ()
     for (y = 0; y < 10; y++)
       ;
 
-#pragma acc parallel loop gang auto /* { dg-error "'auto' conflicts with other OpenACC loop specifiers" } */
+#pragma acc parallel loop gang seq /* { dg-error "'seq' overrides other OpenACC loop specifiers" } */
   for (x = 0; x < 10; x++)
-#pragma acc loop worker auto /* { dg-error "'auto' conflicts with other OpenACC loop specifiers" } */
+#pragma acc loop worker seq /* { dg-error "'seq' overrides other OpenACC loop specifiers" } */
     for (y = 0; y < 10; y++)
 #pragma acc loop vector
       for (z = 0; z < 10; z++)
diff --git a/gcc/testsuite/c-c++-common/goacc/loop-2-kernels.c b/gcc/testsuite/c-c++-common/goacc/loop-2-kernels.c
index 3a11ef5f..2608c12 100644
--- a/gcc/testsuite/c-c++-common/goacc/loop-2-kernels.c
+++ b/gcc/testsuite/c-c++-common/goacc/loop-2-kernels.c
@@ -106,13 +106,13 @@  void K(void)
 #pragma acc loop seq auto // { dg-error "'seq' overrides" }
     for (i = 0; i < 10; i++)
       { }
-#pragma acc loop gang auto // { dg-error "'auto' conflicts" }
+#pragma acc loop gang auto
     for (i = 0; i < 10; i++)
       { }
-#pragma acc loop worker auto // { dg-error "'auto' conflicts" }
+#pragma acc loop worker auto
     for (i = 0; i < 10; i++)
       { }
-#pragma acc loop vector auto // { dg-error "'auto' conflicts" }
+#pragma acc loop vector auto
     for (i = 0; i < 10; i++)
       { }
   }
@@ -177,13 +177,13 @@  void K(void)
 #pragma acc kernels loop seq auto // { dg-error "'seq' overrides" }
   for (i = 0; i < 10; i++)
     { }
-#pragma acc kernels loop gang auto // { dg-error "'auto' conflicts" }
+#pragma acc kernels loop gang auto
   for (i = 0; i < 10; i++)
     { }
-#pragma acc kernels loop worker auto // { dg-error "'auto' conflicts" }
+#pragma acc kernels loop worker auto
   for (i = 0; i < 10; i++)
     { }
-#pragma acc kernels loop vector auto // { dg-error "'auto' conflicts" }
+#pragma acc kernels loop vector auto
   for (i = 0; i < 10; i++)
     { }
 }
diff --git a/gcc/testsuite/c-c++-common/goacc/loop-2-parallel.c b/gcc/testsuite/c-c++-common/goacc/loop-2-parallel.c
index 27f7bbd..457a894 100644
--- a/gcc/testsuite/c-c++-common/goacc/loop-2-parallel.c
+++ b/gcc/testsuite/c-c++-common/goacc/loop-2-parallel.c
@@ -90,13 +90,13 @@  void P(void)
 #pragma acc loop seq auto // { dg-error "'seq' overrides" }
     for (i = 0; i < 10; i++)
       { }
-#pragma acc loop gang auto // { dg-error "'auto' conflicts" }
+#pragma acc loop gang auto
     for (i = 0; i < 10; i++)
       { }
-#pragma acc loop worker auto // { dg-error "'auto' conflicts" }
+#pragma acc loop worker auto
     for (i = 0; i < 10; i++)
       { }
-#pragma acc loop vector auto // { dg-error "'auto' conflicts" }
+#pragma acc loop vector auto
     for (i = 0; i < 10; i++)
       { }
 
@@ -150,13 +150,13 @@  void P(void)
 #pragma acc parallel loop seq auto // { dg-error "'seq' overrides" }
   for (i = 0; i < 10; i++)
     { }
-#pragma acc parallel loop gang auto // { dg-error "'auto' conflicts" }
+#pragma acc parallel loop gang auto
   for (i = 0; i < 10; i++)
     { }
-#pragma acc parallel loop worker auto // { dg-error "'auto' conflicts" }
+#pragma acc parallel loop worker auto
   for (i = 0; i < 10; i++)
     { }
-#pragma acc parallel loop vector auto // { dg-error "'auto' conflicts" }
+#pragma acc parallel loop vector auto
   for (i = 0; i < 10; i++)
     { }
 }
-- 
2.8.1