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 |
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++) > { } > }
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