* [PATCH] openacc: Warnings for strange partitioning choices for parallel regions
@ 2021-02-26 12:21 Julian Brown
2021-04-26 10:39 ` Thomas Schwinge
0 siblings, 1 reply; 7+ messages in thread
From: Julian Brown @ 2021-02-26 12:21 UTC (permalink / raw)
To: gcc-patches; +Cc: Jakub Jelinek, Thomas Schwinge
This patch adds warnings for strange partitioning choices -- specifying
either more or fewer partitioning levels on a parallel directive than the
enclosed offload region actually uses.
We've used a version of this patch on the og8/og9/og10 branches for
quite a while. Versions have been posted for mainline submission as
part of a larger patch several times, e.g. here:
https://gcc.gnu.org/pipermail/gcc-patches/2018-October/507938.html
One motivation for committing this patch is it removes an ongoing
divergence in a large number of test cases between the og* branches and
mainline, namely whether the added warnings are expected as a result of
compiling those test cases, or not.
The warnings themselves are perhaps slightly aggressive, but are intended
to help the user write more efficient code.
Tested with offloading to AMD GCN and (separately) to NVPTX.
OK (now or for stage 1)?
Julian
2021-02-22 Nathan Sidwell <nathan@acm.org>
Tom de Vries <tdevries@suse.de>
Thomas Schwinge <thomas@codesourcery.com>
Julian Brown <julian@codesourcery.com>
Kwok Cheung Yeung <kcy@codesourcnNery.com>
gcc/
(oacc_validate_dims): Emit warnings about strange partitioning
choices.
gcc/testsuite/
* c-c++-common/goacc/acc-icf.c: Update.
* c-c++-common/classify-serial.c: Likewise.
* c-c++-common/kernels-decompose-2.c: Likewise.
* c-c++-common/goacc/parallel-dims-1.c: Likewise.
* c-c++-common/goacc/parallel-reduction.c: Likewise.
* c-c++-common/goacc/pr70688.c: Likewise.
* c-c++-common/goacc/routine-1.c: Likewise.
* c-c++-common/goacc/routine-level-of-parallelism-2.c: Likewise.
* c-c++-common/goacc/uninit-dim-clause.c: Likewise.
* gfortran.dg/goacc/kernels-decompose-2.f95: Likewise.
* gfortran.dg/goacc/parallel-tree.f95: Likewise.
* gfortran.dg/goacc/routine-4.f90: Likewise.
* gfortran.dg/goacc/routine-level-of-parallelism-1.f90: Likewise.
* gfortran.dg/goacc/routine-module-mod-1.f90: Likewise.
* gfortran.dg/goacc/routine-multiple-directives-1.f90: Likewise.
* gfortran.dg/goacc/uninit-dim-clause.f95: Likewise.
libgomp/
* testsuite/libgomp.oacc-c-c++-common/loop-g-1.c: Add -w.
* testsuite/libgomp.oacc-c-c++-common/loop-g-2.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/loop-red-g-1.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/loop-red-w-1.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/loop-red-w-2.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/loop-warn-1.c: New.
* testsuite/libgomp.oacc-c-c++-common/firstprivate-1.c: Update.
* testsuite/libgomp.oacc-c-c++-common/loop-auto-1.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/loop-w-1.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/mode-transitions.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/parallel-dims.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/pr85381-3.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/private-variables.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/reduction-7.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/routine-g-1.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/routine-w-1.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/routine-wv-2.c: Likewise.
* testsuite/libgomp.oacc-fortran/optional-private.f90: Likewise.
* testsuite/libgomp.oacc-fortran/par-reduction-2-1.f: Likewise.
* testsuite/libgomp.oacc-fortran/par-reduction-2-2.f: Likewise.
* testsuite/libgomp.oacc-fortran/parallel-dims.f90: Likewise.
* testsuite/libgomp.oacc-fortran/pr84028.f90: Likewise.
* testsuite/libgomp.oacc-fortran/private-variables.f90: Likewise.
* testsuite/libgomp.oacc-fortran/routine-7.f90: Likewise.
---
gcc/omp-offload.c | 28 ++++++++++++++
gcc/testsuite/c-c++-common/goacc/acc-icf.c | 4 +-
.../c-c++-common/goacc/classify-serial.c | 2 +
.../c-c++-common/goacc/kernels-decompose-2.c | 1 +
.../c-c++-common/goacc/parallel-dims-1.c | 6 ++-
.../c-c++-common/goacc/parallel-reduction.c | 2 +-
gcc/testsuite/c-c++-common/goacc/pr70688.c | 6 +--
gcc/testsuite/c-c++-common/goacc/routine-1.c | 6 +--
.../goacc/routine-level-of-parallelism-2.c | 3 ++
.../c-c++-common/goacc/uninit-dim-clause.c | 3 ++
.../gfortran.dg/goacc/classify-serial.f95 | 2 +
.../gfortran.dg/goacc/kernels-decompose-2.f95 | 1 +
.../gfortran.dg/goacc/parallel-tree.f95 | 3 ++
gcc/testsuite/gfortran.dg/goacc/routine-4.f90 | 4 ++
.../goacc/routine-level-of-parallelism-1.f90 | 4 +-
.../goacc/routine-module-mod-1.f90 | 1 +
.../goacc/routine-multiple-directives-1.f90 | 4 +-
.../gfortran.dg/goacc/uninit-dim-clause.f95 | 3 ++
.../firstprivate-1.c | 2 +
.../libgomp.oacc-c-c++-common/loop-auto-1.c | 2 +-
.../libgomp.oacc-c-c++-common/loop-g-1.c | 1 +
.../libgomp.oacc-c-c++-common/loop-g-2.c | 1 +
.../libgomp.oacc-c-c++-common/loop-red-g-1.c | 1 +
.../libgomp.oacc-c-c++-common/loop-red-w-1.c | 1 +
.../libgomp.oacc-c-c++-common/loop-red-w-2.c | 1 +
.../libgomp.oacc-c-c++-common/loop-w-1.c | 1 +
.../libgomp.oacc-c-c++-common/loop-warn-1.c | 37 +++++++++++++++++++
.../mode-transitions.c | 21 +++++++++++
.../libgomp.oacc-c-c++-common/parallel-dims.c | 12 ++++--
.../libgomp.oacc-c-c++-common/pr85381-3.c | 2 +-
.../private-variables.c | 13 +++++++
.../libgomp.oacc-c-c++-common/reduction-7.c | 8 ++++
.../libgomp.oacc-c-c++-common/routine-g-1.c | 2 +
.../libgomp.oacc-c-c++-common/routine-w-1.c | 1 +
.../libgomp.oacc-c-c++-common/routine-wv-2.c | 3 +-
.../libgomp.oacc-fortran/optional-private.f90 | 3 ++
.../libgomp.oacc-fortran/par-reduction-2-1.f | 4 +-
.../libgomp.oacc-fortran/par-reduction-2-2.f | 4 +-
.../libgomp.oacc-fortran/parallel-dims.f90 | 3 ++
.../libgomp.oacc-fortran/pr84028.f90 | 2 +-
.../private-variables.f90 | 8 ++++
.../libgomp.oacc-fortran/routine-7.f90 | 2 +-
42 files changed, 191 insertions(+), 27 deletions(-)
create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/loop-warn-1.c
diff --git a/gcc/omp-offload.c b/gcc/omp-offload.c
index ba0937fba94..57be342da97 100644
--- a/gcc/omp-offload.c
+++ b/gcc/omp-offload.c
@@ -928,6 +928,34 @@ oacc_validate_dims (tree fn, tree attrs, int *dims, int level, unsigned used)
pos = TREE_CHAIN (pos);
}
+ bool check = true;
+#ifdef ACCEL_COMPILER
+ check = false;
+#endif
+ if (check
+ && !lookup_attribute ("oacc kernels", DECL_ATTRIBUTES (fn)))
+ {
+ static char const *const axes[] =
+ /* Must be kept in sync with GOMP_DIM enumeration. */
+ { "gang", "worker", "vector" };
+ for (ix = level >= 0 ? level : 0; ix != GOMP_DIM_MAX; ix++)
+ if (dims[ix] < 0)
+ ; /* Defaulting axis. */
+ else if ((used & GOMP_DIM_MASK (ix)) && dims[ix] == 1)
+ /* There is partitioned execution, but the user requested a
+ dimension size of 1. They're probably confused. */
+ warning_at (DECL_SOURCE_LOCATION (fn), 0,
+ "region contains %s partitioned code but"
+ " is not %s partitioned", axes[ix], axes[ix]);
+ else if (!(used & GOMP_DIM_MASK (ix)) && dims[ix] != 1)
+ /* The dimension is explicitly partitioned to non-unity, but
+ no use is made within the region. */
+ warning_at (DECL_SOURCE_LOCATION (fn), 0,
+ "region is %s partitioned but"
+ " does not contain %s partitioned code",
+ axes[ix], axes[ix]);
+ }
+
bool changed = targetm.goacc.validate_dims (fn, dims, level, used);
/* Default anything left to 1 or a partitioned default. */
diff --git a/gcc/testsuite/c-c++-common/goacc/acc-icf.c b/gcc/testsuite/c-c++-common/goacc/acc-icf.c
index 98b536c34d0..8601ace5db1 100644
--- a/gcc/testsuite/c-c++-common/goacc/acc-icf.c
+++ b/gcc/testsuite/c-c++-common/goacc/acc-icf.c
@@ -4,7 +4,7 @@
#pragma acc routine gang
int
-routine1 (int n)
+routine1 (int n) /* { dg-bogus "region is worker partitioned but does not contain worker partitioned code" "" { xfail *-*-* } } */
{
int i;
@@ -17,7 +17,7 @@ routine1 (int n)
#pragma acc routine gang
int
-routine2 (int n)
+routine2 (int n) /* { dg-bogus "region is worker partitioned but does not contain worker partitioned code" "" { xfail *-*-* } } */
{
int i;
diff --git a/gcc/testsuite/c-c++-common/goacc/classify-serial.c b/gcc/testsuite/c-c++-common/goacc/classify-serial.c
index 94ace1b3c20..89b5f5ed9f9 100644
--- a/gcc/testsuite/c-c++-common/goacc/classify-serial.c
+++ b/gcc/testsuite/c-c++-common/goacc/classify-serial.c
@@ -15,6 +15,8 @@ extern unsigned int *__restrict c;
void SERIAL ()
{
#pragma acc serial loop copyin (a[0:N], b[0:N]) copyout (c[0:N]) /* { dg-message "optimized: assigned OpenACC gang vector loop parallelism" } */
+/* { dg-warning "region contains gang partitioned code but is not gang partitioned" "" { target *-*-* } .-1 } */
+/* { dg-warning "region contains vector partitioned code but is not vector partitioned" "" { target *-*-* } .-2 } */
for (unsigned int i = 0; i < N; i++)
c[i] = a[i] + b[i];
}
diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-2.c b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-2.c
index ec0f75c4a61..f9080936f25 100644
--- a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-2.c
+++ b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-2.c
@@ -121,6 +121,7 @@ main ()
/*TODO If such a construct is placed in its own part (like it is, here), can't this actually use gang paralelism, instead of "gang-single"?
{ dg-message "note: beginning 'gang-single' part in OpenACC 'kernels' region" "" { target *-*-* } l_part$c_part } */
/* { dg-optimized "assigned OpenACC gang worker vector loop parallelism" "" { target *-*-* } l_part$c_part } */
+ /* { dg-warning "region contains gang partitioned code but is not gang partitioned" "" { target *-*-* } .-6 } */
#pragma acc loop independent /* { dg-line l_loop_j[incr c_loop_j] } */
/* { dg-message "note: parallelized loop nest in OpenACC 'kernels' region" "" { target *-*-* } l_loop_j$c_loop_j } */
diff --git a/gcc/testsuite/c-c++-common/goacc/parallel-dims-1.c b/gcc/testsuite/c-c++-common/goacc/parallel-dims-1.c
index 57f682f7a0d..6cdbebe0128 100644
--- a/gcc/testsuite/c-c++-common/goacc/parallel-dims-1.c
+++ b/gcc/testsuite/c-c++-common/goacc/parallel-dims-1.c
@@ -3,9 +3,11 @@
void f(int i)
{
-#pragma acc kernels num_gangs(i) num_workers(i) vector_length(i)
+#pragma acc kernels \
+ num_gangs(i) num_workers(i) vector_length(i)
;
-#pragma acc parallel num_gangs(i) num_workers(i) vector_length(i)
+#pragma acc parallel /* { dg-bogus "region is (gang|worker|vector) partitioned" "" { xfail *-*-* } } */ \
+ num_gangs(i) num_workers(i) vector_length(i)
;
}
diff --git a/gcc/testsuite/c-c++-common/goacc/parallel-reduction.c b/gcc/testsuite/c-c++-common/goacc/parallel-reduction.c
index d7cc9470127..9a142c4e02d 100644
--- a/gcc/testsuite/c-c++-common/goacc/parallel-reduction.c
+++ b/gcc/testsuite/c-c++-common/goacc/parallel-reduction.c
@@ -6,7 +6,7 @@ main ()
#pragma acc data copy (dummy)
{
-#pragma acc parallel num_gangs (10) copy (sum) reduction (+:sum)
+#pragma acc parallel num_gangs (10) copy (sum) reduction (+:sum) /* { dg-warning "gang partitioned" } */
{
int v = 5;
sum += 10 + v;
diff --git a/gcc/testsuite/c-c++-common/goacc/pr70688.c b/gcc/testsuite/c-c++-common/goacc/pr70688.c
index 5a236654069..3f5584a503e 100644
--- a/gcc/testsuite/c-c++-common/goacc/pr70688.c
+++ b/gcc/testsuite/c-c++-common/goacc/pr70688.c
@@ -21,7 +21,7 @@ parallel_reduction ()
#pragma acc data copy (dummy)
{
-#pragma acc parallel num_gangs (10) copy (sum) reduction (+:sum)
+#pragma acc parallel num_gangs (10) copy (sum) reduction (+:sum) /* { dg-warning "region is gang partitioned" } */
{
int v = 5;
sum += 10 + v;
@@ -36,11 +36,11 @@ main ()
{
int i, s = 0;
-#pragma acc parallel num_gangs (10) copy (s) reduction (+:s)
+#pragma acc parallel num_gangs (10) copy (s) reduction (+:s) /* { dg-warning "region is gang partitioned" } */
for (i = 0; i < n; i++)
s += i+1;
-#pragma acc parallel num_gangs (10) reduction (+:s) copy (s)
+#pragma acc parallel num_gangs (10) reduction (+:s) copy (s) /* { dg-warning "region is gang partitioned" } */
for (i = 0; i < n; i++)
s += i+1;
diff --git a/gcc/testsuite/c-c++-common/goacc/routine-1.c b/gcc/testsuite/c-c++-common/goacc/routine-1.c
index a75692246b6..b90e2c11c9d 100644
--- a/gcc/testsuite/c-c++-common/goacc/routine-1.c
+++ b/gcc/testsuite/c-c++-common/goacc/routine-1.c
@@ -1,16 +1,16 @@
#pragma acc routine gang
-void gang (void)
+void gang (void) /* { dg-warning "partitioned" 3 } */
{
}
#pragma acc routine worker
-void worker (void)
+void worker (void) /* { dg-warning "partitioned" 2 } */
{
}
#pragma acc routine vector
-void vector (void)
+void vector (void) /* { dg-warning "partitioned" 1 } */
{
}
diff --git a/gcc/testsuite/c-c++-common/goacc/routine-level-of-parallelism-2.c b/gcc/testsuite/c-c++-common/goacc/routine-level-of-parallelism-2.c
index a066f2b9c2b..5f3cc20c7dc 100644
--- a/gcc/testsuite/c-c++-common/goacc/routine-level-of-parallelism-2.c
+++ b/gcc/testsuite/c-c++-common/goacc/routine-level-of-parallelism-2.c
@@ -4,6 +4,9 @@
#pragma acc routine gang
void g_1 (void)
+/* { dg-warning "does not contain gang partitioned code" "" { target *-*-* } .-1 } */
+/* { dg-warning "does not contain worker partitioned code" "" { target *-*-* } .-2 } */
+/* { dg-warning "does not contain vector partitioned code" "" { target *-*-* } .-3 } */
{
}
#pragma acc routine (g_1) gang
diff --git a/gcc/testsuite/c-c++-common/goacc/uninit-dim-clause.c b/gcc/testsuite/c-c++-common/goacc/uninit-dim-clause.c
index 827dac71519..9429e5bfa01 100644
--- a/gcc/testsuite/c-c++-common/goacc/uninit-dim-clause.c
+++ b/gcc/testsuite/c-c++-common/goacc/uninit-dim-clause.c
@@ -5,12 +5,15 @@ void acc_parallel()
int i, j, k;
#pragma acc parallel num_gangs(i) /* { dg-warning "is used uninitialized" } */
+ /* { dg-warning "does not contain gang partitioned code" "" { target *-*-* } .-1 } */
;
#pragma acc parallel num_workers(j) /* { dg-warning "is used uninitialized" } */
+ /* { dg-warning "does not contain worker partitioned code" "" { target *-*-* } .-1 } */
;
#pragma acc parallel vector_length(k) /* { dg-warning "is used uninitialized" } */
+ /* { dg-warning "does not contain vector partitioned code" "" { target *-*-* } .-1 } */
;
}
diff --git a/gcc/testsuite/gfortran.dg/goacc/classify-serial.f95 b/gcc/testsuite/gfortran.dg/goacc/classify-serial.f95
index 51061afd2c6..d7052bacbe8 100644
--- a/gcc/testsuite/gfortran.dg/goacc/classify-serial.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/classify-serial.f95
@@ -15,6 +15,8 @@ program main
call setup(a, b)
!$acc serial loop copyin (a(0:n-1), b(0:n-1)) copyout (c(0:n-1)) ! { dg-message "optimized: assigned OpenACC gang vector loop parallelism" }
+ ! { dg-warning "not gang partitioned" "" { target *-*-* } .-1 }
+ ! { dg-warning "not vector partitioned" "" { target *-*-* } .-2 }
do i = 0, n - 1
c(i) = a(i) + b(i)
end do
diff --git a/gcc/testsuite/gfortran.dg/goacc/kernels-decompose-2.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-decompose-2.f95
index 22f65e5c694..09bc6626030 100644
--- a/gcc/testsuite/gfortran.dg/goacc/kernels-decompose-2.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/kernels-decompose-2.f95
@@ -119,6 +119,7 @@ program main
!$acc end kernels
!$acc kernels
+ ! { dg-warning "not gang partitioned" "" { target *-*-* } .-1 }
y = f_g (a(5)) ! { dg-line l_part[incr c_part] }
!TODO If such a construct is placed in its own part (like it is, here), can't this actually use gang paralelism, instead of "gang-single"?
! { dg-message "note: beginning 'gang-single' part in OpenACC 'kernels' region" "" { target *-*-* } l_part$c_part }
diff --git a/gcc/testsuite/gfortran.dg/goacc/parallel-tree.f95 b/gcc/testsuite/gfortran.dg/goacc/parallel-tree.f95
index e33653bdd78..83aaf699cdb 100644
--- a/gcc/testsuite/gfortran.dg/goacc/parallel-tree.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/parallel-tree.f95
@@ -12,6 +12,9 @@ program test
!$acc no_create(n) &
!$acc present(o), pcopy(p), pcopyin(r), pcopyout(s), pcreate(t) &
!$acc deviceptr(u), private(v), firstprivate(w)
+ ! { dg-warning "region is gang partitioned but does not contain gang partitioned code" "" { target *-*-* } .-1 }
+ ! { dg-warning "region is worker partitioned but does not contain worker partitioned code" "" { target *-*-* } .-2 }
+ ! { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .-3 }
!$acc end parallel
end program test
diff --git a/gcc/testsuite/gfortran.dg/goacc/routine-4.f90 b/gcc/testsuite/gfortran.dg/goacc/routine-4.f90
index 6714c7b8229..3fb60e71ae0 100644
--- a/gcc/testsuite/gfortran.dg/goacc/routine-4.f90
+++ b/gcc/testsuite/gfortran.dg/goacc/routine-4.f90
@@ -123,6 +123,7 @@ contains
integer, intent (inout) :: a(N)
integer :: i
+ !$acc loop gang worker vector
do i = 1, N
a(i) = a(i) - a(i)
end do
@@ -133,6 +134,7 @@ contains
integer, intent (inout) :: a(N)
integer :: i
+ !$acc loop worker vector
do i = 1, N
a(i) = a(i) - a(i)
end do
@@ -143,6 +145,7 @@ contains
integer, intent (inout) :: a(N)
integer :: i
+ !$acc loop vector
do i = 1, N
a(i) = a(i) - a(i)
end do
@@ -153,6 +156,7 @@ contains
integer, intent (inout) :: a(N)
integer :: i
+ !$acc loop seq
do i = 1, N
a(i) = a(i) - a(i)
end do
diff --git a/gcc/testsuite/gfortran.dg/goacc/routine-level-of-parallelism-1.f90 b/gcc/testsuite/gfortran.dg/goacc/routine-level-of-parallelism-1.f90
index 83b8c24b41d..9986c0c8d50 100644
--- a/gcc/testsuite/gfortran.dg/goacc/routine-level-of-parallelism-1.f90
+++ b/gcc/testsuite/gfortran.dg/goacc/routine-level-of-parallelism-1.f90
@@ -2,8 +2,10 @@
! with the OpenACC routine directive. The C/C++ counterpart is
! '../../c-c++-common/goacc/routine-level-of-parallelism-2.c'.
-subroutine g_1
+subroutine g_1 ! { dg-warning "region is gang partitioned but does not contain gang partitioned code" }
!$acc routine gang
+! { dg-bogus "region is worker partitioned but does not contain worker partitioned code" "worker partitioned" { xfail *-*-* } .-2 }
+! { dg-bogus "region is vector partitioned but does not contain vector partitioned code" "worker partitioned" { xfail *-*-* } .-3 }
end subroutine g_1
subroutine s_1_2a
diff --git a/gcc/testsuite/gfortran.dg/goacc/routine-module-mod-1.f90 b/gcc/testsuite/gfortran.dg/goacc/routine-module-mod-1.f90
index 23c673fe3bd..ce4466ab581 100644
--- a/gcc/testsuite/gfortran.dg/goacc/routine-module-mod-1.f90
+++ b/gcc/testsuite/gfortran.dg/goacc/routine-module-mod-1.f90
@@ -50,6 +50,7 @@ contains
end do
end subroutine w_1
+ ! { dg-warning "does not contain worker partitioned code" "" { target *-*-* } .+1 }
subroutine g_1
implicit none
!$acc routine gang
diff --git a/gcc/testsuite/gfortran.dg/goacc/routine-multiple-directives-1.f90 b/gcc/testsuite/gfortran.dg/goacc/routine-multiple-directives-1.f90
index 6e12ee92155..e39f6b58b25 100644
--- a/gcc/testsuite/gfortran.dg/goacc/routine-multiple-directives-1.f90
+++ b/gcc/testsuite/gfortran.dg/goacc/routine-multiple-directives-1.f90
@@ -12,14 +12,14 @@
!$ACC ROUTINE(s_2)
END SUBROUTINE s_2
- SUBROUTINE v_1
+ SUBROUTINE v_1 ! { dg-warning "region is vector partitioned but does not contain vector partitioned code" }
!$ACC ROUTINE VECTOR
!$ACC ROUTINE VECTOR
!$ACC ROUTINE(v_1) VECTOR
!$ACC ROUTINE VECTOR
END SUBROUTINE v_1
- SUBROUTINE v_2
+ SUBROUTINE v_2 ! { dg-warning "region is vector partitioned but does not contain vector partitioned code" }
!$ACC ROUTINE(v_2) VECTOR
!$ACC ROUTINE VECTOR
!$ACC ROUTINE(v_2) VECTOR
diff --git a/gcc/testsuite/gfortran.dg/goacc/uninit-dim-clause.f95 b/gcc/testsuite/gfortran.dg/goacc/uninit-dim-clause.f95
index c77d47a3900..b513a56430f 100644
--- a/gcc/testsuite/gfortran.dg/goacc/uninit-dim-clause.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/uninit-dim-clause.f95
@@ -5,12 +5,15 @@ subroutine acc_parallel
integer :: i, j, k
!$acc parallel num_gangs(i) ! { dg-warning "is used uninitialized" }
+ ! { dg-warning "does not contain gang partitioned code" "" { target *-*-* } .-1 }
!$acc end parallel
!$acc parallel num_workers(j) ! { dg-warning "is used uninitialized" }
+ ! { dg-warning "does not contain worker partitioned code" "" { target *-*-* } .-1 }
!$acc end parallel
!$acc parallel vector_length(k) ! { dg-warning "is used uninitialized" }
+ ! { dg-warning "does not contain vector partitioned code" "" { target *-*-* } .-1 }
!$acc end parallel
end subroutine acc_parallel
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/firstprivate-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/firstprivate-1.c
index 689a443ca43..14bc3af4a97 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/firstprivate-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/firstprivate-1.c
@@ -117,6 +117,8 @@ void t4 ()
arr[i] = 3;
#pragma acc parallel firstprivate(x) copy(arr) num_gangs(32) num_workers(8) vector_length(32)
+ /* { dg-warning "region is worker partitioned but does not contain worker partitioned code" "worker" { target *-*-* } 119 } */
+ /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } 119 } */
{
#pragma acc loop gang
for (i = 0; i < 32; i++)
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-auto-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-auto-1.c
index 0273c2bddd7..c6ba9ab8acc 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-auto-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-auto-1.c
@@ -150,7 +150,7 @@ int gang_1 (int *ary, int size)
{
clear (ary, size);
-#pragma acc parallel num_gangs (32) num_workers (32) vector_length(32) copy(ary[0:size]) firstprivate (size)
+#pragma acc parallel num_gangs (32) num_workers (32) vector_length(32) copy(ary[0:size]) firstprivate (size)/* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } } */
{
#pragma acc loop auto
for (int jx = 0; jx < size / 64; jx++)
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-g-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-g-1.c
index 98f02e9840a..5831327d0ae 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-g-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-g-1.c
@@ -1,3 +1,4 @@
+/* { dg-additional-options "-w" } */
#include <stdio.h>
#include <openacc.h>
#include <gomp-constants.h>
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-g-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-g-2.c
index 4152a4e6c82..82e8aae88ac 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-g-2.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-g-2.c
@@ -1,3 +1,4 @@
+/* { dg-additional-options "-w" } */
#include <stdio.h>
#include <openacc.h>
#include <gomp-constants.h>
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-g-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-g-1.c
index 7107502e070..2f3a44f3ba7 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-g-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-g-1.c
@@ -1,3 +1,4 @@
+/* { dg-additional-options "-w" } */
#include <stdio.h>
#include <openacc.h>
#include <gomp-constants.h>
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-w-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-w-1.c
index 648f89e1668..7344fa8bf00 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-w-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-w-1.c
@@ -1,3 +1,4 @@
+/* { dg-additional-options "-w" } */
#include <stdio.h>
#include <openacc.h>
#include <gomp-constants.h>
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-w-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-w-2.c
index f9fcf3703af..d99877ab8a9 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-w-2.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-w-2.c
@@ -1,3 +1,4 @@
+/* { dg-additional-options "-w" } */
#include <stdio.h>
#include <openacc.h>
#include <gomp-constants.h>
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-w-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-w-1.c
index 5fe486f50a1..8731c805b79 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-w-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-w-1.c
@@ -16,6 +16,7 @@ int main ()
#pragma acc parallel num_workers(32) vector_length(32) copy(ary) copy(ondev) \
copyout(workersize)
+ /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } 17 } */
{
#pragma acc loop worker
for (unsigned ix = 0; ix < N; ix++)
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-warn-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-warn-1.c
new file mode 100644
index 00000000000..20a022f2758
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-warn-1.c
@@ -0,0 +1,37 @@
+
+/* Check warnings about suboptimal partitioning choices. */
+
+int main ()
+{
+ int ary[10];
+
+#pragma acc parallel copy(ary) num_gangs (1) /* { dg-warning "is not gang partitioned" } */
+ {
+ #pragma acc loop gang
+ for (int i = 0; i < 10; i++)
+ ary[i] = i;
+ }
+
+#pragma acc parallel copy(ary) num_workers (1) /* { dg-warning "is not worker partitioned" } */
+ {
+ #pragma acc loop worker
+ for (int i = 0; i < 10; i++)
+ ary[i] = i;
+ }
+
+#pragma acc parallel copy(ary) num_gangs (8) /* { dg-warning "is gang partitioned" } */
+ {
+ #pragma acc loop worker
+ for (int i = 0; i < 10; i++)
+ ary[i] = i;
+ }
+
+#pragma acc parallel copy(ary) num_workers (8) /* { dg-warning "is worker partitioned" } */
+ {
+ #pragma acc loop gang
+ for (int i = 0; i < 10; i++)
+ ary[i] = i;
+ }
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/mode-transitions.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/mode-transitions.c
index 4474c127992..f62daf031a0 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/mode-transitions.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/mode-transitions.c
@@ -287,6 +287,7 @@ void t7()
int n = 0;
#pragma acc parallel copy(n) \
num_gangs(1) num_workers(1) vector_length(32)
+ /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } 288 } */
{
n++;
}
@@ -310,6 +311,7 @@ void t8()
#pragma acc parallel copy(arr) \
num_gangs(gangs) num_workers(1) vector_length(32)
+ /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } 312 } */
{
int j;
#pragma acc loop gang
@@ -339,6 +341,7 @@ void t9()
#pragma acc parallel copy(arr) \
num_gangs(gangs) num_workers(1) vector_length(32)
+ /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } 342 } */
{
int j;
#pragma acc loop gang
@@ -371,6 +374,7 @@ void t10()
#pragma acc parallel copy(arr) \
num_gangs(gangs) num_workers(1) vector_length(32)
+ /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } 375 } */
{
int j;
#pragma acc loop gang
@@ -404,6 +408,7 @@ void t11()
#pragma acc parallel copy(arr) \
num_gangs(1024) num_workers(1) vector_length(32)
+ /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } 409 } */
{
int j;
@@ -442,6 +447,7 @@ void t12()
#pragma acc parallel copyout(fizz, buzz, fizzbuzz) \
num_gangs(NUM_GANGS) num_workers(1) vector_length(32)
+ /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } 448 } */
{
int j;
@@ -488,6 +494,7 @@ void t13()
#pragma acc parallel copy(arr) \
num_gangs(8) num_workers(8) vector_length(32)
+ /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } 495 } */
{
int j;
#pragma acc loop gang
@@ -613,6 +620,7 @@ void t16()
#pragma acc parallel copy(n, arr) \
num_gangs(8) num_workers(16) vector_length(32)
+ /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } 621 } */
{
int j;
#pragma acc loop gang
@@ -665,6 +673,7 @@ void t17()
#pragma acc parallel copyin(arr_a) copyout(arr_b) \
num_gangs(num_gangs) num_workers(num_workers) vector_length(32)
+ /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } 674 } */
{
int j;
#pragma acc loop gang
@@ -882,6 +891,8 @@ void t21()
#pragma acc parallel copy(arr) \
num_gangs(8) num_workers(8) vector_length(32)
+ /* { dg-warning "region is worker partitioned but does not contain worker partitioned code" "worker" { target *-*-* } 892 } */
+ /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } 892 } */
{
int j;
#pragma acc loop gang
@@ -905,6 +916,8 @@ void t22()
#pragma acc parallel copy(arr) \
num_gangs(8) num_workers(8) vector_length(32)
+ /* { dg-warning "region is worker partitioned but does not contain worker partitioned code" "worker" { target *-*-* } 917 } */
+ /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } 917 } */
{
int j;
#pragma acc loop gang
@@ -931,6 +944,8 @@ void t23()
#pragma acc parallel copy(arr) \
num_gangs(8) num_workers(8) vector_length(32)
+ /* { dg-warning "region is worker partitioned but does not contain worker partitioned code" "worker" { target *-*-* } 945 } */
+ /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } 945 } */
{
int j;
#pragma acc loop gang
@@ -957,6 +972,8 @@ void t24()
#pragma acc parallel copy(arr) \
num_gangs(8) num_workers(8) vector_length(32)
+ /* { dg-warning "region is worker partitioned but does not contain worker partitioned code" "worker" { target *-*-* } 973 } */
+ /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } 973 } */
{
int j;
#pragma acc loop gang
@@ -988,6 +1005,7 @@ void t25()
#pragma acc parallel copy(arr) \
num_gangs(8) num_workers(8) vector_length(32)
+ /* { dg-warning "region is worker partitioned but does not contain worker partitioned code" "worker" { target *-*-* } 1006 } */
{
int j;
#pragma acc loop gang
@@ -1020,6 +1038,7 @@ void t26()
#pragma acc parallel copy(arr) \
num_gangs(8) num_workers(8) vector_length(32)
+ /* { dg-warning "region is worker partitioned but does not contain worker partitioned code" "worker" { target *-*-* } 1039 } */
{
int j;
#pragma acc loop gang
@@ -1070,6 +1089,8 @@ void t27()
#pragma acc parallel copy(n, arr) copyout(ondev) \
num_gangs(ACTUAL_GANGS) num_workers(8) vector_length(32)
+ /* { dg-warning "region is gang partitioned but does not contain gang partitioned code" "gang" { target *-*-* } 1090 } */
+ /* { dg-warning "region is worker partitioned but does not contain worker partitioned code" "worker" { target *-*-* } 1090 } */
{
int j;
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/parallel-dims.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/parallel-dims.c
index cc4c738c1db..003bcac2413 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/parallel-dims.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/parallel-dims.c
@@ -103,7 +103,7 @@ int main ()
int gangs_min, gangs_max, workers_min, workers_max, vectors_min, vectors_max;
gangs_min = workers_min = vectors_min = INT_MAX;
gangs_max = workers_max = vectors_max = INT_MIN;
-#pragma acc parallel copy (gangs_actual) \
+#pragma acc parallel copy (gangs_actual) /* { dg-warning "region contains gang partitioned code but is not gang partitioned" } */ \
num_gangs (GANGS) /* { dg-warning "'num_gangs' value must be positive" "" { target c++ } } */
{
/* We're actually executing with num_gangs (1). */
@@ -132,7 +132,7 @@ int main ()
int gangs_min, gangs_max, workers_min, workers_max, vectors_min, vectors_max;
gangs_min = workers_min = vectors_min = INT_MAX;
gangs_max = workers_max = vectors_max = INT_MIN;
-#pragma acc parallel copy (workers_actual) \
+#pragma acc parallel copy (workers_actual) /* { dg-warning "region contains worker partitioned code but is not worker partitioned" } */ \
num_workers (WORKERS) /* { dg-warning "'num_workers' value must be positive" "" { target c++ } } */
{
/* We're actually executing with num_workers (1). */
@@ -161,7 +161,8 @@ int main ()
int gangs_min, gangs_max, workers_min, workers_max, vectors_min, vectors_max;
gangs_min = workers_min = vectors_min = INT_MAX;
gangs_max = workers_max = vectors_max = INT_MIN;
-#pragma acc parallel copy (vectors_actual) /* { dg-warning "using vector_length \\(32\\), ignoring 1" "" { target openacc_nvidia_accel_selected } } */ \
+#pragma acc parallel copy (vectors_actual) /* { dg-warning "region contains vector partitioned code but is not vector partitioned" } */ \
+ /* { dg-warning "using vector_length \\(32\\), ignoring 1" "" { target openacc_nvidia_accel_selected } 164 } */ \
vector_length (VECTORS) /* { dg-warning "'vector_length' value must be positive" "" { target c++ } } */
{
/* We're actually executing with vector_length (1), just the GCC nvptx
@@ -205,7 +206,7 @@ int main ()
int gangs_min, gangs_max, workers_min, workers_max, vectors_min, vectors_max;
gangs_min = workers_min = vectors_min = INT_MAX;
gangs_max = workers_max = vectors_max = INT_MIN;
-#pragma acc parallel copy (gangs_actual) \
+#pragma acc parallel copy (gangs_actual) /* { dg-warning "region is gang partitioned but does not contain gang partitioned code" } */ \
reduction (min: gangs_min, workers_min, vectors_min) reduction (max: gangs_max, workers_max, vectors_max) \
num_gangs (gangs)
{
@@ -617,6 +618,9 @@ int main ()
gangs_max = workers_max = vectors_max = INT_MIN;
#pragma acc serial copy (vectors_actual) /* { dg-warning "using vector_length \\(32\\), ignoring 1" "" { target openacc_nvidia_accel_selected } } */ \
copy (gangs_min, gangs_max, workers_min, workers_max, vectors_min, vectors_max)
+/* { dg-warning "not gang partitioned" "" { target *-*-* } 619 } */
+/* { dg-warning "not worker partitioned" "" { target *-*-* } 619 } */
+/* { dg-warning "not vector partitioned" "" { target *-*-* } 619 } */
{
if (acc_on_device (acc_device_nvidia))
{
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85381-3.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85381-3.c
index cddbf271906..53fb049f992 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85381-3.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85381-3.c
@@ -1,6 +1,6 @@
/* { dg-do run { target openacc_nvidia_accel_selected } }
{ dg-skip-if "" { *-*-* } { "*" } { "-O2" } } */
-/* { dg-additional-options "-foffload=-fdump-rtl-mach" } */
+/* { dg-additional-options "-foffload=-fdump-rtl-mach -w" } */
int a;
#pragma acc declare create(a)
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/private-variables.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/private-variables.c
index 53f03d17bb2..f0c3447dfef 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/private-variables.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/private-variables.c
@@ -22,6 +22,8 @@ void local_g_1()
arr[i] = 3;
#pragma acc parallel copy(arr) num_gangs(32) num_workers(8) vector_length(32)
+ /* { dg-warning "region is worker partitioned but does not contain worker partitioned code" "worker" { target *-*-* } 24 } */
+ /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } 24 } */
{
int x;
@@ -295,6 +297,8 @@ void loop_g_1()
arr[i] = i;
#pragma acc parallel copy(arr) num_gangs(32) num_workers(8) vector_length(32)
+ /* { dg-warning "region is worker partitioned but does not contain worker partitioned code" "worker" { target *-*-* } 299 } */
+ /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } 299 } */
{
#pragma acc loop gang private(x)
for (i = 0; i < 32; i++)
@@ -320,6 +324,7 @@ void loop_g_2()
arr[i] = i;
#pragma acc parallel copy(arr) num_gangs(32) num_workers(8) vector_length(32)
+ /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } 326 } */
{
#pragma acc loop gang private(x)
for (i = 0; i < 32; i++)
@@ -348,6 +353,7 @@ void loop_g_3()
arr[i] = i;
#pragma acc parallel copy(arr) num_gangs(32) num_workers(8) vector_length(32)
+ /* { dg-warning "region is worker partitioned but does not contain worker partitioned code" "worker" { target *-*-* } 355 } */
{
#pragma acc loop gang private(x)
for (i = 0; i < 32; i++)
@@ -376,6 +382,7 @@ void loop_g_4()
arr[i] = i;
#pragma acc parallel copy(arr) num_gangs(32) num_workers(8) vector_length(32)
+ /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } 384 } */
{
#pragma acc loop gang private(x)
for (i = 0; i < 32; i++)
@@ -408,6 +415,7 @@ void loop_g_5()
arr[i] = i;
#pragma acc parallel copy(arr) num_gangs(32) num_workers(8) vector_length(32)
+ /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } 417 } */
{
#pragma acc loop gang private(x)
for (i = 0; i < 32; i++)
@@ -438,6 +446,7 @@ void loop_g_6()
arr[i] = i;
#pragma acc parallel copy(arr) num_gangs(32) num_workers(8) vector_length(32)
+ /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } 448 } */
{
#pragma acc loop gang private(pt)
for (i = 0; i < 32; i++)
@@ -559,6 +568,7 @@ void loop_w_1()
arr[i] = i;
#pragma acc parallel copy(arr) num_gangs(32) num_workers(8) vector_length(32)
+ /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } 570 } */
{
int j;
@@ -875,6 +885,8 @@ void parallel_g_1()
arr[i] = 3;
#pragma acc parallel private(x) copy(arr) num_gangs(32) num_workers(8) vector_length(32)
+ /* { dg-warning "region is worker partitioned but does not contain worker partitioned code" "worker" { target *-*-* } 887 } */
+ /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } 887 } */
{
#pragma acc loop gang(static:1)
for (i = 0; i < 32; i++)
@@ -904,6 +916,7 @@ void parallel_g_2()
arr[i] = i;
#pragma acc parallel private(x) copy(arr) num_gangs(32) num_workers(2) vector_length(32)
+ /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } 918 } */
{
#pragma acc loop gang
for (i = 0; i < 32; i++)
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/reduction-7.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/reduction-7.c
index c4940b8ad9d..68ae919ea6f 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/reduction-7.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/reduction-7.c
@@ -14,6 +14,8 @@ void g_np_1()
arr[i] = i;
#pragma acc parallel num_gangs(32) num_workers(32) vector_length(32)
+ /* { dg-warning "region is worker partitioned but does not contain worker partitioned code" "worker" { target *-*-* } 16 } */
+ /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } 16 } */
{
#pragma acc loop gang reduction(+:res)
for (i = 0; i < 1024; i++)
@@ -28,6 +30,8 @@ void g_np_1()
res = hres = 1;
#pragma acc parallel num_gangs(32) num_workers(32) vector_length(32)
+ /* { dg-warning "region is worker partitioned but does not contain worker partitioned code" "worker" { target *-*-* } 32 } */
+ /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } 32 } */
{
#pragma acc loop gang reduction(*:res)
for (i = 0; i < 12; i++)
@@ -52,6 +56,7 @@ void gv_np_1()
arr[i] = i;
#pragma acc parallel num_gangs(32) num_workers(32) vector_length(32)
+ /* { dg-warning "region is worker partitioned but does not contain worker partitioned code" "worker" { target *-*-* } 58 } */
{
#pragma acc loop gang vector reduction(+:res)
for (i = 0; i < 1024; i++)
@@ -76,6 +81,7 @@ void gw_np_1()
arr[i] = i;
#pragma acc parallel num_gangs(32) num_workers(32) vector_length(32)
+ /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } 83 } */
{
#pragma acc loop gang worker reduction(+:res)
for (i = 0; i < 1024; i++)
@@ -239,6 +245,7 @@ void v_p_1()
#pragma acc parallel num_gangs(32) num_workers(32) vector_length(32) \
private(res) copyout(out)
+ /* { dg-warning "region is worker partitioned but does not contain worker partitioned code" "worker" { target *-*-* } 246 } */
{
#pragma acc loop gang
for (j = 0; j < 32; j++)
@@ -315,6 +322,7 @@ void w_p_1()
#pragma acc parallel num_gangs(32) num_workers(32) vector_length(32) \
private(res) copyout(out)
+ /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } 323 } */
{
#pragma acc loop gang
for (j = 0; j < 32; j++)
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-g-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-g-1.c
index a164f576bc3..8c3b938506f 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-g-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-g-1.c
@@ -6,6 +6,8 @@
#pragma acc routine gang
void __attribute__ ((noinline)) gang (int ary[N])
+/* { dg-warning "region is worker partitioned but does not contain worker partitioned code" "worker" { target *-*-* } 8 } */
+/* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } 8 } */
{
#pragma acc loop gang
for (unsigned ix = 0; ix < N; ix++)
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-w-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-w-1.c
index acd9884cbd6..bf9228105c5 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-w-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-w-1.c
@@ -6,6 +6,7 @@
#pragma acc routine worker
void __attribute__ ((noinline)) worker (int ary[N])
+/* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } 8 } */
{
#pragma acc loop worker
for (unsigned ix = 0; ix < N; ix++)
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-wv-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-wv-2.c
index 9769ee72430..de167c55c76 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-wv-2.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-wv-2.c
@@ -42,7 +42,8 @@ int DoWorkVec (int nw)
ary[ix][jx] = 0xdeadbeef;
printf ("spawning %d ...", nw); fflush (stdout);
-
+
+ /* { dg-warning "region contains vector partitioned code but is not vector partitioned" "vector" { target openacc_radeon_accel_selected } 47 } */
#pragma acc parallel num_workers(nw) vector_length (NUM_VECTORS) copy (ary)
{
WorkVec ((int *)ary, WIDTH, HEIGHT, nw, NUM_VECTORS);
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/optional-private.f90 b/libgomp/testsuite/libgomp.oacc-fortran/optional-private.f90
index 0320bbb3bc9..6bc91b7a0bb 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/optional-private.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/optional-private.f90
@@ -30,6 +30,8 @@ contains
end do
!$acc parallel copy(arr) num_gangs(32) num_workers(8) vector_length(32)
+ ! { dg-warning "region is worker partitioned but does not contain worker partitioned code" "worker" { target *-*-* } 32 }
+ ! { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } 32 }
!$acc loop gang private(x)
do i = 1, 32
x = i * 2;
@@ -55,6 +57,7 @@ contains
end do
!$acc parallel copy(arr) num_gangs(32) num_workers(8) vector_length(32)
+ ! { dg-warning "region is worker partitioned but does not contain worker partitioned code" "worker" { target *-*-* } 59 }
!$acc loop gang private(pt)
do i = 0, 31
pt%x = i
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/par-reduction-2-1.f b/libgomp/testsuite/libgomp.oacc-fortran/par-reduction-2-1.f
index aa1bb634ba6..ff31116d980 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/par-reduction-2-1.f
+++ b/libgomp/testsuite/libgomp.oacc-fortran/par-reduction-2-1.f
@@ -14,7 +14,7 @@
RES2 = 0
!$ACC PARALLEL NUM_GANGS(256) NUM_WORKERS(32) VECTOR_LENGTH(32)
-!$ACC& REDUCTION(+:RES1) COPY(RES1, RES2) ASYNC(1)
+!$ACC& REDUCTION(+:RES1) COPY(RES1, RES2) ASYNC(1) ! { dg-warning "region is (gang|worker|vector) partitioned" }
res1 = res1 + 5
!$ACC ATOMIC
@@ -36,7 +36,7 @@
RES2 = 1
!$ACC PARALLEL NUM_GANGS(8) NUM_WORKERS(32) VECTOR_LENGTH(32)
-!$ACC& REDUCTION(*:RES1) COPY(RES1, RES2) ASYNC(1)
+!$ACC& REDUCTION(*:RES1) COPY(RES1, RES2) ASYNC(1) ! { dg-warning "region is (gang|worker|vector) partitioned" }
res1 = res1 * 5
!$ACC ATOMIC
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/par-reduction-2-2.f b/libgomp/testsuite/libgomp.oacc-fortran/par-reduction-2-2.f
index 5694de1aee8..47c5ff3474d 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/par-reduction-2-2.f
+++ b/libgomp/testsuite/libgomp.oacc-fortran/par-reduction-2-2.f
@@ -14,7 +14,7 @@
RES2 = 0
!$ACC PARALLEL NUM_GANGS(256) NUM_WORKERS(32) VECTOR_LENGTH(32)
-!$ACC& REDUCTION(+:RES1) COPY(RES1, RES2) ASYNC(1)
+!$ACC& REDUCTION(+:RES1) COPY(RES1, RES2) ASYNC(1) ! { dg-warning "region is (gang|worker|vector) partitioned" }
res1 = res1 + 5
!$ACC ATOMIC
@@ -36,7 +36,7 @@
RES2 = 1
!$ACC PARALLEL NUM_GANGS(8) NUM_WORKERS(32) VECTOR_LENGTH(32)
-!$ACC& REDUCTION(*:RES1) COPY(RES1, RES2) ASYNC(1)
+!$ACC& REDUCTION(*:RES1) COPY(RES1, RES2) ASYNC(1) ! { dg-warning "region is (gang|worker|vector) partitioned" }
res1 = res1 * 5
!$ACC ATOMIC
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/parallel-dims.f90 b/libgomp/testsuite/libgomp.oacc-fortran/parallel-dims.f90
index 1bfcd6ce099..b515c62391b 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/parallel-dims.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/parallel-dims.f90
@@ -84,6 +84,9 @@ program main
vectors_max = -huge(gangs_max) - 1 ! INT_MIN
!$acc serial copy (vectors_actual) &
!$acc copy (gangs_min, gangs_max, workers_min, workers_max, vectors_min, vectors_max) ! { dg-warning "using vector_length \\(32\\), ignoring 1" "" { target openacc_nvidia_accel_selected } }
+ ! { dg-warning "not gang partitioned" "" { target *-*-* } 86 }
+ ! { dg-warning "not worker partitioned" "" { target *-*-* } 86 }
+ ! { dg-warning "not vector partitioned" "" { target *-*-* } 86 }
if (acc_on_device (acc_device_nvidia)) then
! The GCC nvptx back end enforces vector_length (32).
! It's unclear if that's actually permissible here;
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/pr84028.f90 b/libgomp/testsuite/libgomp.oacc-fortran/pr84028.f90
index 2b361220bb6..8cb76a93d7f 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/pr84028.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/pr84028.f90
@@ -5,7 +5,7 @@ program foo
a = 1
- !$acc parallel num_gangs(1) num_workers(2)
+ !$acc parallel num_gangs(1) num_workers(2) ! { dg-warning "region is worker partitioned" }
if (any(a(1:3,1:3,1:3).ne.1)) STOP 1
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/private-variables.f90 b/libgomp/testsuite/libgomp.oacc-fortran/private-variables.f90
index 472a6a14fff..fbff5ccd46e 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/private-variables.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/private-variables.f90
@@ -13,6 +13,8 @@ subroutine t1()
end do
!$acc parallel copy(arr) num_gangs(32) num_workers(8) vector_length(32)
+ ! { dg-warning "region is worker partitioned but does not contain worker partitioned code" "worker" { target *-*-* } 15 }
+ ! { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } 15 }
!$acc loop gang private(x)
do i = 1, 32
x = i * 2;
@@ -37,6 +39,7 @@ subroutine t2()
end do
!$acc parallel copy(arr) num_gangs(32) num_workers(8) vector_length(32)
+ ! { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } 41 }
!$acc loop gang private(x)
do i = 0, 31
x = i * 2;
@@ -65,6 +68,7 @@ subroutine t3()
end do
!$acc parallel copy(arr) num_gangs(32) num_workers(8) vector_length(32)
+ ! { dg-warning "region is worker partitioned but does not contain worker partitioned code" "worker" { target *-*-* } 70 }
!$acc loop gang private(x)
do i = 0, 31
x = i * 2;
@@ -98,6 +102,7 @@ subroutine t4()
end do
!$acc parallel copy(arr) num_gangs(32) num_workers(8) vector_length(32)
+ ! { dg-warning "region is worker partitioned but does not contain worker partitioned code" "worker" { target *-*-* } 104 }
!$acc loop gang private(pt)
do i = 0, 31
pt%x = i
@@ -208,6 +213,7 @@ subroutine t7()
end do
!$acc parallel copy(arr) num_gangs(32) num_workers(8) vector_length(32)
+ ! { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } 215 }
!$acc loop gang private(x)
do i = 0, 31
!$acc loop worker private(x)
@@ -507,6 +513,8 @@ subroutine t14()
end do
!$acc parallel private(x) copy(arr) num_gangs(n) num_workers(8) vector_length(32)
+ ! { dg-warning "region is worker partitioned but does not contain worker partitioned code" "worker" { target *-*-* } 515 }
+ ! { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } 515 }
!$acc loop gang(static:1)
do i = 1, n
x = i * 2;
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/routine-7.f90 b/libgomp/testsuite/libgomp.oacc-fortran/routine-7.f90
index 1009f4a81e5..51d52456733 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/routine-7.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/routine-7.f90
@@ -100,7 +100,7 @@ subroutine gang (a)
integer, intent (inout) :: a(N)
integer :: i
- !$acc loop gang
+ !$acc loop gang worker vector
do i = 1, N
a(i) = a(i) - i
end do
--
2.29.2
^ permalink raw reply [flat|nested] 7+ messages in thread
* Re: [PATCH] openacc: Warnings for strange partitioning choices for parallel regions
2021-02-26 12:21 [PATCH] openacc: Warnings for strange partitioning choices for parallel regions Julian Brown
@ 2021-04-26 10:39 ` Thomas Schwinge
2021-04-26 12:32 ` Tobias Burnus
0 siblings, 1 reply; 7+ messages in thread
From: Thomas Schwinge @ 2021-04-26 10:39 UTC (permalink / raw)
To: Julian Brown, gcc-patches; +Cc: Jakub Jelinek, Tobias Burnus
[-- Attachment #1: Type: text/plain, Size: 2449 bytes --]
Hi!
On 2021-02-26T04:21:54-0800, Julian Brown <julian@codesourcery.com> wrote:
> This patch adds warnings for strange partitioning choices -- specifying
> either more or fewer partitioning levels on a parallel directive than the
> enclosed offload region actually uses.
Thanks!
> We've used a version of this patch on the og8/og9/og10 branches for
> quite a while. Versions have been posted for mainline submission as
> part of a larger patch several times, e.g. here:
>
> https://gcc.gnu.org/pipermail/gcc-patches/2018-October/507938.html
>
> One motivation for committing this patch is it removes an ongoing
> divergence in a large number of test cases between the og* branches and
> mainline, namely whether the added warnings are expected as a result of
> compiling those test cases, or not.
>
> The warnings themselves are perhaps slightly aggressive, but are intended
> to help the user write more efficient code.
As I'd said in
<http://mid.mail-archive.com/87zhoefcpf.fsf@euler.schwinge.homeip.net>:
| The general intention is good, but I've seen cases where I considered
| these diagnostics to be too noisy. See also the several 'dg-bogus'
| with XFAIL
... as well as 'dg-additional-options "-w"' being added.
I've spent some more time on this; in particular, 'dg-bogus' with XFAIL
setup, to indicate individual items that I think need to be addressed.
(Please let me know if anybody makes a different categorization regarding
these.) Of course, not all of our testcases are representative of
real-world code (but I've anyway removed '-w' in favor of 'dg-warning' or
'dg-bogus' with XFAIL, for coverage), but I still came to the conclusion
that this is too noisy to enable by default (as had been proposed), and
thus for the time being have hidden it behind '-Wopenacc-parallelism'.
(Adding an option is good anyway, so that users are able to disable
unwanted diagnostics; we need to do that for other OpenACC diagnostics,
too.) The goal is to eventually enable '-Wopenacc-parallelism', with
'-Wall' and some "special cases" with '-Wextra', I suppose.
I've pushed "Add '-Wopenacc-parallelism'" to master branch in commit
22cff118f7526bec195ed6e41452980820fdf3a8, see attached.
Grüße
Thomas
-----------------
Mentor Graphics (Deutschland) GmbH, Arnulfstrasse 201, 80634 München Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Frank Thürauf
[-- Warning: decoded text below may be mangled, UTF-8 assumed --]
[-- Attachment #2: 0001-Add-Wopenacc-parallelism.patch --]
[-- Type: text/x-diff, Size: 88327 bytes --]
From 22cff118f7526bec195ed6e41452980820fdf3a8 Mon Sep 17 00:00:00 2001
From: Thomas Schwinge <thomas@codesourcery.com>
Date: Fri, 23 Apr 2021 12:23:51 +0200
Subject: [PATCH] Add '-Wopenacc-parallelism'
... to diagnose potentially suboptimal choices regarding OpenACC parallelism.
Not enabled by default: too noisy ("*potentially* suboptimal choices"); see
XFAILed 'dg-bogus'es.
gcc/c-family/
* c.opt (Wopenacc-parallelism): New.
gcc/fortran/
* lang.opt (Wopenacc-parallelism): New.
gcc/
* omp-offload.c (oacc_validate_dims): Implement
'-Wopenacc-parallelism'.
* doc/invoke.texi (-Wopenacc-parallelism): Document.
gcc/testsuite/
* c-c++-common/goacc/diag-parallelism-1.c: New.
* c-c++-common/goacc/acc-icf.c: Specify '-Wopenacc-parallelism',
and match diagnostics, as appropriate.
* c-c++-common/goacc/classify-kernels-unparallelized.c: Likewise.
* c-c++-common/goacc/classify-kernels.c: Likewise.
* c-c++-common/goacc/classify-parallel.c: Likewise.
* c-c++-common/goacc/classify-routine.c: Likewise.
* c-c++-common/goacc/classify-serial.c: Likewise.
* c-c++-common/goacc/kernels-decompose-1.c: Likewise.
* c-c++-common/goacc/kernels-decompose-2.c: Likewise.
* c-c++-common/goacc/parallel-dims-1.c: Likewise.
* c-c++-common/goacc/parallel-reduction.c: Likewise.
* c-c++-common/goacc/pr70688.c: Likewise.
* c-c++-common/goacc/routine-1.c: Likewise.
* c-c++-common/goacc/routine-level-of-parallelism-2.c: Likewise.
* c-c++-common/goacc/uninit-dim-clause.c: Likewise.
* gfortran.dg/goacc/classify-kernels-unparallelized.f95: Likewise.
* gfortran.dg/goacc/classify-kernels.f95: Likewise.
* gfortran.dg/goacc/classify-parallel.f95: Likewise.
* gfortran.dg/goacc/classify-routine.f95: Likewise.
* gfortran.dg/goacc/classify-serial.f95: Likewise.
* gfortran.dg/goacc/kernels-decompose-1.f95: Likewise.
* gfortran.dg/goacc/kernels-decompose-2.f95: Likewise.
* gfortran.dg/goacc/parallel-tree.f95: Likewise.
* gfortran.dg/goacc/routine-4.f90: Likewise.
* gfortran.dg/goacc/routine-level-of-parallelism-1.f90: Likewise.
* gfortran.dg/goacc/routine-module-mod-1.f90: Likewise.
* gfortran.dg/goacc/routine-multiple-directives-1.f90: Likewise.
* gfortran.dg/goacc/uninit-dim-clause.f95: Likewise.
libgomp/
* testsuite/libgomp.oacc-c-c++-common/firstprivate-1.c: Specify
'-Wopenacc-parallelism', and match diagnostics, as appropriate.
* testsuite/libgomp.oacc-c-c++-common/loop-auto-1.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/loop-red-w-1.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/loop-red-w-2.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/loop-w-1.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/mode-transitions.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/par-reduction-1.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/par-reduction-2.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/parallel-dims.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/parallel-reduction.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/pr85381-3.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/private-variables.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/reduction-5.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/reduction-7.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/routine-g-1.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/routine-w-1.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/routine-wv-2.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/static-variable-1.c:
Likewise.
* testsuite/libgomp.oacc-fortran/optional-private.f90: Likewise.
* testsuite/libgomp.oacc-fortran/par-reduction-2-1.f: Likewise.
* testsuite/libgomp.oacc-fortran/par-reduction-2-2.f: Likewise.
* testsuite/libgomp.oacc-fortran/parallel-dims.f90: Likewise.
* testsuite/libgomp.oacc-fortran/parallel-reduction.f90: Likewise.
* testsuite/libgomp.oacc-fortran/pr84028.f90: Likewise.
* testsuite/libgomp.oacc-fortran/private-variables.f90: Likewise.
* testsuite/libgomp.oacc-fortran/reduction-1.f90: Likewise.
* testsuite/libgomp.oacc-fortran/reduction-5.f90: Likewise.
* testsuite/libgomp.oacc-fortran/reduction-6.f90: Likewise.
* testsuite/libgomp.oacc-fortran/routine-7.f90: Likewise.
Co-Authored-By: Nathan Sidwell <nathan@codesourcery.com>
Co-Authored-By: Tom de Vries <vries@codesourcery.com>
Co-Authored-By: Julian Brown <julian@codesourcery.com>
Co-Authored-By: Kwok Cheung Yeung <kcy@codesourcery.com>
---
gcc/c-family/c.opt | 4 +
gcc/doc/invoke.texi | 10 +-
gcc/fortran/lang.opt | 4 +
gcc/omp-offload.c | 29 ++++
gcc/testsuite/c-c++-common/goacc/acc-icf.c | 7 +
.../goacc/classify-kernels-unparallelized.c | 3 +
.../c-c++-common/goacc/classify-kernels.c | 3 +
.../c-c++-common/goacc/classify-parallel.c | 3 +
.../c-c++-common/goacc/classify-routine.c | 3 +
.../c-c++-common/goacc/classify-serial.c | 8 ++
.../c-c++-common/goacc/diag-parallelism-1.c | 124 ++++++++++++++++++
.../c-c++-common/goacc/kernels-decompose-1.c | 3 +
.../c-c++-common/goacc/kernels-decompose-2.c | 4 +
.../c-c++-common/goacc/parallel-dims-1.c | 7 +
.../c-c++-common/goacc/parallel-reduction.c | 4 +
gcc/testsuite/c-c++-common/goacc/pr70688.c | 6 +
gcc/testsuite/c-c++-common/goacc/routine-1.c | 8 ++
.../goacc/routine-level-of-parallelism-2.c | 6 +
.../c-c++-common/goacc/uninit-dim-clause.c | 6 +
.../goacc/classify-kernels-unparallelized.f95 | 3 +
.../gfortran.dg/goacc/classify-kernels.f95 | 3 +
.../gfortran.dg/goacc/classify-parallel.f95 | 3 +
.../gfortran.dg/goacc/classify-routine.f95 | 3 +
.../gfortran.dg/goacc/classify-serial.f95 | 6 +
.../gfortran.dg/goacc/kernels-decompose-1.f95 | 3 +
.../gfortran.dg/goacc/kernels-decompose-2.f95 | 4 +
.../gfortran.dg/goacc/parallel-tree.f95 | 6 +
gcc/testsuite/gfortran.dg/goacc/routine-4.f90 | 9 ++
.../goacc/routine-level-of-parallelism-1.f90 | 6 +
.../goacc/routine-module-mod-1.f90 | 4 +
.../goacc/routine-multiple-directives-1.f90 | 5 +
.../gfortran.dg/goacc/uninit-dim-clause.f95 | 6 +
.../firstprivate-1.c | 5 +
.../libgomp.oacc-c-c++-common/loop-auto-1.c | 4 +
.../libgomp.oacc-c-c++-common/loop-red-w-1.c | 4 +
.../libgomp.oacc-c-c++-common/loop-red-w-2.c | 4 +
.../libgomp.oacc-c-c++-common/loop-w-1.c | 4 +
.../mode-transitions.c | 24 ++++
.../par-reduction-1.c | 9 ++
.../par-reduction-2.c | 9 ++
.../libgomp.oacc-c-c++-common/parallel-dims.c | 10 ++
.../parallel-reduction.c | 5 +-
.../libgomp.oacc-c-c++-common/pr85381-3.c | 6 +
.../private-variables.c | 16 +++
.../libgomp.oacc-c-c++-common/reduction-5.c | 4 +-
.../libgomp.oacc-c-c++-common/reduction-7.c | 11 ++
.../libgomp.oacc-c-c++-common/routine-g-1.c | 5 +
.../libgomp.oacc-c-c++-common/routine-w-1.c | 4 +
.../libgomp.oacc-c-c++-common/routine-wv-2.c | 4 +
.../static-variable-1.c | 15 +++
.../libgomp.oacc-fortran/optional-private.f90 | 6 +
.../libgomp.oacc-fortran/par-reduction-2-1.f | 9 ++
.../libgomp.oacc-fortran/par-reduction-2-2.f | 9 ++
.../libgomp.oacc-fortran/parallel-dims.f90 | 6 +
.../parallel-reduction.f90 | 5 +
.../libgomp.oacc-fortran/pr84028.f90 | 4 +
.../private-variables.f90 | 11 ++
.../libgomp.oacc-fortran/reduction-1.f90 | 4 +
.../libgomp.oacc-fortran/reduction-5.f90 | 4 +
.../libgomp.oacc-fortran/reduction-6.f90 | 4 +
.../libgomp.oacc-fortran/routine-7.f90 | 6 +
61 files changed, 511 insertions(+), 3 deletions(-)
create mode 100644 gcc/testsuite/c-c++-common/goacc/diag-parallelism-1.c
diff --git a/gcc/c-family/c.opt b/gcc/c-family/c.opt
index 3f8b72cdc00..f1b4c3f34ec 100644
--- a/gcc/c-family/c.opt
+++ b/gcc/c-family/c.opt
@@ -1037,6 +1037,10 @@ Wold-style-definition
C ObjC Var(warn_old_style_definition) Init(-1) Warning
Warn if an old-style parameter definition is used.
+Wopenacc-parallelism
+C C++ Var(warn_openacc_parallelism) Warning
+Warn about potentially suboptimal choices related to OpenACC parallelism.
+
Wopenmp-simd
C C++ Var(warn_openmp_simd) Warning LangEnabledBy(C C++,Wall)
Warn if a simd directive is overridden by the vectorizer cost model.
diff --git a/gcc/doc/invoke.texi b/gcc/doc/invoke.texi
index e98b0962b9f..40cacc6f8e7 100644
--- a/gcc/doc/invoke.texi
+++ b/gcc/doc/invoke.texi
@@ -364,7 +364,9 @@ Objective-C and Objective-C++ Dialects}.
-Wmissing-include-dirs -Wmissing-noreturn -Wno-missing-profile @gol
-Wno-multichar -Wmultistatement-macros -Wnonnull -Wnonnull-compare @gol
-Wnormalized=@r{[}none@r{|}id@r{|}nfc@r{|}nfkc@r{]} @gol
--Wnull-dereference -Wno-odr -Wopenmp-simd @gol
+-Wnull-dereference -Wno-odr @gol
+-Wopenacc-parallelism @gol
+-Wopenmp-simd @gol
-Wno-overflow -Woverlength-strings -Wno-override-init-side-effects @gol
-Wpacked -Wno-packed-bitfield-compat -Wpacked-not-aligned -Wpadded @gol
-Wparentheses -Wno-pedantic-ms-format @gol
@@ -8749,6 +8751,12 @@ Do not warn about compile-time overflow in constant expressions.
Warn about One Definition Rule violations during link-time optimization.
Enabled by default.
+@item -Wopenacc-parallelism
+@opindex Wopenacc-parallelism
+@opindex Wno-openacc-parallelism
+@cindex OpenACC accelerator programming
+Warn about potentially suboptimal choices related to OpenACC parallelism.
+
@item -Wopenmp-simd
@opindex Wopenmp-simd
@opindex Wno-openmp-simd
diff --git a/gcc/fortran/lang.opt b/gcc/fortran/lang.opt
index 388ef8c0fdb..6db01c736be 100644
--- a/gcc/fortran/lang.opt
+++ b/gcc/fortran/lang.opt
@@ -285,6 +285,10 @@ Wuse-without-only
Fortran Var(warn_use_without_only) Warning
Warn about USE statements that have no ONLY qualifier.
+Wopenacc-parallelism
+Fortran
+; Documented in C
+
Wopenmp-simd
Fortran
; Documented in C
diff --git a/gcc/omp-offload.c b/gcc/omp-offload.c
index ba0937fba94..16124613fa7 100644
--- a/gcc/omp-offload.c
+++ b/gcc/omp-offload.c
@@ -928,6 +928,35 @@ oacc_validate_dims (tree fn, tree attrs, int *dims, int level, unsigned used)
pos = TREE_CHAIN (pos);
}
+ bool check = true;
+#ifdef ACCEL_COMPILER
+ check = false;
+#endif
+ if (check
+ && warn_openacc_parallelism
+ && !lookup_attribute ("oacc kernels", DECL_ATTRIBUTES (fn)))
+ {
+ static char const *const axes[] =
+ /* Must be kept in sync with GOMP_DIM enumeration. */
+ { "gang", "worker", "vector" };
+ for (ix = level >= 0 ? level : 0; ix != GOMP_DIM_MAX; ix++)
+ if (dims[ix] < 0)
+ ; /* Defaulting axis. */
+ else if ((used & GOMP_DIM_MASK (ix)) && dims[ix] == 1)
+ /* There is partitioned execution, but the user requested a
+ dimension size of 1. They're probably confused. */
+ warning_at (DECL_SOURCE_LOCATION (fn), OPT_Wopenacc_parallelism,
+ "region contains %s partitioned code but"
+ " is not %s partitioned", axes[ix], axes[ix]);
+ else if (!(used & GOMP_DIM_MASK (ix)) && dims[ix] != 1)
+ /* The dimension is explicitly partitioned to non-unity, but
+ no use is made within the region. */
+ warning_at (DECL_SOURCE_LOCATION (fn), OPT_Wopenacc_parallelism,
+ "region is %s partitioned but"
+ " does not contain %s partitioned code",
+ axes[ix], axes[ix]);
+ }
+
bool changed = targetm.goacc.validate_dims (fn, dims, level, used);
/* Default anything left to 1 or a partitioned default. */
diff --git a/gcc/testsuite/c-c++-common/goacc/acc-icf.c b/gcc/testsuite/c-c++-common/goacc/acc-icf.c
index 98b536c34d0..9cf119bf89c 100644
--- a/gcc/testsuite/c-c++-common/goacc/acc-icf.c
+++ b/gcc/testsuite/c-c++-common/goacc/acc-icf.c
@@ -2,7 +2,12 @@
/* { dg-additional-options "-fopenacc -O2 -fdump-ipa-icf" } */
+/* { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting
+ aspects of that functionality. */
+
#pragma acc routine gang
+/* { dg-bogus "warning: region is worker partitioned but does not contain worker partitioned code" "TODO default 'gang' 'vector'" { xfail *-*-* } .+3 }
+ TODO It's the compiler's own decision to not use 'worker' parallelism here, so it doesn't make sense to bother the user about it. */
int
routine1 (int n)
{
@@ -16,6 +21,8 @@ routine1 (int n)
}
#pragma acc routine gang
+/* { dg-bogus "warning: region is worker partitioned but does not contain worker partitioned code" "TODO default 'gang' 'vector'" { xfail *-*-* } .+3 }
+ TODO It's the compiler's own decision to not use 'worker' parallelism here, so it doesn't make sense to bother the user about it. */
int
routine2 (int n)
{
diff --git a/gcc/testsuite/c-c++-common/goacc/classify-kernels-unparallelized.c b/gcc/testsuite/c-c++-common/goacc/classify-kernels-unparallelized.c
index d4c4b2ca237..218f6248062 100644
--- a/gcc/testsuite/c-c++-common/goacc/classify-kernels-unparallelized.c
+++ b/gcc/testsuite/c-c++-common/goacc/classify-kernels-unparallelized.c
@@ -7,6 +7,9 @@
{ dg-additional-options "-fdump-tree-parloops1-all" }
{ dg-additional-options "-fdump-tree-oaccdevlow" } */
+/* { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting
+ aspects of that functionality. */
+
#define N 1024
extern unsigned int *__restrict a;
diff --git a/gcc/testsuite/c-c++-common/goacc/classify-kernels.c b/gcc/testsuite/c-c++-common/goacc/classify-kernels.c
index 16e9b9e31d1..95a150ca9ac 100644
--- a/gcc/testsuite/c-c++-common/goacc/classify-kernels.c
+++ b/gcc/testsuite/c-c++-common/goacc/classify-kernels.c
@@ -7,6 +7,9 @@
{ dg-additional-options "-fdump-tree-parloops1-all" }
{ dg-additional-options "-fdump-tree-oaccdevlow" } */
+/* { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting
+ aspects of that functionality. */
+
#define N 1024
extern unsigned int *__restrict a;
diff --git a/gcc/testsuite/c-c++-common/goacc/classify-parallel.c b/gcc/testsuite/c-c++-common/goacc/classify-parallel.c
index 933d7664386..230e70c66cd 100644
--- a/gcc/testsuite/c-c++-common/goacc/classify-parallel.c
+++ b/gcc/testsuite/c-c++-common/goacc/classify-parallel.c
@@ -6,6 +6,9 @@
{ dg-additional-options "-fdump-tree-ompexp" }
{ dg-additional-options "-fdump-tree-oaccdevlow" } */
+/* { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting
+ aspects of that functionality. */
+
#define N 1024
extern unsigned int *__restrict a;
diff --git a/gcc/testsuite/c-c++-common/goacc/classify-routine.c b/gcc/testsuite/c-c++-common/goacc/classify-routine.c
index 0b9ba6ea69f..81fe3696baa 100644
--- a/gcc/testsuite/c-c++-common/goacc/classify-routine.c
+++ b/gcc/testsuite/c-c++-common/goacc/classify-routine.c
@@ -6,6 +6,9 @@
{ dg-additional-options "-fdump-tree-ompexp" }
{ dg-additional-options "-fdump-tree-oaccdevlow" } */
+/* { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting
+ aspects of that functionality. */
+
#define N 1024
extern unsigned int *__restrict a;
diff --git a/gcc/testsuite/c-c++-common/goacc/classify-serial.c b/gcc/testsuite/c-c++-common/goacc/classify-serial.c
index 94ace1b3c20..ae052ae6a1c 100644
--- a/gcc/testsuite/c-c++-common/goacc/classify-serial.c
+++ b/gcc/testsuite/c-c++-common/goacc/classify-serial.c
@@ -6,6 +6,9 @@
{ dg-additional-options "-fdump-tree-ompexp" }
{ dg-additional-options "-fdump-tree-oaccdevlow" } */
+/* { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting
+ aspects of that functionality. */
+
#define N 1024
extern unsigned int *__restrict a;
@@ -15,6 +18,11 @@ extern unsigned int *__restrict c;
void SERIAL ()
{
#pragma acc serial loop copyin (a[0:N], b[0:N]) copyout (c[0:N]) /* { dg-message "optimized: assigned OpenACC gang vector loop parallelism" } */
+ /* { dg-bogus "warning: region contains gang partitioned code but is not gang partitioned" "TODO 'serial'" { xfail *-*-* } .-1 }
+ { dg-bogus "warning: region contains worker partitioned code but is not worker partitioned" "" { target *-*-* } .-2 }
+ { dg-bogus "warning: region contains vector partitioned code but is not vector partitioned" "TODO 'serial'" { xfail *-*-* } .-3 }
+ TODO Should we really diagnose this if the user explicitly requested 'serial'?
+ TODO Should we instead diagnose ('-Wextra' category?) that the user may enable use of parallelism if replacing 'serial' with 'parallel', if applicable? */
for (unsigned int i = 0; i < N; i++)
c[i] = a[i] + b[i];
}
diff --git a/gcc/testsuite/c-c++-common/goacc/diag-parallelism-1.c b/gcc/testsuite/c-c++-common/goacc/diag-parallelism-1.c
new file mode 100644
index 00000000000..6ee7bd078dd
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/diag-parallelism-1.c
@@ -0,0 +1,124 @@
+/* Diagnostics about potentially suboptimal choices related to OpenACC
+ parallelism.
+
+ { dg-additional-options "-Wopenacc-parallelism" }
+*/
+
+
+//TODO 'kernels'
+
+//TODO 'serial'
+
+//TODO 'routine'
+
+//TODO Fortran
+
+
+static void f1 ()
+{
+ int ary[10];
+
+
+#pragma acc parallel num_gangs (1)
+ /* { dg-warning "region contains gang partitioned code but is not gang partitioned" "" { target *-*-* } .-1 } */
+ {
+ #pragma acc loop gang
+ for (int i = 0; i < 10; i++)
+ ary[i] = i;
+ }
+
+#pragma acc parallel num_workers (1)
+ /* { dg-warning "region contains worker partitioned code but is not worker partitioned" "" { target *-*-* } .-1 } */
+ {
+ #pragma acc loop worker
+ for (int i = 0; i < 10; i++)
+ ary[i] = i;
+ }
+
+#pragma acc parallel vector_length (1)
+ /* { dg-warning "region contains vector partitioned code but is not vector partitioned" "" { target *-*-* } .-1 } */
+ {
+ #pragma acc loop vector
+ for (int i = 0; i < 10; i++)
+ ary[i] = i;
+ }
+}
+
+
+static void f2 ()
+{
+ int ary[10];
+
+
+#pragma acc parallel num_gangs (8)
+ /* { dg-warning "region is gang partitioned but does not contain gang partitioned code" "" { target *-*-* } .-1 } */
+ {
+ #pragma acc loop worker
+ for (int i = 0; i < 10; i++)
+ ary[i] = i;
+ }
+
+#pragma acc parallel num_gangs (8)
+ /* { dg-warning "region is gang partitioned but does not contain gang partitioned code" "" { target *-*-* } .-1 } */
+ {
+ #pragma acc loop vector
+ for (int i = 0; i < 10; i++)
+ ary[i] = i;
+ }
+
+#pragma acc parallel num_gangs (8)
+ /* { dg-warning "region is gang partitioned but does not contain gang partitioned code" "" { target *-*-* } .-1 } */
+ {
+ #pragma acc loop worker vector
+ for (int i = 0; i < 10; i++)
+ ary[i] = i;
+ }
+
+#pragma acc parallel num_workers (8)
+ /* { dg-warning "region is worker partitioned but does not contain worker partitioned code" "" { target *-*-* } .-1 } */
+ {
+ #pragma acc loop gang
+ for (int i = 0; i < 10; i++)
+ ary[i] = i;
+ }
+
+#pragma acc parallel num_workers (8)
+ /* { dg-warning "region is worker partitioned but does not contain worker partitioned code" "" { target *-*-* } .-1 } */
+ {
+ #pragma acc loop vector
+ for (int i = 0; i < 10; i++)
+ ary[i] = i;
+ }
+
+#pragma acc parallel num_workers (8)
+ /* { dg-warning "region is worker partitioned but does not contain worker partitioned code" "" { target *-*-* } .-1 } */
+ {
+ #pragma acc loop gang vector
+ for (int i = 0; i < 10; i++)
+ ary[i] = i;
+ }
+
+#pragma acc parallel vector_length (8)
+ /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .-1 } */
+ {
+ #pragma acc loop gang
+ for (int i = 0; i < 10; i++)
+ ary[i] = i;
+ }
+
+#pragma acc parallel vector_length (8)
+ /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .-1 } */
+ {
+ #pragma acc loop worker
+ for (int i = 0; i < 10; i++)
+ ary[i] = i;
+ }
+
+#pragma acc parallel vector_length (8)
+ /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .-1 } */
+ {
+ #pragma acc loop gang worker
+ for (int i = 0; i < 10; i++)
+ ary[i] = i;
+ }
+}
diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-1.c b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-1.c
index 87219c88fac..f549cbadfa7 100644
--- a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-1.c
+++ b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-1.c
@@ -5,6 +5,9 @@
/* { dg-additional-options "--param=openacc-kernels=decompose" }
{ dg-additional-options "-fdump-tree-omp_oacc_kernels_decompose" } */
+/* { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting
+ aspects of that functionality. */
+
/* See also '../../gfortran.dg/goacc/kernels-decompose-1.f95'. */
/* It's only with Tcl 8.5 (released in 2007) that "the variable 'varName'
diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-2.c b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-2.c
index 3781e75d0f2..cdf85d4bafa 100644
--- a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-2.c
+++ b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-2.c
@@ -4,6 +4,9 @@
/* { dg-additional-options "--param=openacc-kernels=decompose" }
/* { dg-additional-options "-O2" } for 'parloops'. */
+/* { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting
+ aspects of that functionality. */
+
/* See also '../../gfortran.dg/goacc/kernels-decompose-2.f95'. */
/* It's only with Tcl 8.5 (released in 2007) that "the variable 'varName'
@@ -116,6 +119,7 @@ main ()
}
#pragma acc kernels
+ /* { dg-bogus "warning: region contains gang partitioned code but is not gang partitioned" "TODO 'kernels'" { xfail *-*-* } .-1 } */
{
y = f_g (a[5]); /* { dg-line l_part[incr c_part] } */
/*TODO If such a construct is placed in its own part (like it is, here), can't this actually use gang paralelism, instead of "gang-single"?
diff --git a/gcc/testsuite/c-c++-common/goacc/parallel-dims-1.c b/gcc/testsuite/c-c++-common/goacc/parallel-dims-1.c
index 57f682f7a0d..2a8d35d493d 100644
--- a/gcc/testsuite/c-c++-common/goacc/parallel-dims-1.c
+++ b/gcc/testsuite/c-c++-common/goacc/parallel-dims-1.c
@@ -1,11 +1,18 @@
/* Valid use of OpenACC parallelism dimensions clauses: num_gangs, num_workers,
vector_length. */
+/* { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting
+ aspects of that functionality. */
+
void f(int i)
{
#pragma acc kernels num_gangs(i) num_workers(i) vector_length(i)
;
#pragma acc parallel num_gangs(i) num_workers(i) vector_length(i)
+ /* { dg-bogus "warning: region is gang partitioned but does not contain gang partitioned code" "TODO runtime" { xfail *-*-* } .-1 }
+ { dg-bogus "warning: region is worker partitioned but does not contain worker partitioned code" "TODO runtime" { xfail *-*-* } .-2 }
+ { dg-bogus "warning: region is vector partitioned but does not contain vector partitioned code" "TODO runtime" { xfail *-*-* } .-3 }
+ TODO 'region is [...] partitioned' isn't correct for 'i == 1'. */
;
}
diff --git a/gcc/testsuite/c-c++-common/goacc/parallel-reduction.c b/gcc/testsuite/c-c++-common/goacc/parallel-reduction.c
index d7cc9470127..c5c0edc98ee 100644
--- a/gcc/testsuite/c-c++-common/goacc/parallel-reduction.c
+++ b/gcc/testsuite/c-c++-common/goacc/parallel-reduction.c
@@ -1,3 +1,6 @@
+/* { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting
+ aspects of that functionality. */
+
int
main ()
{
@@ -7,6 +10,7 @@ main ()
#pragma acc data copy (dummy)
{
#pragma acc parallel num_gangs (10) copy (sum) reduction (+:sum)
+ /* { dg-bogus "warning: region is gang partitioned but does not contain gang partitioned code" "TODO 'reduction'" { xfail *-*-* } .-1 } */
{
int v = 5;
sum += 10 + v;
diff --git a/gcc/testsuite/c-c++-common/goacc/pr70688.c b/gcc/testsuite/c-c++-common/goacc/pr70688.c
index 5a236654069..e96a8530772 100644
--- a/gcc/testsuite/c-c++-common/goacc/pr70688.c
+++ b/gcc/testsuite/c-c++-common/goacc/pr70688.c
@@ -1,3 +1,6 @@
+/* { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting
+ aspects of that functionality. */
+
const int n = 100;
int
@@ -22,6 +25,7 @@ parallel_reduction ()
#pragma acc data copy (dummy)
{
#pragma acc parallel num_gangs (10) copy (sum) reduction (+:sum)
+ /* { dg-bogus "warning: region is gang partitioned but does not contain gang partitioned code" "TODO 'reduction'" { xfail *-*-* } .-1 } */
{
int v = 5;
sum += 10 + v;
@@ -37,10 +41,12 @@ main ()
int i, s = 0;
#pragma acc parallel num_gangs (10) copy (s) reduction (+:s)
+ /* { dg-bogus "warning: region is gang partitioned but does not contain gang partitioned code" "TODO 'reduction'" { xfail *-*-* } .-1 } */
for (i = 0; i < n; i++)
s += i+1;
#pragma acc parallel num_gangs (10) reduction (+:s) copy (s)
+ /* { dg-bogus "warning: region is gang partitioned but does not contain gang partitioned code" "TODO 'reduction'" { xfail *-*-* } .-1 } */
for (i = 0; i < n; i++)
s += i+1;
diff --git a/gcc/testsuite/c-c++-common/goacc/routine-1.c b/gcc/testsuite/c-c++-common/goacc/routine-1.c
index a75692246b6..051f7937dee 100644
--- a/gcc/testsuite/c-c++-common/goacc/routine-1.c
+++ b/gcc/testsuite/c-c++-common/goacc/routine-1.c
@@ -1,15 +1,23 @@
+/* { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting
+ aspects of that functionality. */
#pragma acc routine gang
+/* { dg-warning "region is gang partitioned but does not contain gang partitioned code" "" { target *-*-* } .+3 }
+ { dg-warning "region is worker partitioned but does not contain worker partitioned code" "" { target *-*-* } .+2 }
+ { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .+1 } */
void gang (void)
{
}
#pragma acc routine worker
+/* { dg-warning "region is worker partitioned but does not contain worker partitioned code" "" { target *-*-* } .+2 }
+ { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .+1 } */
void worker (void)
{
}
#pragma acc routine vector
+/* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .+1 } */
void vector (void)
{
}
diff --git a/gcc/testsuite/c-c++-common/goacc/routine-level-of-parallelism-2.c b/gcc/testsuite/c-c++-common/goacc/routine-level-of-parallelism-2.c
index a066f2b9c2b..33678fe8d0a 100644
--- a/gcc/testsuite/c-c++-common/goacc/routine-level-of-parallelism-2.c
+++ b/gcc/testsuite/c-c++-common/goacc/routine-level-of-parallelism-2.c
@@ -2,7 +2,13 @@
with the OpenACC 'routine' directive. The Fortran counterpart is
'../../gfortran.dg/goacc/routine-level-of-parallelism-1.f90'. */
+/* { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting
+ aspects of that functionality. */
+
#pragma acc routine gang
+/* { dg-warning "region is gang partitioned but does not contain gang partitioned code" "" { target *-*-* } .+3 }
+ { dg-warning "region is worker partitioned but does not contain worker partitioned code" "" { target *-*-* } .+2 }
+ { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .+1 } */
void g_1 (void)
{
}
diff --git a/gcc/testsuite/c-c++-common/goacc/uninit-dim-clause.c b/gcc/testsuite/c-c++-common/goacc/uninit-dim-clause.c
index 827dac71519..d9b143bb108 100644
--- a/gcc/testsuite/c-c++-common/goacc/uninit-dim-clause.c
+++ b/gcc/testsuite/c-c++-common/goacc/uninit-dim-clause.c
@@ -1,16 +1,22 @@
/* { dg-additional-options "-Wuninitialized" } */
+/* { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting
+ aspects of that functionality. */
+
void acc_parallel()
{
int i, j, k;
#pragma acc parallel num_gangs(i) /* { dg-warning "is used uninitialized" } */
+ /* { dg-warning "region is gang partitioned but does not contain gang partitioned code" "" { target *-*-* } .-1 } */
;
#pragma acc parallel num_workers(j) /* { dg-warning "is used uninitialized" } */
+ /* { dg-warning "region is worker partitioned but does not contain worker partitioned code" "" { target *-*-* } .-1 } */
;
#pragma acc parallel vector_length(k) /* { dg-warning "is used uninitialized" } */
+ /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .-1 } */
;
}
diff --git a/gcc/testsuite/gfortran.dg/goacc/classify-kernels-unparallelized.f95 b/gcc/testsuite/gfortran.dg/goacc/classify-kernels-unparallelized.f95
index 6cca3d6eefb..cb5251a2aeb 100644
--- a/gcc/testsuite/gfortran.dg/goacc/classify-kernels-unparallelized.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/classify-kernels-unparallelized.f95
@@ -7,6 +7,9 @@
! { dg-additional-options "-fdump-tree-parloops1-all" }
! { dg-additional-options "-fdump-tree-oaccdevlow" }
+! { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting
+! aspects of that functionality.
+
program main
implicit none
integer, parameter :: n = 1024
diff --git a/gcc/testsuite/gfortran.dg/goacc/classify-kernels.f95 b/gcc/testsuite/gfortran.dg/goacc/classify-kernels.f95
index 715a983bb26..07aaf065e1d 100644
--- a/gcc/testsuite/gfortran.dg/goacc/classify-kernels.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/classify-kernels.f95
@@ -7,6 +7,9 @@
! { dg-additional-options "-fdump-tree-parloops1-all" }
! { dg-additional-options "-fdump-tree-oaccdevlow" }
+! { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting
+! aspects of that functionality.
+
program main
implicit none
integer, parameter :: n = 1024
diff --git a/gcc/testsuite/gfortran.dg/goacc/classify-parallel.f95 b/gcc/testsuite/gfortran.dg/goacc/classify-parallel.f95
index 01f06bbcc27..a41e0e68b38 100644
--- a/gcc/testsuite/gfortran.dg/goacc/classify-parallel.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/classify-parallel.f95
@@ -6,6 +6,9 @@
! { dg-additional-options "-fdump-tree-ompexp" }
! { dg-additional-options "-fdump-tree-oaccdevlow" }
+! { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting
+! aspects of that functionality.
+
program main
implicit none
integer, parameter :: n = 1024
diff --git a/gcc/testsuite/gfortran.dg/goacc/classify-routine.f95 b/gcc/testsuite/gfortran.dg/goacc/classify-routine.f95
index 401d5270391..52cc870dfba 100644
--- a/gcc/testsuite/gfortran.dg/goacc/classify-routine.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/classify-routine.f95
@@ -6,6 +6,9 @@
! { dg-additional-options "-fdump-tree-ompexp" }
! { dg-additional-options "-fdump-tree-oaccdevlow" }
+! { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting
+! aspects of that functionality.
+
subroutine ROUTINE
!$acc routine worker
integer, parameter :: n = 1024
diff --git a/gcc/testsuite/gfortran.dg/goacc/classify-serial.f95 b/gcc/testsuite/gfortran.dg/goacc/classify-serial.f95
index 51061afd2c6..6d402e683de 100644
--- a/gcc/testsuite/gfortran.dg/goacc/classify-serial.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/classify-serial.f95
@@ -6,6 +6,9 @@
! { dg-additional-options "-fdump-tree-ompexp" }
! { dg-additional-options "-fdump-tree-oaccdevlow" }
+! { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting
+! aspects of that functionality.
+
program main
implicit none
integer, parameter :: n = 1024
@@ -15,6 +18,9 @@ program main
call setup(a, b)
!$acc serial loop copyin (a(0:n-1), b(0:n-1)) copyout (c(0:n-1)) ! { dg-message "optimized: assigned OpenACC gang vector loop parallelism" }
+ ! { dg-bogus "warning: region contains gang partitioned code but is not gang partitioned" "TODO 'serial'" { xfail *-*-* } .-1 }
+ ! { dg-bogus "warning: region contains worker partitioned code but is not worker partitioned" "" { target *-*-* } .-2 }
+ ! { dg-bogus "warning: region contains vector partitioned code but is not vector partitioned" "TODO 'serial'" { xfail *-*-* } .-3 }
do i = 0, n - 1
c(i) = a(i) + b(i)
end do
diff --git a/gcc/testsuite/gfortran.dg/goacc/kernels-decompose-1.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-decompose-1.f95
index e2523504ef5..ddaf7f8e43d 100644
--- a/gcc/testsuite/gfortran.dg/goacc/kernels-decompose-1.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/kernels-decompose-1.f95
@@ -5,6 +5,9 @@
! { dg-additional-options "--param=openacc-kernels=decompose" }
! { dg-additional-options "-fdump-tree-omp_oacc_kernels_decompose" }
+! { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting
+! aspects of that functionality.
+
! See also '../../c-c++-common/goacc/kernels-decompose-1.c'.
! It's only with Tcl 8.5 (released in 2007) that "the variable 'varName'
diff --git a/gcc/testsuite/gfortran.dg/goacc/kernels-decompose-2.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-decompose-2.f95
index cc12b77817b..a3cec8325c7 100644
--- a/gcc/testsuite/gfortran.dg/goacc/kernels-decompose-2.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/kernels-decompose-2.f95
@@ -4,6 +4,9 @@
! { dg-additional-options "--param=openacc-kernels=decompose" }
! { dg-additional-options "-O2" } for 'parloops'.
+! { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting
+! aspects of that functionality.
+
! See also '../../c-c++-common/goacc/kernels-decompose-2.c'.
! It's only with Tcl 8.5 (released in 2007) that "the variable 'varName'
@@ -119,6 +122,7 @@ program main
!$acc end kernels
!$acc kernels
+ ! { dg-bogus "warning: region contains gang partitioned code but is not gang partitioned" "TODO 'kernels'" { xfail *-*-* } .-1 }
y = f_g (a(5)) ! { dg-line l_part[incr c_part] }
!TODO If such a construct is placed in its own part (like it is, here), can't this actually use gang paralelism, instead of "gang-single"?
! { dg-message "note: beginning 'gang-single' part in OpenACC 'kernels' region" "" { target *-*-* } l_part$c_part }
diff --git a/gcc/testsuite/gfortran.dg/goacc/parallel-tree.f95 b/gcc/testsuite/gfortran.dg/goacc/parallel-tree.f95
index e33653bdd78..c51cb2a0504 100644
--- a/gcc/testsuite/gfortran.dg/goacc/parallel-tree.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/parallel-tree.f95
@@ -2,6 +2,9 @@
! test for tree-dump-original and spaces-commas
+! { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting
+! aspects of that functionality.
+
program test
implicit none
integer :: q, i, j, k, m, n, o, p, r, s, t, u, v, w
@@ -12,6 +15,9 @@ program test
!$acc no_create(n) &
!$acc present(o), pcopy(p), pcopyin(r), pcopyout(s), pcreate(t) &
!$acc deviceptr(u), private(v), firstprivate(w)
+ ! { dg-warning "region is gang partitioned but does not contain gang partitioned code" "" { target *-*-* } .-1 }
+ ! { dg-warning "region is worker partitioned but does not contain worker partitioned code" "" { target *-*-* } .-2 }
+ ! { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .-3 }
!$acc end parallel
end program test
diff --git a/gcc/testsuite/gfortran.dg/goacc/routine-4.f90 b/gcc/testsuite/gfortran.dg/goacc/routine-4.f90
index 6714c7b8229..53b1fbe5039 100644
--- a/gcc/testsuite/gfortran.dg/goacc/routine-4.f90
+++ b/gcc/testsuite/gfortran.dg/goacc/routine-4.f90
@@ -1,5 +1,8 @@
! Test invalid calls to routines.
+! { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting
+! aspects of that functionality.
+
module param
integer, parameter :: N = 32
end module param
@@ -120,6 +123,9 @@ contains
subroutine gang (a) ! { dg-message "declared here" 3 }
!$acc routine gang
+ ! { dg-warning "region is gang partitioned but does not contain gang partitioned code" "" { target *-*-* } .-2 }
+ ! { dg-warning "region is worker partitioned but does not contain worker partitioned code" "" { target *-*-* } .-3 }
+ ! { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .-4 }
integer, intent (inout) :: a(N)
integer :: i
@@ -130,6 +136,8 @@ contains
subroutine worker (a) ! { dg-message "declared here" 2 }
!$acc routine worker
+ ! { dg-warning "region is worker partitioned but does not contain worker partitioned code" "" { target *-*-* } .-2 }
+ ! { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .-3 }
integer, intent (inout) :: a(N)
integer :: i
@@ -140,6 +148,7 @@ contains
subroutine vector (a) ! { dg-message "declared here" }
!$acc routine vector
+ ! { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .-2 }
integer, intent (inout) :: a(N)
integer :: i
diff --git a/gcc/testsuite/gfortran.dg/goacc/routine-level-of-parallelism-1.f90 b/gcc/testsuite/gfortran.dg/goacc/routine-level-of-parallelism-1.f90
index 83b8c24b41d..99929f136f6 100644
--- a/gcc/testsuite/gfortran.dg/goacc/routine-level-of-parallelism-1.f90
+++ b/gcc/testsuite/gfortran.dg/goacc/routine-level-of-parallelism-1.f90
@@ -2,8 +2,14 @@
! with the OpenACC routine directive. The C/C++ counterpart is
! '../../c-c++-common/goacc/routine-level-of-parallelism-2.c'.
+! { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting
+! aspects of that functionality.
+
subroutine g_1
!$acc routine gang
+ ! { dg-warning "region is gang partitioned but does not contain gang partitioned code" "" { target *-*-* } .-2 }
+ ! { dg-warning "region is worker partitioned but does not contain worker partitioned code" "" { target *-*-* } .-3 }
+ ! { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .-4 }
end subroutine g_1
subroutine s_1_2a
diff --git a/gcc/testsuite/gfortran.dg/goacc/routine-module-mod-1.f90 b/gcc/testsuite/gfortran.dg/goacc/routine-module-mod-1.f90
index 23c673fe3bd..22106508141 100644
--- a/gcc/testsuite/gfortran.dg/goacc/routine-module-mod-1.f90
+++ b/gcc/testsuite/gfortran.dg/goacc/routine-module-mod-1.f90
@@ -2,6 +2,9 @@
! { dg-additional-options "-fopt-info-optimized-omp" }
+! { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting
+! aspects of that functionality.
+
module routine_module_mod_1
contains
subroutine s_1
@@ -53,6 +56,7 @@ contains
subroutine g_1
implicit none
!$acc routine gang
+ ! { dg-bogus "warning: region is worker partitioned but does not contain worker partitioned code" "TODO default 'gang' 'vector'" { xfail *-*-* } .-3 }
integer :: i
diff --git a/gcc/testsuite/gfortran.dg/goacc/routine-multiple-directives-1.f90 b/gcc/testsuite/gfortran.dg/goacc/routine-multiple-directives-1.f90
index 6e12ee92155..622a9d9ccce 100644
--- a/gcc/testsuite/gfortran.dg/goacc/routine-multiple-directives-1.f90
+++ b/gcc/testsuite/gfortran.dg/goacc/routine-multiple-directives-1.f90
@@ -1,5 +1,8 @@
! Check for valid cases of multiple OpenACC 'routine' directives.
+! { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting
+! aspects of that functionality.
+
SUBROUTINE s_1
!$ACC ROUTINE(s_1)
!$ACC ROUTINE(s_1) SEQ
@@ -17,12 +20,14 @@
!$ACC ROUTINE VECTOR
!$ACC ROUTINE(v_1) VECTOR
!$ACC ROUTINE VECTOR
+! { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .-5 }
END SUBROUTINE v_1
SUBROUTINE v_2
!$ACC ROUTINE(v_2) VECTOR
!$ACC ROUTINE VECTOR
!$ACC ROUTINE(v_2) VECTOR
+! { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .-4 }
END SUBROUTINE v_2
SUBROUTINE sub_1
diff --git a/gcc/testsuite/gfortran.dg/goacc/uninit-dim-clause.f95 b/gcc/testsuite/gfortran.dg/goacc/uninit-dim-clause.f95
index c77d47a3900..59e923f90fb 100644
--- a/gcc/testsuite/gfortran.dg/goacc/uninit-dim-clause.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/uninit-dim-clause.f95
@@ -1,16 +1,22 @@
! { dg-additional-options "-Wuninitialized" }
+! { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting
+! aspects of that functionality.
+
subroutine acc_parallel
implicit none
integer :: i, j, k
!$acc parallel num_gangs(i) ! { dg-warning "is used uninitialized" }
+ ! { dg-warning "region is gang partitioned but does not contain gang partitioned code" "" { target *-*-* } .-1 }
!$acc end parallel
!$acc parallel num_workers(j) ! { dg-warning "is used uninitialized" }
+ ! { dg-warning "region is worker partitioned but does not contain worker partitioned code" "" { target *-*-* } .-1 }
!$acc end parallel
!$acc parallel vector_length(k) ! { dg-warning "is used uninitialized" }
+ ! { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .-1 }
!$acc end parallel
end subroutine acc_parallel
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/firstprivate-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/firstprivate-1.c
index 689a443ca43..0990e3db224 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/firstprivate-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/firstprivate-1.c
@@ -1,3 +1,6 @@
+/* { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting
+ aspects of that functionality. */
+
#include <openacc.h>
@@ -117,6 +120,8 @@ void t4 ()
arr[i] = 3;
#pragma acc parallel firstprivate(x) copy(arr) num_gangs(32) num_workers(8) vector_length(32)
+ /* { dg-warning "region is worker partitioned but does not contain worker partitioned code" "" { target *-*-* } .-1 } */
+ /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .-2 } */
{
#pragma acc loop gang
for (i = 0; i < 32; i++)
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-auto-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-auto-1.c
index 0273c2bddd7..c13cab780cb 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-auto-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-auto-1.c
@@ -3,6 +3,9 @@
/* { dg-additional-options "-fopenacc-dim=32" } */
+/* { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting
+ aspects of that functionality. */
+
#include <stdio.h>
#include <openacc.h>
#include <gomp-constants.h>
@@ -151,6 +154,7 @@ int gang_1 (int *ary, int size)
clear (ary, size);
#pragma acc parallel num_gangs (32) num_workers (32) vector_length(32) copy(ary[0:size]) firstprivate (size)
+ /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .-1 } */
{
#pragma acc loop auto
for (int jx = 0; jx < size / 64; jx++)
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-w-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-w-1.c
index 648f89e1668..2f749e04ae0 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-w-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-w-1.c
@@ -1,3 +1,6 @@
+/* { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting
+ aspects of that functionality. */
+
#include <stdio.h>
#include <openacc.h>
#include <gomp-constants.h>
@@ -12,6 +15,7 @@ int main ()
#pragma acc parallel num_workers(32) vector_length(32) copy(ondev) \
copyout(workersize)
+ /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .-2 } */
{
#pragma acc loop worker reduction(+:t)
for (unsigned ix = 0; ix < N; ix++)
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-w-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-w-2.c
index f9fcf3703af..9727e22d3c2 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-w-2.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-w-2.c
@@ -1,3 +1,6 @@
+/* { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting
+ aspects of that functionality. */
+
#include <stdio.h>
#include <openacc.h>
#include <gomp-constants.h>
@@ -12,6 +15,7 @@ int main ()
#pragma acc parallel num_workers(32) vector_length(32) copy(q) copy(ondev) \
copyout(workersize)
+ /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .-2 } */
{
int t = q;
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-w-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-w-1.c
index 5fe486f50a1..d639e14a67c 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-w-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-w-1.c
@@ -1,3 +1,6 @@
+/* { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting
+ aspects of that functionality. */
+
#include <stdio.h>
#include <openacc.h>
#include <gomp-constants.h>
@@ -16,6 +19,7 @@ int main ()
#pragma acc parallel num_workers(32) vector_length(32) copy(ary) copy(ondev) \
copyout(workersize)
+ /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } .-2 } */
{
#pragma acc loop worker
for (unsigned ix = 0; ix < N; ix++)
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/mode-transitions.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/mode-transitions.c
index 4474c127992..6c989abedf5 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/mode-transitions.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/mode-transitions.c
@@ -1,3 +1,6 @@
+/* { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting
+ aspects of that functionality. */
+
/* Miscellaneous test cases for gang/worker/vector mode transitions. */
#include <assert.h>
@@ -287,6 +290,7 @@ void t7()
int n = 0;
#pragma acc parallel copy(n) \
num_gangs(1) num_workers(1) vector_length(32)
+ /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .-2 } */
{
n++;
}
@@ -310,6 +314,7 @@ void t8()
#pragma acc parallel copy(arr) \
num_gangs(gangs) num_workers(1) vector_length(32)
+ /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .-2 } */
{
int j;
#pragma acc loop gang
@@ -339,6 +344,7 @@ void t9()
#pragma acc parallel copy(arr) \
num_gangs(gangs) num_workers(1) vector_length(32)
+ /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .-2 } */
{
int j;
#pragma acc loop gang
@@ -371,6 +377,7 @@ void t10()
#pragma acc parallel copy(arr) \
num_gangs(gangs) num_workers(1) vector_length(32)
+ /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .-2 } */
{
int j;
#pragma acc loop gang
@@ -404,6 +411,7 @@ void t11()
#pragma acc parallel copy(arr) \
num_gangs(1024) num_workers(1) vector_length(32)
+ /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .-2 } */
{
int j;
@@ -442,6 +450,7 @@ void t12()
#pragma acc parallel copyout(fizz, buzz, fizzbuzz) \
num_gangs(NUM_GANGS) num_workers(1) vector_length(32)
+ /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .-2 } */
{
int j;
@@ -488,6 +497,7 @@ void t13()
#pragma acc parallel copy(arr) \
num_gangs(8) num_workers(8) vector_length(32)
+ /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .-2 } */
{
int j;
#pragma acc loop gang
@@ -613,6 +623,7 @@ void t16()
#pragma acc parallel copy(n, arr) \
num_gangs(8) num_workers(16) vector_length(32)
+ /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .-2 } */
{
int j;
#pragma acc loop gang
@@ -665,6 +676,7 @@ void t17()
#pragma acc parallel copyin(arr_a) copyout(arr_b) \
num_gangs(num_gangs) num_workers(num_workers) vector_length(32)
+ /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .-2 } */
{
int j;
#pragma acc loop gang
@@ -882,6 +894,8 @@ void t21()
#pragma acc parallel copy(arr) \
num_gangs(8) num_workers(8) vector_length(32)
+ /* { dg-warning "region is worker partitioned but does not contain worker partitioned code" "" { target *-*-* } .-2 } */
+ /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .-3 } */
{
int j;
#pragma acc loop gang
@@ -905,6 +919,8 @@ void t22()
#pragma acc parallel copy(arr) \
num_gangs(8) num_workers(8) vector_length(32)
+ /* { dg-warning "region is worker partitioned but does not contain worker partitioned code" "" { target *-*-* } .-2 } */
+ /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .-3 } */
{
int j;
#pragma acc loop gang
@@ -931,6 +947,8 @@ void t23()
#pragma acc parallel copy(arr) \
num_gangs(8) num_workers(8) vector_length(32)
+ /* { dg-warning "region is worker partitioned but does not contain worker partitioned code" "" { target *-*-* } .-2 } */
+ /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .-3 } */
{
int j;
#pragma acc loop gang
@@ -957,6 +975,8 @@ void t24()
#pragma acc parallel copy(arr) \
num_gangs(8) num_workers(8) vector_length(32)
+ /* { dg-warning "region is worker partitioned but does not contain worker partitioned code" "" { target *-*-* } .-2 } */
+ /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .-3 } */
{
int j;
#pragma acc loop gang
@@ -988,6 +1008,7 @@ void t25()
#pragma acc parallel copy(arr) \
num_gangs(8) num_workers(8) vector_length(32)
+ /* { dg-warning "region is worker partitioned but does not contain worker partitioned code" "" { target *-*-* } .-2 } */
{
int j;
#pragma acc loop gang
@@ -1020,6 +1041,7 @@ void t26()
#pragma acc parallel copy(arr) \
num_gangs(8) num_workers(8) vector_length(32)
+ /* { dg-warning "region is worker partitioned but does not contain worker partitioned code" "" { target *-*-* } .-2 } */
{
int j;
#pragma acc loop gang
@@ -1070,6 +1092,8 @@ void t27()
#pragma acc parallel copy(n, arr) copyout(ondev) \
num_gangs(ACTUAL_GANGS) num_workers(8) vector_length(32)
+ /* { dg-bogus "warning: region is gang partitioned but does not contain gang partitioned code" "TODO 'atomic'" { xfail *-*-* } .-2 } */
+ /* { dg-warning "region is worker partitioned but does not contain worker partitioned code" "" { target *-*-* } .-3 } */
{
int j;
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/par-reduction-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/par-reduction-1.c
index 0c1ecc68e9d..8157cff109a 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/par-reduction-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/par-reduction-1.c
@@ -1,3 +1,6 @@
+/* { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting
+ aspects of that functionality. */
+
#include <assert.h>
/* Test of reduction on parallel directive. */
@@ -16,6 +19,9 @@ main (int argc, char *argv[])
#endif
#pragma acc parallel num_gangs(GANGS) num_workers(32) vector_length(32) \
reduction(+:res1) copy(res2, res1)
+ /* { dg-bogus "warning: region is gang partitioned but does not contain gang partitioned code" "TODO 'reduction', 'atomic'" { xfail { ! openacc_host_selected } } .-2 } */
+ /* { dg-warning "region is worker partitioned but does not contain worker partitioned code" "" { target *-*-* } .-3 } */
+ /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .-4 } */
{
res1 += 5;
@@ -37,6 +43,9 @@ main (int argc, char *argv[])
#endif
#pragma acc parallel num_gangs(GANGS) num_workers(32) vector_length(32) \
reduction(*:res1) copy(res1, res2)
+ /* { dg-bogus "warning: region is gang partitioned but does not contain gang partitioned code" "TODO 'reduction', 'atomic'" { xfail { ! openacc_host_selected } } .-2 } */
+ /* { dg-warning "region is worker partitioned but does not contain worker partitioned code" "" { target *-*-* } .-3 } */
+ /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .-4 } */
{
res1 *= 5;
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/par-reduction-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/par-reduction-2.c
index e3ed281610d..8b7900e66ff 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/par-reduction-2.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/par-reduction-2.c
@@ -1,6 +1,9 @@
/* Test of reduction on parallel directive (with async). */
/* See also Fortran variants in "../libgomp.oacc-fortran/par-reduction-2*". */
+/* { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting
+ aspects of that functionality. */
+
#include <assert.h>
#include <openacc.h>
@@ -16,6 +19,9 @@ main (int argc, char *argv[])
#endif
#pragma acc parallel num_gangs(GANGS) num_workers(32) vector_length(32) \
reduction(+:res1) copy(res1, res2) async(1)
+ /* { dg-bogus "warning: region is gang partitioned but does not contain gang partitioned code" "TODO 'reduction', 'atomic'" { xfail { ! openacc_host_selected } } .-2 } */
+ /* { dg-warning "region is worker partitioned but does not contain worker partitioned code" "" { target *-*-* } .-3 } */
+ /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .-4 } */
{
res1 += 5;
@@ -39,6 +45,9 @@ main (int argc, char *argv[])
#endif
#pragma acc parallel num_gangs(GANGS) num_workers(32) vector_length(32) \
reduction(*:res1) copy(res1, res2) async(1)
+ /* { dg-bogus "warning: region is gang partitioned but does not contain gang partitioned code" "TODO 'reduction', 'atomic'" { xfail { ! openacc_host_selected } } .-2 } */
+ /* { dg-warning "region is worker partitioned but does not contain worker partitioned code" "" { target *-*-* } .-3 } */
+ /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .-4 } */
{
res1 *= 5;
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/parallel-dims.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/parallel-dims.c
index cc4c738c1db..ef4917aafff 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/parallel-dims.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/parallel-dims.c
@@ -1,6 +1,9 @@
/* OpenACC parallelism dimensions clauses: num_gangs, num_workers,
vector_length. */
+/* { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting
+ aspects of that functionality. */
+
/* See also '../libgomp.oacc-fortran/parallel-dims.f90'. */
#include <limits.h>
@@ -105,6 +108,7 @@ int main ()
gangs_max = workers_max = vectors_max = INT_MIN;
#pragma acc parallel copy (gangs_actual) \
num_gangs (GANGS) /* { dg-warning "'num_gangs' value must be positive" "" { target c++ } } */
+ /* { dg-warning "region contains gang partitioned code but is not gang partitioned" "" { target *-*-* } .-2 } */
{
/* We're actually executing with num_gangs (1). */
gangs_actual = 1;
@@ -134,6 +138,7 @@ int main ()
gangs_max = workers_max = vectors_max = INT_MIN;
#pragma acc parallel copy (workers_actual) \
num_workers (WORKERS) /* { dg-warning "'num_workers' value must be positive" "" { target c++ } } */
+ /* { dg-warning "region contains worker partitioned code but is not worker partitioned" "" { target *-*-* } .-2 } */
{
/* We're actually executing with num_workers (1). */
workers_actual = 1;
@@ -163,6 +168,7 @@ int main ()
gangs_max = workers_max = vectors_max = INT_MIN;
#pragma acc parallel copy (vectors_actual) /* { dg-warning "using vector_length \\(32\\), ignoring 1" "" { target openacc_nvidia_accel_selected } } */ \
vector_length (VECTORS) /* { dg-warning "'vector_length' value must be positive" "" { target c++ } } */
+ /* { dg-warning "region contains vector partitioned code but is not vector partitioned" "" { target *-*-* } .-2 } */
{
/* We're actually executing with vector_length (1), just the GCC nvptx
back end enforces vector_length (32). */
@@ -208,6 +214,7 @@ int main ()
#pragma acc parallel copy (gangs_actual) \
reduction (min: gangs_min, workers_min, vectors_min) reduction (max: gangs_max, workers_max, vectors_max) \
num_gangs (gangs)
+ /* { dg-bogus "warning: region is gang partitioned but does not contain gang partitioned code" "TODO 'reduction'" { xfail *-*-* } .-3 } */
{
if (acc_on_device (acc_device_host))
{
@@ -617,6 +624,9 @@ int main ()
gangs_max = workers_max = vectors_max = INT_MIN;
#pragma acc serial copy (vectors_actual) /* { dg-warning "using vector_length \\(32\\), ignoring 1" "" { target openacc_nvidia_accel_selected } } */ \
copy (gangs_min, gangs_max, workers_min, workers_max, vectors_min, vectors_max)
+ /* { dg-bogus "warning: region contains gang partitioned code but is not gang partitioned" "TODO 'serial'" { xfail *-*-* } .-2 }
+ { dg-bogus "warning: region contains worker partitioned code but is not worker partitioned" "TODO 'serial'" { xfail *-*-* } .-3 }
+ { dg-bogus "warning: region contains vector partitioned code but is not vector partitioned" "TODO 'serial'" { xfail *-*-* } .-4 } */
{
if (acc_on_device (acc_device_nvidia))
{
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/parallel-reduction.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/parallel-reduction.c
index 80745816fff..b15ee8b22ff 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/parallel-reduction.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/parallel-reduction.c
@@ -1,4 +1,5 @@
-/* { dg-do run } */
+/* { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting
+ aspects of that functionality. */
#include <stdlib.h>
#include <openacc.h>
@@ -15,6 +16,7 @@ main ()
#pragma acc data copy (dummy)
{
#pragma acc parallel num_gangs (N) reduction (+:s1) copy(s1)
+ /* { dg-bogus "warning: region is gang partitioned but does not contain gang partitioned code" "TODO 'reduction'" { xfail *-*-* } .-1 } */
{
s1++;
}
@@ -35,6 +37,7 @@ main ()
s2 = 0;
#pragma acc parallel num_gangs (10) reduction (+:s1, s2) copy(s1, s2)
+ /* { dg-bogus "warning: region is gang partitioned but does not contain gang partitioned code" "TODO 'reduction'" { xfail *-*-* } .-1 } */
{
s1++;
s2 += N;
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85381-3.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85381-3.c
index cddbf271906..2a7e48116c3 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85381-3.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85381-3.c
@@ -2,10 +2,14 @@
{ dg-skip-if "" { *-*-* } { "*" } { "-O2" } } */
/* { dg-additional-options "-foffload=-fdump-rtl-mach" } */
+/* { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting
+ aspects of that functionality. */
+
int a;
#pragma acc declare create(a)
#pragma acc routine vector
+/* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .+2 } */
void __attribute__((noinline, noclone))
foo_v (void)
{
@@ -13,6 +17,8 @@ foo_v (void)
}
#pragma acc routine worker
+/* { dg-warning "region is worker partitioned but does not contain worker partitioned code" "" { target *-*-* } .+3 }
+ { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .+2 } */
void __attribute__((noinline, noclone))
foo_w (void)
{
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/private-variables.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/private-variables.c
index 53f03d17bb2..3cc6f150f63 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/private-variables.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/private-variables.c
@@ -1,3 +1,6 @@
+/* { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting
+ aspects of that functionality. */
+
#include <assert.h>
#include <openacc.h>
@@ -22,6 +25,8 @@ void local_g_1()
arr[i] = 3;
#pragma acc parallel copy(arr) num_gangs(32) num_workers(8) vector_length(32)
+ /* { dg-warning "region is worker partitioned but does not contain worker partitioned code" "" { target *-*-* } .-1 } */
+ /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .-2 } */
{
int x;
@@ -295,6 +300,8 @@ void loop_g_1()
arr[i] = i;
#pragma acc parallel copy(arr) num_gangs(32) num_workers(8) vector_length(32)
+ /* { dg-warning "region is worker partitioned but does not contain worker partitioned code" "" { target *-*-* } .-1 } */
+ /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .-2 } */
{
#pragma acc loop gang private(x)
for (i = 0; i < 32; i++)
@@ -320,6 +327,7 @@ void loop_g_2()
arr[i] = i;
#pragma acc parallel copy(arr) num_gangs(32) num_workers(8) vector_length(32)
+ /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .-1 } */
{
#pragma acc loop gang private(x)
for (i = 0; i < 32; i++)
@@ -348,6 +356,7 @@ void loop_g_3()
arr[i] = i;
#pragma acc parallel copy(arr) num_gangs(32) num_workers(8) vector_length(32)
+ /* { dg-warning "region is worker partitioned but does not contain worker partitioned code" "" { target *-*-* } .-1 } */
{
#pragma acc loop gang private(x)
for (i = 0; i < 32; i++)
@@ -376,6 +385,7 @@ void loop_g_4()
arr[i] = i;
#pragma acc parallel copy(arr) num_gangs(32) num_workers(8) vector_length(32)
+ /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .-1 } */
{
#pragma acc loop gang private(x)
for (i = 0; i < 32; i++)
@@ -408,6 +418,7 @@ void loop_g_5()
arr[i] = i;
#pragma acc parallel copy(arr) num_gangs(32) num_workers(8) vector_length(32)
+ /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .-1 } */
{
#pragma acc loop gang private(x)
for (i = 0; i < 32; i++)
@@ -438,6 +449,7 @@ void loop_g_6()
arr[i] = i;
#pragma acc parallel copy(arr) num_gangs(32) num_workers(8) vector_length(32)
+ /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .-1 } */
{
#pragma acc loop gang private(pt)
for (i = 0; i < 32; i++)
@@ -559,6 +571,7 @@ void loop_w_1()
arr[i] = i;
#pragma acc parallel copy(arr) num_gangs(32) num_workers(8) vector_length(32)
+ /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .-1 } */
{
int j;
@@ -875,6 +888,8 @@ void parallel_g_1()
arr[i] = 3;
#pragma acc parallel private(x) copy(arr) num_gangs(32) num_workers(8) vector_length(32)
+ /* { dg-warning "region is worker partitioned but does not contain worker partitioned code" "" { target *-*-* } .-1 } */
+ /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .-2 } */
{
#pragma acc loop gang(static:1)
for (i = 0; i < 32; i++)
@@ -904,6 +919,7 @@ void parallel_g_2()
arr[i] = i;
#pragma acc parallel private(x) copy(arr) num_gangs(32) num_workers(2) vector_length(32)
+ /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .-1 } */
{
#pragma acc loop gang
for (i = 0; i < 32; i++)
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/reduction-5.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/reduction-5.c
index 1cec9f152c4..bae1dee6ad2 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/reduction-5.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/reduction-5.c
@@ -1,4 +1,5 @@
-/* { dg-do run } */
+/* { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting
+ aspects of that functionality. */
/* Multiple reductions. */
@@ -45,6 +46,7 @@ main (void)
/* Nvptx targets require a vector_length or 32 in to allow spinlocks with
gangs. */
check_reduction (num_workers (nw) vector_length (vl), worker);
+ /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .-1 } */
check_reduction (vector_length (vl), vector);
check_reduction (num_gangs (ng) num_workers (nw) vector_length (vl), gang
worker vector);
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/reduction-7.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/reduction-7.c
index c4940b8ad9d..c2fb922a7f1 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/reduction-7.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/reduction-7.c
@@ -1,5 +1,8 @@
/* Tests of reduction on loop directive. */
+/* { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting
+ aspects of that functionality. */
+
#include <assert.h>
@@ -14,6 +17,8 @@ void g_np_1()
arr[i] = i;
#pragma acc parallel num_gangs(32) num_workers(32) vector_length(32)
+ /* { dg-warning "region is worker partitioned but does not contain worker partitioned code" "" { target *-*-* } .-1 } */
+ /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .-2 } */
{
#pragma acc loop gang reduction(+:res)
for (i = 0; i < 1024; i++)
@@ -28,6 +33,8 @@ void g_np_1()
res = hres = 1;
#pragma acc parallel num_gangs(32) num_workers(32) vector_length(32)
+ /* { dg-warning "region is worker partitioned but does not contain worker partitioned code" "" { target *-*-* } .-1 } */
+ /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .-2 } */
{
#pragma acc loop gang reduction(*:res)
for (i = 0; i < 12; i++)
@@ -52,6 +59,7 @@ void gv_np_1()
arr[i] = i;
#pragma acc parallel num_gangs(32) num_workers(32) vector_length(32)
+ /* { dg-warning "region is worker partitioned but does not contain worker partitioned code" "" { target *-*-* } .-1 } */
{
#pragma acc loop gang vector reduction(+:res)
for (i = 0; i < 1024; i++)
@@ -76,6 +84,7 @@ void gw_np_1()
arr[i] = i;
#pragma acc parallel num_gangs(32) num_workers(32) vector_length(32)
+ /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .-1 } */
{
#pragma acc loop gang worker reduction(+:res)
for (i = 0; i < 1024; i++)
@@ -239,6 +248,7 @@ void v_p_1()
#pragma acc parallel num_gangs(32) num_workers(32) vector_length(32) \
private(res) copyout(out)
+ /* { dg-warning "region is worker partitioned but does not contain worker partitioned code" "" { target *-*-* } .-2 } */
{
#pragma acc loop gang
for (j = 0; j < 32; j++)
@@ -315,6 +325,7 @@ void w_p_1()
#pragma acc parallel num_gangs(32) num_workers(32) vector_length(32) \
private(res) copyout(out)
+ /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .-2 } */
{
#pragma acc loop gang
for (j = 0; j < 32; j++)
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-g-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-g-1.c
index a164f576bc3..1536ce0ecb2 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-g-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-g-1.c
@@ -1,3 +1,6 @@
+/* { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting
+ aspects of that functionality. */
+
#include <stdio.h>
#include <openacc.h>
#include <gomp-constants.h>
@@ -6,6 +9,8 @@
#pragma acc routine gang
void __attribute__ ((noinline)) gang (int ary[N])
+/* { dg-warning "region is worker partitioned but does not contain worker partitioned code" "" { target *-*-* } .-1 } */
+/* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .-2 } */
{
#pragma acc loop gang
for (unsigned ix = 0; ix < N; ix++)
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-w-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-w-1.c
index acd9884cbd6..b9137d81935 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-w-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-w-1.c
@@ -1,3 +1,6 @@
+/* { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting
+ aspects of that functionality. */
+
#include <stdio.h>
#include <openacc.h>
#include <gomp-constants.h>
@@ -6,6 +9,7 @@
#pragma acc routine worker
void __attribute__ ((noinline)) worker (int ary[N])
+/* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .-1 } */
{
#pragma acc loop worker
for (unsigned ix = 0; ix < N; ix++)
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-wv-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-wv-2.c
index 9769ee72430..a03a2c2b163 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-wv-2.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-wv-2.c
@@ -1,3 +1,6 @@
+/* { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting
+ aspects of that functionality. */
+
#include <stdio.h>
#include <openacc.h>
#include <gomp-constants.h>
@@ -44,6 +47,7 @@ int DoWorkVec (int nw)
printf ("spawning %d ...", nw); fflush (stdout);
#pragma acc parallel num_workers(nw) vector_length (NUM_VECTORS) copy (ary)
+ /* { dg-warning "region contains vector partitioned code but is not vector partitioned" "" { target openacc_radeon_accel_selected } .-1 } */
{
WorkVec ((int *)ary, WIDTH, HEIGHT, nw, NUM_VECTORS);
}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-variable-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-variable-1.c
index 1d415cdcf76..0c071c37346 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-variable-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-variable-1.c
@@ -9,6 +9,9 @@
variables" (only visible to members of the GitHub OpenACC organization).
*/
+/* { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting
+ aspects of that functionality. */
+
#undef NDEBUG
#include <assert.h>
@@ -63,6 +66,9 @@ static void t0_c(void)
static const int t0_r_var_init = 61;
#pragma acc routine gang
+/* { dg-bogus "warning: region is gang partitioned but does not contain gang partitioned code" "TODO 'atomic'" { xfail *-*-* } .+4 } */
+/* { dg-warning "region is worker partitioned but does not contain worker partitioned code" "" { target *-*-* } .+3 } */
+/* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .+2 } */
__attribute__((noinline))
static int t0_r_r(void)
{
@@ -123,6 +129,7 @@ static void t1_c(void)
{
int result = 0;
int num_gangs_actual = -1;
+ /* { dg-bogus "warning: region is gang partitioned but does not contain gang partitioned code" "TODO 'atomic'" { xfail *-*-* } .+1 } */
#pragma acc parallel \
num_gangs(num_gangs_request) \
reduction(max:num_gangs_actual) \
@@ -153,6 +160,9 @@ static void t1_c(void)
static const int t1_r2_var_init = 166;
#pragma acc routine gang
+/* { dg-bogus "warning: region is gang partitioned but does not contain gang partitioned code" "TODO 'atomic'" { xfail *-*-* } .+4 } */
+/* { dg-warning "region is worker partitioned but does not contain worker partitioned code" "" { target *-*-* } .+3 } */
+/* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .+2 } */
__attribute__((noinline))
static int t1_r2_r(void)
{
@@ -245,6 +255,9 @@ static void t1_r2(void)
static const int t2_var_init_2 = -55;
#pragma acc routine gang
+/* { dg-bogus "warning: region is gang partitioned but does not contain gang partitioned code" "TODO 'atomic'" { xfail *-*-* } .+4 } */
+/* { dg-warning "region is worker partitioned but does not contain worker partitioned code" "" { target *-*-* } .+3 } */
+/* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .+2 } */
__attribute__((noinline))
static int t2_r(void)
{
@@ -286,6 +299,7 @@ static void t2(void)
itself, meaning that all 'i = 0' execution has finished (on the
device) before 'i = 1' is started (on the device), etc. */
+ /* { dg-bogus "warning: region is gang partitioned but does not contain gang partitioned code" "TODO 'atomic'" { xfail *-*-* } .+1 } */
#pragma acc parallel \
present(results_1) \
num_gangs(num_gangs_request_1) \
@@ -308,6 +322,7 @@ static void t2(void)
results_2[i][__builtin_goacc_parlevel_id(GOMP_DIM_GANG)] += t2_r();
}
+ /* { dg-bogus "warning: region is gang partitioned but does not contain gang partitioned code" "TODO 'atomic'" { xfail *-*-* } .+1 } */
#pragma acc parallel \
present(results_3) \
num_gangs(num_gangs_request_3) \
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/optional-private.f90 b/libgomp/testsuite/libgomp.oacc-fortran/optional-private.f90
index 0320bbb3bc9..4d36d869b0c 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/optional-private.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/optional-private.f90
@@ -4,6 +4,9 @@
! { dg-do run }
+! { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting
+! aspects of that functionality.
+
program main
implicit none
@@ -30,6 +33,8 @@ contains
end do
!$acc parallel copy(arr) num_gangs(32) num_workers(8) vector_length(32)
+ ! { dg-warning "region is worker partitioned but does not contain worker partitioned code" "" { target *-*-* } .-1 }
+ ! { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .-2 }
!$acc loop gang private(x)
do i = 1, 32
x = i * 2;
@@ -55,6 +60,7 @@ contains
end do
!$acc parallel copy(arr) num_gangs(32) num_workers(8) vector_length(32)
+ ! { dg-warning "region is worker partitioned but does not contain worker partitioned code" "" { target *-*-* } .-1 }
!$acc loop gang private(pt)
do i = 0, 31
pt%x = i
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/par-reduction-2-1.f b/libgomp/testsuite/libgomp.oacc-fortran/par-reduction-2-1.f
index aa1bb634ba6..408404cb7e6 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/par-reduction-2-1.f
+++ b/libgomp/testsuite/libgomp.oacc-fortran/par-reduction-2-1.f
@@ -4,6 +4,9 @@
! { dg-do run }
+! { dg-additional-options "-Wopenacc-parallelism" } for
+! testing/documenting aspects of that functionality.
+
PROGRAM MAIN
IMPLICIT NONE
INCLUDE "openacc_lib.h"
@@ -15,6 +18,9 @@
!$ACC PARALLEL NUM_GANGS(256) NUM_WORKERS(32) VECTOR_LENGTH(32)
!$ACC& REDUCTION(+:RES1) COPY(RES1, RES2) ASYNC(1)
+! { dg-bogus "warning: region is gang partitioned but does not contain gang partitioned code" "TODO 'reduction', 'atomic'" { xfail *-*-* } .-1 }
+! { dg-warning "region is worker partitioned but does not contain worker partitioned code" "" { target *-*-* } .-2 }
+! { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .-3 }
res1 = res1 + 5
!$ACC ATOMIC
@@ -37,6 +43,9 @@
!$ACC PARALLEL NUM_GANGS(8) NUM_WORKERS(32) VECTOR_LENGTH(32)
!$ACC& REDUCTION(*:RES1) COPY(RES1, RES2) ASYNC(1)
+! { dg-bogus "warning: region is gang partitioned but does not contain gang partitioned code" "TODO 'reduction', 'atomic'" { xfail *-*-* } .-1 }
+! { dg-warning "region is worker partitioned but does not contain worker partitioned code" "" { target *-*-* } .-2 }
+! { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .-3 }
res1 = res1 * 5
!$ACC ATOMIC
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/par-reduction-2-2.f b/libgomp/testsuite/libgomp.oacc-fortran/par-reduction-2-2.f
index 5694de1aee8..91ad99b256f 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/par-reduction-2-2.f
+++ b/libgomp/testsuite/libgomp.oacc-fortran/par-reduction-2-2.f
@@ -4,6 +4,9 @@
! { dg-do run }
+! { dg-additional-options "-Wopenacc-parallelism" } for
+! testing/documenting aspects of that functionality.
+
PROGRAM MAIN
USE OPENACC
IMPLICIT NONE
@@ -15,6 +18,9 @@
!$ACC PARALLEL NUM_GANGS(256) NUM_WORKERS(32) VECTOR_LENGTH(32)
!$ACC& REDUCTION(+:RES1) COPY(RES1, RES2) ASYNC(1)
+! { dg-bogus "warning: region is gang partitioned but does not contain gang partitioned code" "TODO 'reduction', 'atomic'" { xfail *-*-* } .-1 }
+! { dg-warning "region is worker partitioned but does not contain worker partitioned code" "" { target *-*-* } .-2 }
+! { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .-3 }
res1 = res1 + 5
!$ACC ATOMIC
@@ -37,6 +43,9 @@
!$ACC PARALLEL NUM_GANGS(8) NUM_WORKERS(32) VECTOR_LENGTH(32)
!$ACC& REDUCTION(*:RES1) COPY(RES1, RES2) ASYNC(1)
+! { dg-bogus "warning: region is gang partitioned but does not contain gang partitioned code" "TODO 'reduction', 'atomic'" { xfail *-*-* } .-1 }
+! { dg-warning "region is worker partitioned but does not contain worker partitioned code" "" { target *-*-* } .-2 }
+! { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .-3 }
res1 = res1 * 5
!$ACC ATOMIC
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/parallel-dims.f90 b/libgomp/testsuite/libgomp.oacc-fortran/parallel-dims.f90
index 1bfcd6ce099..fae99b338a5 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/parallel-dims.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/parallel-dims.f90
@@ -5,6 +5,9 @@
! { dg-do run }
! { dg-prune-output "command-line option '-fintrinsic-modules-path=.*' is valid for Fortran but not for C" }
+! { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting
+! aspects of that functionality.
+
! See also '../libgomp.oacc-c-c++-common/parallel-dims.c'.
module acc_routines
@@ -84,6 +87,9 @@ program main
vectors_max = -huge(gangs_max) - 1 ! INT_MIN
!$acc serial copy (vectors_actual) &
!$acc copy (gangs_min, gangs_max, workers_min, workers_max, vectors_min, vectors_max) ! { dg-warning "using vector_length \\(32\\), ignoring 1" "" { target openacc_nvidia_accel_selected } }
+ ! { dg-bogus "warning: region contains gang partitioned code but is not gang partitioned" "TODO 'serial'" { xfail *-*-* } .-1 }
+ ! { dg-bogus "warning: region contains worker partitioned code but is not worker partitioned" "TODO 'serial'" { xfail *-*-* } .-2 }
+ ! { dg-bogus "warning: region contains vector partitioned code but is not vector partitioned" "TODO 'serial'" { xfail *-*-* } .-3 }
if (acc_on_device (acc_device_nvidia)) then
! The GCC nvptx back end enforces vector_length (32).
! It's unclear if that's actually permissible here;
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/parallel-reduction.f90 b/libgomp/testsuite/libgomp.oacc-fortran/parallel-reduction.f90
index b926f4fe898..edf748efbb4 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/parallel-reduction.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/parallel-reduction.f90
@@ -1,5 +1,8 @@
! { dg-do run }
+! { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting
+! aspects of that functionality.
+
program reduction
implicit none
integer, parameter :: n = 10
@@ -10,6 +13,7 @@ program reduction
s2 = 0
!$acc parallel reduction(+:s1,s2) num_gangs (n) copy(s1)
+ ! { dg-bogus "warning: region is gang partitioned but does not contain gang partitioned code" "TODO 'reduction'" { xfail *-*-* } .-1 }
s1 = s1 + 1
s2 = s2 + 1
!$acc end parallel
@@ -40,6 +44,7 @@ subroutine redsub(s1, s2, n)
integer :: s1, s2, n
!$acc parallel reduction(+:s1,s2) num_gangs (10) copy(s1)
+ ! { dg-bogus "warning: region is gang partitioned but does not contain gang partitioned code" "TODO 'reduction'" { xfail *-*-* } .-1 }
s1 = s1 + 1
s2 = s2 + 1
!$acc end parallel
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/pr84028.f90 b/libgomp/testsuite/libgomp.oacc-fortran/pr84028.f90
index 2b361220bb6..e60322143ca 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/pr84028.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/pr84028.f90
@@ -1,11 +1,15 @@
! { dg-do run }
+! { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting
+! aspects of that functionality.
+
program foo
integer :: a(3,3,3), ll, lll
a = 1
!$acc parallel num_gangs(1) num_workers(2)
+ ! { dg-warning "region is worker partitioned but does not contain worker partitioned code" "" { target *-*-* } .-1 }
if (any(a(1:3,1:3,1:3).ne.1)) STOP 1
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/private-variables.f90 b/libgomp/testsuite/libgomp.oacc-fortran/private-variables.f90
index 472a6a14fff..81043a22fd8 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/private-variables.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/private-variables.f90
@@ -2,6 +2,9 @@
! { dg-do run }
+! { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting
+! aspects of that functionality.
+
! Test of gang-private variables declared on loop directive.
@@ -13,6 +16,8 @@ subroutine t1()
end do
!$acc parallel copy(arr) num_gangs(32) num_workers(8) vector_length(32)
+ ! { dg-warning "region is worker partitioned but does not contain worker partitioned code" "" { target *-*-* } .-1 }
+ ! { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .-2 }
!$acc loop gang private(x)
do i = 1, 32
x = i * 2;
@@ -37,6 +42,7 @@ subroutine t2()
end do
!$acc parallel copy(arr) num_gangs(32) num_workers(8) vector_length(32)
+ ! { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .-1 }
!$acc loop gang private(x)
do i = 0, 31
x = i * 2;
@@ -65,6 +71,7 @@ subroutine t3()
end do
!$acc parallel copy(arr) num_gangs(32) num_workers(8) vector_length(32)
+ ! { dg-warning "region is worker partitioned but does not contain worker partitioned code" "" { target *-*-* } .-1 }
!$acc loop gang private(x)
do i = 0, 31
x = i * 2;
@@ -98,6 +105,7 @@ subroutine t4()
end do
!$acc parallel copy(arr) num_gangs(32) num_workers(8) vector_length(32)
+ ! { dg-warning "region is worker partitioned but does not contain worker partitioned code" "" { target *-*-* } .-1 }
!$acc loop gang private(pt)
do i = 0, 31
pt%x = i
@@ -208,6 +216,7 @@ subroutine t7()
end do
!$acc parallel copy(arr) num_gangs(32) num_workers(8) vector_length(32)
+ ! { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .-1 }
!$acc loop gang private(x)
do i = 0, 31
!$acc loop worker private(x)
@@ -507,6 +516,8 @@ subroutine t14()
end do
!$acc parallel private(x) copy(arr) num_gangs(n) num_workers(8) vector_length(32)
+ ! { dg-warning "region is worker partitioned but does not contain worker partitioned code" "" { target *-*-* } .-1 }
+ ! { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .-2 }
!$acc loop gang(static:1)
do i = 1, n
x = i * 2;
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/reduction-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/reduction-1.f90
index 40c1a86ab92..95c3ed7ee9c 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/reduction-1.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/reduction-1.f90
@@ -1,5 +1,8 @@
! { dg-do run }
+! { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting
+! aspects of that functionality.
+
! Integer reductions
program reduction_1
@@ -279,6 +282,7 @@ program reduction_1
!$acc end parallel
!$acc parallel vector_length(vl) copy(rv)
+ ! { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .-1 }
!$acc loop reduction(ior:rv) gang
do i = 1, n
rv = ior (rv, array(i))
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/reduction-5.f90 b/libgomp/testsuite/libgomp.oacc-fortran/reduction-5.f90
index a529e22d9a2..5d1c1d99fce 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/reduction-5.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/reduction-5.f90
@@ -1,5 +1,8 @@
! { dg-do run }
+! { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting
+! aspects of that functionality.
+
! subroutine reduction
program reduction
@@ -45,6 +48,7 @@ subroutine redsub_worker(sum, n, c)
sum = 0
!$acc parallel copyin (n, c) num_workers(4) vector_length (32) copy(sum)
+ ! { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .-1 }
!$acc loop reduction(+:sum) worker
do i = 1, n
sum = sum + c
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/reduction-6.f90 b/libgomp/testsuite/libgomp.oacc-fortran/reduction-6.f90
index 249ed2edcbc..6908d16418c 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/reduction-6.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/reduction-6.f90
@@ -1,6 +1,9 @@
! { dg-do run }
! { dg-additional-options "-cpp" }
+! { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting
+! aspects of that functionality.
+
program reduction
implicit none
@@ -28,6 +31,7 @@ program reduction
!$acc end parallel
!$acc parallel num_workers (4) vector_length (32)
+ ! { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .-1 }
!$acc loop reduction(+:ws1, ws2) worker
do i = 1, n
ws1 = ws1 + 1
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/routine-7.f90 b/libgomp/testsuite/libgomp.oacc-fortran/routine-7.f90
index 1009f4a81e5..c34de3a4963 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/routine-7.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/routine-7.f90
@@ -2,6 +2,10 @@
! { dg-do run }
! { dg-additional-options "-cpp" }
+! { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting
+! aspects of that functionality.
+!TODO { dg-additional-options "-fno-inline" } for stable results regarding OpenACC 'routine'.
+
#define M 8
#define N 32
@@ -97,6 +101,8 @@ end subroutine worker
subroutine gang (a)
!$acc routine gang
+ ! { dg-warning "region is worker partitioned but does not contain worker partitioned code" "" { target *-*-* } .-2 }
+ ! { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .-3 }
integer, intent (inout) :: a(N)
integer :: i
--
2.30.2
^ permalink raw reply [flat|nested] 7+ messages in thread
* Re: [PATCH] openacc: Warnings for strange partitioning choices for parallel regions
2021-04-26 10:39 ` Thomas Schwinge
@ 2021-04-26 12:32 ` Tobias Burnus
2021-04-26 19:54 ` Thomas Schwinge
2021-04-26 20:00 ` [Patch, committed] OpenACC: Fix pattern in dg-bogus in Fortran testcases (was: " Tobias Burnus
0 siblings, 2 replies; 7+ messages in thread
From: Tobias Burnus @ 2021-04-26 12:32 UTC (permalink / raw)
To: Thomas Schwinge, Julian Brown, gcc-patches; +Cc: Jakub Jelinek, Tobias Burnus
Hi Thomas,
first, can you add an entry regarding this flag
to https://gcc.gnu.org/gcc-12/changes.html ?
Secondly, I now see FAILs like:
FAIL: gfortran.dg/goacc/classify-serial.f95 -O (test for excess errors)
Excess errors:
gfortran.dg/goacc/classify-serial.f95:20:132: Warning: region contains gang partitioned code but is not gang partitioned [-Wopenacc-parallelism]
gfortran.dg/goacc/classify-serial.f95:20:132: Warning: region contains vector partitioned code but is not vector partitioned [-Wopenacc-parallelism]
FAIL: gfortran.dg/goacc/kernels-decompose-2.f95 -O (test for excess errors)
Excess errors:
gfortran.dg/goacc/kernels-decompose-2.f95:124:15: Warning: region contains gang partitioned code but is not gang partitioned [-Wopenacc-parallelism]
FAIL: gfortran.dg/goacc/routine-module-mod-1.f90 (test for excess errors)
Excess errors:
gfortran.dg/goacc/routine-module-mod-1.f90:56:16: Warning: region is worker partitioned but does not contain worker partitioned code [-Wopenacc-parallelism]
Tobias
On 26.04.21 12:39, Thomas Schwinge wrote:
> Hi!
>
> On 2021-02-26T04:21:54-0800, Julian Brown <julian@codesourcery.com> wrote:
>> This patch adds warnings for strange partitioning choices -- specifying
>> either more or fewer partitioning levels on a parallel directive than the
>> enclosed offload region actually uses.
> Thanks!
>
>> We've used a version of this patch on the og8/og9/og10 branches for
>> quite a while. Versions have been posted for mainline submission as
>> part of a larger patch several times, e.g. here:
>>
>> https://gcc.gnu.org/pipermail/gcc-patches/2018-October/507938.html
>>
>> One motivation for committing this patch is it removes an ongoing
>> divergence in a large number of test cases between the og* branches and
>> mainline, namely whether the added warnings are expected as a result of
>> compiling those test cases, or not.
>>
>> The warnings themselves are perhaps slightly aggressive, but are intended
>> to help the user write more efficient code.
> As I'd said in
> <http://mid.mail-archive.com/87zhoefcpf.fsf@euler.schwinge.homeip.net>:
>
> | The general intention is good, but I've seen cases where I considered
> | these diagnostics to be too noisy. See also the several 'dg-bogus'
> | with XFAIL
>
> ... as well as 'dg-additional-options "-w"' being added.
>
> I've spent some more time on this; in particular, 'dg-bogus' with XFAIL
> setup, to indicate individual items that I think need to be addressed.
> (Please let me know if anybody makes a different categorization regarding
> these.) Of course, not all of our testcases are representative of
> real-world code (but I've anyway removed '-w' in favor of 'dg-warning' or
> 'dg-bogus' with XFAIL, for coverage), but I still came to the conclusion
> that this is too noisy to enable by default (as had been proposed), and
> thus for the time being have hidden it behind '-Wopenacc-parallelism'.
> (Adding an option is good anyway, so that users are able to disable
> unwanted diagnostics; we need to do that for other OpenACC diagnostics,
> too.) The goal is to eventually enable '-Wopenacc-parallelism', with
> '-Wall' and some "special cases" with '-Wextra', I suppose.
>
> I've pushed "Add '-Wopenacc-parallelism'" to master branch in commit
> 22cff118f7526bec195ed6e41452980820fdf3a8, see attached.
>
>
> Grüße
> Thomas
>
>
-----------------
Mentor Graphics (Deutschland) GmbH, Arnulfstrasse 201, 80634 München Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Frank Thürauf
^ permalink raw reply [flat|nested] 7+ messages in thread
* Re: [PATCH] openacc: Warnings for strange partitioning choices for parallel regions
2021-04-26 12:32 ` Tobias Burnus
@ 2021-04-26 19:54 ` Thomas Schwinge
2021-04-26 20:51 ` Tobias Burnus
2021-04-26 20:00 ` [Patch, committed] OpenACC: Fix pattern in dg-bogus in Fortran testcases (was: " Tobias Burnus
1 sibling, 1 reply; 7+ messages in thread
From: Thomas Schwinge @ 2021-04-26 19:54 UTC (permalink / raw)
To: Tobias Burnus; +Cc: Jakub Jelinek, Julian Brown, gcc-patches, fortran
Hi!
On 2021-04-26T14:32:10+0200, Tobias Burnus <tobias@codesourcery.com> wrote:
> first, can you add an entry regarding this flag
> to https://gcc.gnu.org/gcc-12/changes.html ?
Is that a standard thing that all new command-line flags do get
highlighted there? (I wouldn't have considered that one important
enough?)
> Secondly, I now see FAILs like:
>
> FAIL: gfortran.dg/goacc/classify-serial.f95 -O (test for excess errors)
> Excess errors:
> gfortran.dg/goacc/classify-serial.f95:20:132: Warning: region contains gang partitioned code but is not gang partitioned [-Wopenacc-parallelism]
> gfortran.dg/goacc/classify-serial.f95:20:132: Warning: region contains vector partitioned code but is not vector partitioned [-Wopenacc-parallelism]
Eek! I do reproduce that -- but only in a GCC build configuration
without offloading enabled! For example, for
'gfortran.dg/goacc/classify-serial.f95', we've got the following diff in
diagnostics in a GCC build configuration without vs. with offloading
enabled:
{+[...]/gcc/testsuite/gfortran.dg/goacc/classify-serial.f95: In function 'MAIN__._omp_fn.0':+}
[...]/gcc/testsuite/gfortran.dg/goacc/classify-serial.f95:20:132: [-Warning:-]{+warning:+} region contains gang partitioned code but is not gang partitioned [-Wopenacc-parallelism]
[...]/gcc/testsuite/gfortran.dg/goacc/classify-serial.f95:20:132: [-Warning:-]{+warning:+} region contains vector partitioned code but is not vector partitioned [-Wopenacc-parallelism]
[...]/gcc/testsuite/gfortran.dg/goacc/classify-serial.f95:20:132: optimized: assigned OpenACC gang vector loop parallelism
PASS: gfortran.dg/goacc/classify-serial.f95 -O (test for warnings, line 20)
[-XPASS:-]{+XFAIL:+} gfortran.dg/goacc/classify-serial.f95 -O TODO 'serial' (test for bogus messages, line 20)
PASS: gfortran.dg/goacc/classify-serial.f95 -O (test for bogus messages, line 20)
[-XPASS:-]{+XFAIL:+} gfortran.dg/goacc/classify-serial.f95 -O TODO 'serial' (test for bogus messages, line 20)
[-FAIL:-]{+PASS:+} gfortran.dg/goacc/classify-serial.f95 -O (test for excess errors)[-Excess errors:-]
[-[...]/gcc/testsuite/gfortran.dg/goacc/classify-serial.f95:20:132: Warning: region contains gang partitioned code but is not gang partitioned [-Wopenacc-parallelism]-]
[-[...]/gcc/testsuite/gfortran.dg/goacc/classify-serial.f95:20:132: Warning: region contains vector partitioned code but is not vector partitioned [-Wopenacc-parallelism]-]
PASS: gfortran.dg/goacc/classify-serial.f95 -O scan-tree-dump-times ompexp "(?n)__attribute__\\(\\(oacc serial, omp target entrypoint\\)\\)" 1
PASS: gfortran.dg/goacc/classify-serial.f95 -O scan-tree-dump-times oaccdevlow "(?n)Function is OpenACC serial offload" 1
PASS: gfortran.dg/goacc/classify-serial.f95 -O scan-tree-dump-times oaccdevlow "(?n)Compute dimensions \\[1, 1, 1\\]" 1
Notice upper-case "Warning" vs. lower-case "warning". That's for Fortran
only; for C, C++, it's lower-case "warning" for both build variants.
It's of course easy to fix up the regexp, but should we maybe rather
figure out how to unify the reporting?
I do understand that Fortran has some upper-case diagnostics, but I don't
understand the difference/relevance of GCC build configuration without
vs. with offloading enabled, how is that coming becoming relevant?
Also notice that the preamble "In function 'MAIN__._omp_fn.0':" only
appears in the GCC build configuration with offloading. Looking at
'c-c++-common/goacc/classify-serial.c', for C, we've got "In function
'SERIAL._omp_fn.0':" for both GCC build configurations, and for C++,
we've got "In function '_Z6SERIALv._omp_fn.0':" without offloading
enabled vs. "In function 'SERIAL() [clone ._omp_fn.0]':" with offloading
enabled.
So there generally seems to be some difference in function outlining
setup without vs. with offloading enabled? ..., and for Fortran, with
offloading enabled, that causes the outlined function to be
"de-Fortranified", thus the lower-case "warning" diagnostic?
> FAIL: gfortran.dg/goacc/kernels-decompose-2.f95 -O (test for excess errors)
> Excess errors:
> gfortran.dg/goacc/kernels-decompose-2.f95:124:15: Warning: region contains gang partitioned code but is not gang partitioned [-Wopenacc-parallelism]
>
> FAIL: gfortran.dg/goacc/routine-module-mod-1.f90 (test for excess errors)
> Excess errors:
> gfortran.dg/goacc/routine-module-mod-1.f90:56:16: Warning: region is worker partitioned but does not contain worker partitioned code [-Wopenacc-parallelism]
..., and similarly, several more in 'libgomp.oacc-fortran'.
Grüße
Thomas
> On 26.04.21 12:39, Thomas Schwinge wrote:
>
>> Hi!
>>
>> On 2021-02-26T04:21:54-0800, Julian Brown <julian@codesourcery.com> wrote:
>>> This patch adds warnings for strange partitioning choices -- specifying
>>> either more or fewer partitioning levels on a parallel directive than the
>>> enclosed offload region actually uses.
>> Thanks!
>>
>>> We've used a version of this patch on the og8/og9/og10 branches for
>>> quite a while. Versions have been posted for mainline submission as
>>> part of a larger patch several times, e.g. here:
>>>
>>> https://gcc.gnu.org/pipermail/gcc-patches/2018-October/507938.html
>>>
>>> One motivation for committing this patch is it removes an ongoing
>>> divergence in a large number of test cases between the og* branches and
>>> mainline, namely whether the added warnings are expected as a result of
>>> compiling those test cases, or not.
>>>
>>> The warnings themselves are perhaps slightly aggressive, but are intended
>>> to help the user write more efficient code.
>> As I'd said in
>> <http://mid.mail-archive.com/87zhoefcpf.fsf@euler.schwinge.homeip.net>:
>>
>> | The general intention is good, but I've seen cases where I considered
>> | these diagnostics to be too noisy. See also the several 'dg-bogus'
>> | with XFAIL
>>
>> ... as well as 'dg-additional-options "-w"' being added.
>>
>> I've spent some more time on this; in particular, 'dg-bogus' with XFAIL
>> setup, to indicate individual items that I think need to be addressed.
>> (Please let me know if anybody makes a different categorization regarding
>> these.) Of course, not all of our testcases are representative of
>> real-world code (but I've anyway removed '-w' in favor of 'dg-warning' or
>> 'dg-bogus' with XFAIL, for coverage), but I still came to the conclusion
>> that this is too noisy to enable by default (as had been proposed), and
>> thus for the time being have hidden it behind '-Wopenacc-parallelism'.
>> (Adding an option is good anyway, so that users are able to disable
>> unwanted diagnostics; we need to do that for other OpenACC diagnostics,
>> too.) The goal is to eventually enable '-Wopenacc-parallelism', with
>> '-Wall' and some "special cases" with '-Wextra', I suppose.
>>
>> I've pushed "Add '-Wopenacc-parallelism'" to master branch in commit
>> 22cff118f7526bec195ed6e41452980820fdf3a8, see attached.
>>
>>
>> Grüße
>> Thomas
>>
>>
> -----------------
> Mentor Graphics (Deutschland) GmbH, Arnulfstrasse 201, 80634 München Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Frank Thürauf
-----------------
Mentor Graphics (Deutschland) GmbH, Arnulfstrasse 201, 80634 München Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Frank Thürauf
^ permalink raw reply [flat|nested] 7+ messages in thread
* [Patch, committed] OpenACC: Fix pattern in dg-bogus in Fortran testcases (was: Re: [PATCH] openacc: Warnings for strange partitioning choices for parallel regions)
2021-04-26 12:32 ` Tobias Burnus
2021-04-26 19:54 ` Thomas Schwinge
@ 2021-04-26 20:00 ` Tobias Burnus
1 sibling, 0 replies; 7+ messages in thread
From: Tobias Burnus @ 2021-04-26 20:00 UTC (permalink / raw)
To: Tobias Burnus, Thomas Schwinge, Julian Brown, gcc-patches; +Cc: Jakub Jelinek
[-- Attachment #1: Type: text/plain, Size: 464 bytes --]
On 26.04.21 14:32, Tobias Burnus wrote:
> Secondly, I now see FAILs like:
>
> FAIL: gfortran.dg/goacc/classify-serial.f95 -O (test for excess errors)
> [...]
Now fixed myself to reduce the clutter.
Committed as r12-130-g5a26ba75de623f75fb44cddc2a9c982d31c96213
Tobias
-----------------
Mentor Graphics (Deutschland) GmbH, Arnulfstrasse 201, 80634 München Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Frank Thürauf
[-- Attachment #2: committed.diff --]
[-- Type: text/x-patch, Size: 10052 bytes --]
commit 5a26ba75de623f75fb44cddc2a9c982d31c96213
Author: Tobias Burnus <tobias@codesourcery.com>
Date: Mon Apr 26 21:57:31 2021 +0200
OpenACC: Fix pattern in dg-bogus in Fortran testcases
libgomp/ChangeLog:
* testsuite/libgomp.oacc-fortran/par-reduction-2-1.f:
Correct spelling in dg-bogus to match -Wopenacc-parallelism.
* testsuite/libgomp.oacc-fortran/par-reduction-2-2.f: Likewise.
* testsuite/libgomp.oacc-fortran/parallel-dims.f90: Likewise.
* testsuite/libgomp.oacc-fortran/parallel-reduction.f90: Likewise.
gcc/testsuite/ChangeLog:
* gfortran.dg/goacc/classify-serial.f95:
Correct spelling in dg-bogus to match -Wopenacc-parallelism.
* gfortran.dg/goacc/kernels-decompose-2.f95: Likewise.
* gfortran.dg/goacc/routine-module-mod-1.f90: Likewise.
diff --git a/gcc/testsuite/gfortran.dg/goacc/classify-serial.f95 b/gcc/testsuite/gfortran.dg/goacc/classify-serial.f95
index 6d402e683de..386b95e78e4 100644
--- a/gcc/testsuite/gfortran.dg/goacc/classify-serial.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/classify-serial.f95
@@ -18,9 +18,9 @@ program main
call setup(a, b)
!$acc serial loop copyin (a(0:n-1), b(0:n-1)) copyout (c(0:n-1)) ! { dg-message "optimized: assigned OpenACC gang vector loop parallelism" }
- ! { dg-bogus "warning: region contains gang partitioned code but is not gang partitioned" "TODO 'serial'" { xfail *-*-* } .-1 }
- ! { dg-bogus "warning: region contains worker partitioned code but is not worker partitioned" "" { target *-*-* } .-2 }
- ! { dg-bogus "warning: region contains vector partitioned code but is not vector partitioned" "TODO 'serial'" { xfail *-*-* } .-3 }
+ ! { dg-bogus "Warning: region contains gang partitioned code but is not gang partitioned" "TODO 'serial'" { xfail *-*-* } .-1 }
+ ! { dg-bogus "Warning: region contains worker partitioned code but is not worker partitioned" "" { target *-*-* } .-2 }
+ ! { dg-bogus "Warning: region contains vector partitioned code but is not vector partitioned" "TODO 'serial'" { xfail *-*-* } .-3 }
do i = 0, n - 1
c(i) = a(i) + b(i)
end do
diff --git a/gcc/testsuite/gfortran.dg/goacc/kernels-decompose-2.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-decompose-2.f95
index a3cec8325c7..1277ad9ddda 100644
--- a/gcc/testsuite/gfortran.dg/goacc/kernels-decompose-2.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/kernels-decompose-2.f95
@@ -122,7 +122,7 @@ program main
!$acc end kernels
!$acc kernels
- ! { dg-bogus "warning: region contains gang partitioned code but is not gang partitioned" "TODO 'kernels'" { xfail *-*-* } .-1 }
+ ! { dg-bogus "Warning: region contains gang partitioned code but is not gang partitioned" "TODO 'kernels'" { xfail *-*-* } .-1 }
y = f_g (a(5)) ! { dg-line l_part[incr c_part] }
!TODO If such a construct is placed in its own part (like it is, here), can't this actually use gang paralelism, instead of "gang-single"?
! { dg-message "note: beginning 'gang-single' part in OpenACC 'kernels' region" "" { target *-*-* } l_part$c_part }
diff --git a/gcc/testsuite/gfortran.dg/goacc/routine-module-mod-1.f90 b/gcc/testsuite/gfortran.dg/goacc/routine-module-mod-1.f90
index 22106508141..198ae01914b 100644
--- a/gcc/testsuite/gfortran.dg/goacc/routine-module-mod-1.f90
+++ b/gcc/testsuite/gfortran.dg/goacc/routine-module-mod-1.f90
@@ -56,7 +56,7 @@ contains
subroutine g_1
implicit none
!$acc routine gang
- ! { dg-bogus "warning: region is worker partitioned but does not contain worker partitioned code" "TODO default 'gang' 'vector'" { xfail *-*-* } .-3 }
+ ! { dg-bogus "Warning: region is worker partitioned but does not contain worker partitioned code" "TODO default 'gang' 'vector'" { xfail *-*-* } .-3 }
integer :: i
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/par-reduction-2-1.f b/libgomp/testsuite/libgomp.oacc-fortran/par-reduction-2-1.f
index 408404cb7e6..baa472b3e87 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/par-reduction-2-1.f
+++ b/libgomp/testsuite/libgomp.oacc-fortran/par-reduction-2-1.f
@@ -18,7 +18,7 @@
!$ACC PARALLEL NUM_GANGS(256) NUM_WORKERS(32) VECTOR_LENGTH(32)
!$ACC& REDUCTION(+:RES1) COPY(RES1, RES2) ASYNC(1)
-! { dg-bogus "warning: region is gang partitioned but does not contain gang partitioned code" "TODO 'reduction', 'atomic'" { xfail *-*-* } .-1 }
+! { dg-bogus "Warning: region is gang partitioned but does not contain gang partitioned code" "TODO 'reduction', 'atomic'" { xfail *-*-* } .-1 }
! { dg-warning "region is worker partitioned but does not contain worker partitioned code" "" { target *-*-* } .-2 }
! { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .-3 }
res1 = res1 + 5
@@ -43,7 +43,7 @@
!$ACC PARALLEL NUM_GANGS(8) NUM_WORKERS(32) VECTOR_LENGTH(32)
!$ACC& REDUCTION(*:RES1) COPY(RES1, RES2) ASYNC(1)
-! { dg-bogus "warning: region is gang partitioned but does not contain gang partitioned code" "TODO 'reduction', 'atomic'" { xfail *-*-* } .-1 }
+! { dg-bogus "Warning: region is gang partitioned but does not contain gang partitioned code" "TODO 'reduction', 'atomic'" { xfail *-*-* } .-1 }
! { dg-warning "region is worker partitioned but does not contain worker partitioned code" "" { target *-*-* } .-2 }
! { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .-3 }
res1 = res1 * 5
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/par-reduction-2-2.f b/libgomp/testsuite/libgomp.oacc-fortran/par-reduction-2-2.f
index 91ad99b256f..4ebab6c7b18 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/par-reduction-2-2.f
+++ b/libgomp/testsuite/libgomp.oacc-fortran/par-reduction-2-2.f
@@ -18,7 +18,7 @@
!$ACC PARALLEL NUM_GANGS(256) NUM_WORKERS(32) VECTOR_LENGTH(32)
!$ACC& REDUCTION(+:RES1) COPY(RES1, RES2) ASYNC(1)
-! { dg-bogus "warning: region is gang partitioned but does not contain gang partitioned code" "TODO 'reduction', 'atomic'" { xfail *-*-* } .-1 }
+! { dg-bogus "Warning: region is gang partitioned but does not contain gang partitioned code" "TODO 'reduction', 'atomic'" { xfail *-*-* } .-1 }
! { dg-warning "region is worker partitioned but does not contain worker partitioned code" "" { target *-*-* } .-2 }
! { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .-3 }
res1 = res1 + 5
@@ -43,7 +43,7 @@
!$ACC PARALLEL NUM_GANGS(8) NUM_WORKERS(32) VECTOR_LENGTH(32)
!$ACC& REDUCTION(*:RES1) COPY(RES1, RES2) ASYNC(1)
-! { dg-bogus "warning: region is gang partitioned but does not contain gang partitioned code" "TODO 'reduction', 'atomic'" { xfail *-*-* } .-1 }
+! { dg-bogus "Warning: region is gang partitioned but does not contain gang partitioned code" "TODO 'reduction', 'atomic'" { xfail *-*-* } .-1 }
! { dg-warning "region is worker partitioned but does not contain worker partitioned code" "" { target *-*-* } .-2 }
! { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .-3 }
res1 = res1 * 5
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/parallel-dims.f90 b/libgomp/testsuite/libgomp.oacc-fortran/parallel-dims.f90
index fae99b338a5..d9847969369 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/parallel-dims.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/parallel-dims.f90
@@ -87,9 +87,9 @@ program main
vectors_max = -huge(gangs_max) - 1 ! INT_MIN
!$acc serial copy (vectors_actual) &
!$acc copy (gangs_min, gangs_max, workers_min, workers_max, vectors_min, vectors_max) ! { dg-warning "using vector_length \\(32\\), ignoring 1" "" { target openacc_nvidia_accel_selected } }
- ! { dg-bogus "warning: region contains gang partitioned code but is not gang partitioned" "TODO 'serial'" { xfail *-*-* } .-1 }
- ! { dg-bogus "warning: region contains worker partitioned code but is not worker partitioned" "TODO 'serial'" { xfail *-*-* } .-2 }
- ! { dg-bogus "warning: region contains vector partitioned code but is not vector partitioned" "TODO 'serial'" { xfail *-*-* } .-3 }
+ ! { dg-bogus "Warning: region contains gang partitioned code but is not gang partitioned" "TODO 'serial'" { xfail *-*-* } .-1 }
+ ! { dg-bogus "Warning: region contains worker partitioned code but is not worker partitioned" "TODO 'serial'" { xfail *-*-* } .-2 }
+ ! { dg-bogus "Warning: region contains vector partitioned code but is not vector partitioned" "TODO 'serial'" { xfail *-*-* } .-3 }
if (acc_on_device (acc_device_nvidia)) then
! The GCC nvptx back end enforces vector_length (32).
! It's unclear if that's actually permissible here;
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/parallel-reduction.f90 b/libgomp/testsuite/libgomp.oacc-fortran/parallel-reduction.f90
index edf748efbb4..af25f0dac9b 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/parallel-reduction.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/parallel-reduction.f90
@@ -13,7 +13,7 @@ program reduction
s2 = 0
!$acc parallel reduction(+:s1,s2) num_gangs (n) copy(s1)
- ! { dg-bogus "warning: region is gang partitioned but does not contain gang partitioned code" "TODO 'reduction'" { xfail *-*-* } .-1 }
+ ! { dg-bogus "Warning: region is gang partitioned but does not contain gang partitioned code" "TODO 'reduction'" { xfail *-*-* } .-1 }
s1 = s1 + 1
s2 = s2 + 1
!$acc end parallel
@@ -44,7 +44,7 @@ subroutine redsub(s1, s2, n)
integer :: s1, s2, n
!$acc parallel reduction(+:s1,s2) num_gangs (10) copy(s1)
- ! { dg-bogus "warning: region is gang partitioned but does not contain gang partitioned code" "TODO 'reduction'" { xfail *-*-* } .-1 }
+ ! { dg-bogus "Warning: region is gang partitioned but does not contain gang partitioned code" "TODO 'reduction'" { xfail *-*-* } .-1 }
s1 = s1 + 1
s2 = s2 + 1
!$acc end parallel
^ permalink raw reply [flat|nested] 7+ messages in thread
* Re: [PATCH] openacc: Warnings for strange partitioning choices for parallel regions
2021-04-26 19:54 ` Thomas Schwinge
@ 2021-04-26 20:51 ` Tobias Burnus
2021-04-26 21:23 ` [Patch, committed] OpenACC: Fix pattern in dg-bogus in Fortran testcases again (Re: [PATCH] openacc: Warnings for strange partitioning choices for parallel regions) Tobias Burnus
0 siblings, 1 reply; 7+ messages in thread
From: Tobias Burnus @ 2021-04-26 20:51 UTC (permalink / raw)
To: Thomas Schwinge; +Cc: Jakub Jelinek, Julian Brown, gcc-patches, fortran
On 26.04.21 21:54, Thomas Schwinge wrote:
> On 2021-04-26T14:32:10+0200, Tobias Burnus <tobias@codesourcery.com> wrote:
>> first, can you add an entry regarding this flag
>> to https://gcc.gnu.org/gcc-12/changes.html ?
> Is that a standard thing that all new command-line flags do get
> highlighted there? (I wouldn't have considered that one important
> enough?)
I think options which one expects to to use as user – especially when
they are for consideration with -W... do make sense. In general, as a
user I am happy to have a more detailed information about changes than
'something has changed'.
>> Secondly, I now see FAILs like:
>>
>> FAIL: gfortran.dg/goacc/classify-serial.f95 -O (test for excess errors)
>> Excess errors:
>> gfortran.dg/goacc/classify-serial.f95:20:132: Warning: region contains gang partitioned code but is not gang partitioned [-Wopenacc-parallelism]
>> gfortran.dg/goacc/classify-serial.f95:20:132: Warning: region contains vector partitioned code but is not vector partitioned [-Wopenacc-parallelism]
> Eek! I do reproduce that -- but only in a GCC build configuration
> without offloading enabled!
Aha, that explains why you did commit it like that – now I have fixed
the nonoffload version – and broken the offloaded one.
> Notice upper-case "Warning" vs. lower-case "warning". That's for Fortran
> only; for C, C++, it's lower-case "warning" for both build variants.
>
> It's of course easy to fix up the regexp, but should we maybe rather
> figure out how to unify the reporting?
>
> I do understand that Fortran has some upper-case diagnostics, but I don't
> understand the difference/relevance of GCC build configuration without
> vs. with offloading enabled, how is that coming becoming relevant?
The default is 'global_dc = &global_diagnostic_context' which gets
overridden in fortran/error.c.
I have no idea why there is a difference but wonder whether it has
something to do with taking the LTO path (flag_generate_offload) and
some overriding of the settings.
I believe this is done in tree.c's free_lang_data:
if (in_lto_p
|| (!flag_generate_lto && !flag_generate_offload))
...
return 0;
}
/* Reset diagnostic machinery. */
tree_diagnostics_defaults (global_dc);
And in cgraphunit.c's symbol_table::finalize_compilation_unit
if (!in_lto_p && g->have_offload)
flag_generate_offload = 1;
And g->have_offload is guarded by ENABLE_OFFLOAD.
> Also notice that the preamble "In function 'MAIN__._omp_fn.0':" only
> appears in the GCC build configuration with offloading. Looking at
> 'c-c++-common/goacc/classify-serial.c', for C, we've got "In function
> 'SERIAL._omp_fn.0':" for both GCC build configurations, and for C++,
> we've got "In function '_Z6SERIALv._omp_fn.0':" without offloading
> enabled vs. "In function 'SERIAL() [clone ._omp_fn.0]':" with offloading
> enabled.
I note that:
langhooks.c: (context->printer, _("In function %qs"),
and in cp/error.c's cp_print_error_function:
pp_printf (context->printer, function_category (fndecl),
cxx_printable_name_translate (fndecl, 2));
My bet is that here the same happens as above – the global_dc gets
reset – and the diagnostic becomes different.
Tobias
-----------------
Mentor Graphics (Deutschland) GmbH, Arnulfstrasse 201, 80634 München Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Frank Thürauf
^ permalink raw reply [flat|nested] 7+ messages in thread
* [Patch, committed] OpenACC: Fix pattern in dg-bogus in Fortran testcases again (Re: [PATCH] openacc: Warnings for strange partitioning choices for parallel regions)
2021-04-26 20:51 ` Tobias Burnus
@ 2021-04-26 21:23 ` Tobias Burnus
0 siblings, 0 replies; 7+ messages in thread
From: Tobias Burnus @ 2021-04-26 21:23 UTC (permalink / raw)
To: Thomas Schwinge; +Cc: Jakub Jelinek, Julian Brown, gcc-patches, fortran
[-- Attachment #1: Type: text/plain, Size: 1349 bytes --]
Hi all,
as discussed below/in this thread, the reset of the diagnostic
contest affects the output with ENABLE_OFFLOAD. The attached
patch now uses [Ww]arning in dg-bogus.
It might be possible to reduce the differences between
ENABLE_OFFLOAD being true or false, but it does not seem to
be simple and may not be worth the effort.
In any case, this patch should permit to have no FAIL
with ENABLE_OFFLOAD set and unset. – And should make the
testers happy.
Committed as r12-135-gbd7ebe9da745a62184052dd1b15f4dd10fbdc9f4
Tobias
PS: 'Warning:' with Fortran diagnostic and ENABLE_OFFSET unset,
'warning:' with default diagnostic and ENABLE_OFFSET set – due to:
On 26.04.21 22:51, Tobias Burnus wrote:
> I believe this is done in tree.c's free_lang_data:
> if (in_lto_p
> || (!flag_generate_lto && !flag_generate_offload))
> ...
> return 0;
> }
>
> /* Reset diagnostic machinery. */
> tree_diagnostics_defaults (global_dc);
>
> And in cgraphunit.c's symbol_table::finalize_compilation_unit
> if (!in_lto_p && g->have_offload)
> flag_generate_offload = 1;
>
> And g->have_offload is guarded by ENABLE_OFFLOAD.
-----------------
Mentor Graphics (Deutschland) GmbH, Arnulfstrasse 201, 80634 München Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Frank Thürauf
[-- Attachment #2: committed.diff --]
[-- Type: text/x-patch, Size: 11621 bytes --]
commit bd7ebe9da745a62184052dd1b15f4dd10fbdc9f4
Author: Tobias Burnus <tobias@codesourcery.com>
Date: Mon Apr 26 22:59:21 2021 +0200
OpenACC: Fix pattern in dg-bogus in Fortran testcases again
It turned out that a compiler built without offloading support
and one with can produce slightly different diagnostic.
Offloading support implies ENABLE_OFFLOAD which implies that
g->have_offload is set when offloading is actually needed.
In cgraphunit.c, the latter causes flag_generate_offload = 1,
which in turn affects tree.c's free_lang_data.
The result is that the front-end specific diagnostic gets reset
('tree_diagnostics_defaults (global_dc)'), which affects in this
case 'Warning' vs. 'warning' via the Fortran frontend.
Result: 'Warning:' vs. 'warning:'.
Side note: Other FE also override the diagnostic, leading to
similar differences, e.g. the C++ FE outputs mangled function
names differently, cf. patch thread.
libgomp/ChangeLog:
* testsuite/libgomp.oacc-fortran/par-reduction-2-1.f:
Use [Ww]arning in dg-bogus as FE diagnostic and default
diagnostic differ and the result depends on ENABLE_OFFLOAD.
* testsuite/libgomp.oacc-fortran/par-reduction-2-2.f: Likewise.
* testsuite/libgomp.oacc-fortran/parallel-dims.f90: Likewise.
* testsuite/libgomp.oacc-fortran/parallel-reduction.f90: Likewise.
gcc/testsuite/ChangeLog:
* gfortran.dg/goacc/classify-serial.f95:
Use [Ww]arning in dg-bogus as FE diagnostic and default
diagnostic differ and the result depends on ENABLE_OFFLOAD.
* gfortran.dg/goacc/kernels-decompose-2.f95: Likewise.
* gfortran.dg/goacc/routine-module-mod-1.f90: Likewise.
---
gcc/testsuite/gfortran.dg/goacc/classify-serial.f95 | 6 +++---
gcc/testsuite/gfortran.dg/goacc/kernels-decompose-2.f95 | 2 +-
gcc/testsuite/gfortran.dg/goacc/routine-module-mod-1.f90 | 2 +-
libgomp/testsuite/libgomp.oacc-fortran/par-reduction-2-1.f | 4 ++--
libgomp/testsuite/libgomp.oacc-fortran/par-reduction-2-2.f | 4 ++--
libgomp/testsuite/libgomp.oacc-fortran/parallel-dims.f90 | 6 +++---
libgomp/testsuite/libgomp.oacc-fortran/parallel-reduction.f90 | 4 ++--
7 files changed, 14 insertions(+), 14 deletions(-)
diff --git a/gcc/testsuite/gfortran.dg/goacc/classify-serial.f95 b/gcc/testsuite/gfortran.dg/goacc/classify-serial.f95
index 386b95e78e4..6dcb1b170f8 100644
--- a/gcc/testsuite/gfortran.dg/goacc/classify-serial.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/classify-serial.f95
@@ -18,9 +18,9 @@ program main
call setup(a, b)
!$acc serial loop copyin (a(0:n-1), b(0:n-1)) copyout (c(0:n-1)) ! { dg-message "optimized: assigned OpenACC gang vector loop parallelism" }
- ! { dg-bogus "Warning: region contains gang partitioned code but is not gang partitioned" "TODO 'serial'" { xfail *-*-* } .-1 }
- ! { dg-bogus "Warning: region contains worker partitioned code but is not worker partitioned" "" { target *-*-* } .-2 }
- ! { dg-bogus "Warning: region contains vector partitioned code but is not vector partitioned" "TODO 'serial'" { xfail *-*-* } .-3 }
+ ! { dg-bogus "\[Ww\]arning: region contains gang partitioned code but is not gang partitioned" "TODO 'serial'" { xfail *-*-* } .-1 }
+ ! { dg-bogus "\[Ww\]arning: region contains worker partitioned code but is not worker partitioned" "" { target *-*-* } .-2 }
+ ! { dg-bogus "\[Ww\]arning: region contains vector partitioned code but is not vector partitioned" "TODO 'serial'" { xfail *-*-* } .-3 }
do i = 0, n - 1
c(i) = a(i) + b(i)
end do
diff --git a/gcc/testsuite/gfortran.dg/goacc/kernels-decompose-2.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-decompose-2.f95
index 1277ad9ddda..238482b91a4 100644
--- a/gcc/testsuite/gfortran.dg/goacc/kernels-decompose-2.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/kernels-decompose-2.f95
@@ -122,7 +122,7 @@ program main
!$acc end kernels
!$acc kernels
- ! { dg-bogus "Warning: region contains gang partitioned code but is not gang partitioned" "TODO 'kernels'" { xfail *-*-* } .-1 }
+ ! { dg-bogus "\[Ww\]arning: region contains gang partitioned code but is not gang partitioned" "TODO 'kernels'" { xfail *-*-* } .-1 }
y = f_g (a(5)) ! { dg-line l_part[incr c_part] }
!TODO If such a construct is placed in its own part (like it is, here), can't this actually use gang paralelism, instead of "gang-single"?
! { dg-message "note: beginning 'gang-single' part in OpenACC 'kernels' region" "" { target *-*-* } l_part$c_part }
diff --git a/gcc/testsuite/gfortran.dg/goacc/routine-module-mod-1.f90 b/gcc/testsuite/gfortran.dg/goacc/routine-module-mod-1.f90
index 198ae01914b..835619c6509 100644
--- a/gcc/testsuite/gfortran.dg/goacc/routine-module-mod-1.f90
+++ b/gcc/testsuite/gfortran.dg/goacc/routine-module-mod-1.f90
@@ -56,7 +56,7 @@ contains
subroutine g_1
implicit none
!$acc routine gang
- ! { dg-bogus "Warning: region is worker partitioned but does not contain worker partitioned code" "TODO default 'gang' 'vector'" { xfail *-*-* } .-3 }
+ ! { dg-bogus "\[Ww\]arning: region is worker partitioned but does not contain worker partitioned code" "TODO default 'gang' 'vector'" { xfail *-*-* } .-3 }
integer :: i
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/par-reduction-2-1.f b/libgomp/testsuite/libgomp.oacc-fortran/par-reduction-2-1.f
index baa472b3e87..1fe1b471107 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/par-reduction-2-1.f
+++ b/libgomp/testsuite/libgomp.oacc-fortran/par-reduction-2-1.f
@@ -18,7 +18,7 @@
!$ACC PARALLEL NUM_GANGS(256) NUM_WORKERS(32) VECTOR_LENGTH(32)
!$ACC& REDUCTION(+:RES1) COPY(RES1, RES2) ASYNC(1)
-! { dg-bogus "Warning: region is gang partitioned but does not contain gang partitioned code" "TODO 'reduction', 'atomic'" { xfail *-*-* } .-1 }
+! { dg-bogus "\[Ww\]arning: region is gang partitioned but does not contain gang partitioned code" "TODO 'reduction', 'atomic'" { xfail *-*-* } .-1 }
! { dg-warning "region is worker partitioned but does not contain worker partitioned code" "" { target *-*-* } .-2 }
! { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .-3 }
res1 = res1 + 5
@@ -43,7 +43,7 @@
!$ACC PARALLEL NUM_GANGS(8) NUM_WORKERS(32) VECTOR_LENGTH(32)
!$ACC& REDUCTION(*:RES1) COPY(RES1, RES2) ASYNC(1)
-! { dg-bogus "Warning: region is gang partitioned but does not contain gang partitioned code" "TODO 'reduction', 'atomic'" { xfail *-*-* } .-1 }
+! { dg-bogus "\[Ww\]arning: region is gang partitioned but does not contain gang partitioned code" "TODO 'reduction', 'atomic'" { xfail *-*-* } .-1 }
! { dg-warning "region is worker partitioned but does not contain worker partitioned code" "" { target *-*-* } .-2 }
! { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .-3 }
res1 = res1 * 5
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/par-reduction-2-2.f b/libgomp/testsuite/libgomp.oacc-fortran/par-reduction-2-2.f
index 4ebab6c7b18..0221ab81281 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/par-reduction-2-2.f
+++ b/libgomp/testsuite/libgomp.oacc-fortran/par-reduction-2-2.f
@@ -18,7 +18,7 @@
!$ACC PARALLEL NUM_GANGS(256) NUM_WORKERS(32) VECTOR_LENGTH(32)
!$ACC& REDUCTION(+:RES1) COPY(RES1, RES2) ASYNC(1)
-! { dg-bogus "Warning: region is gang partitioned but does not contain gang partitioned code" "TODO 'reduction', 'atomic'" { xfail *-*-* } .-1 }
+! { dg-bogus "\[Ww\]arning: region is gang partitioned but does not contain gang partitioned code" "TODO 'reduction', 'atomic'" { xfail *-*-* } .-1 }
! { dg-warning "region is worker partitioned but does not contain worker partitioned code" "" { target *-*-* } .-2 }
! { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .-3 }
res1 = res1 + 5
@@ -43,7 +43,7 @@
!$ACC PARALLEL NUM_GANGS(8) NUM_WORKERS(32) VECTOR_LENGTH(32)
!$ACC& REDUCTION(*:RES1) COPY(RES1, RES2) ASYNC(1)
-! { dg-bogus "Warning: region is gang partitioned but does not contain gang partitioned code" "TODO 'reduction', 'atomic'" { xfail *-*-* } .-1 }
+! { dg-bogus "\[Ww\]arning: region is gang partitioned but does not contain gang partitioned code" "TODO 'reduction', 'atomic'" { xfail *-*-* } .-1 }
! { dg-warning "region is worker partitioned but does not contain worker partitioned code" "" { target *-*-* } .-2 }
! { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .-3 }
res1 = res1 * 5
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/parallel-dims.f90 b/libgomp/testsuite/libgomp.oacc-fortran/parallel-dims.f90
index d9847969369..f69ab5a6642 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/parallel-dims.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/parallel-dims.f90
@@ -87,9 +87,9 @@ program main
vectors_max = -huge(gangs_max) - 1 ! INT_MIN
!$acc serial copy (vectors_actual) &
!$acc copy (gangs_min, gangs_max, workers_min, workers_max, vectors_min, vectors_max) ! { dg-warning "using vector_length \\(32\\), ignoring 1" "" { target openacc_nvidia_accel_selected } }
- ! { dg-bogus "Warning: region contains gang partitioned code but is not gang partitioned" "TODO 'serial'" { xfail *-*-* } .-1 }
- ! { dg-bogus "Warning: region contains worker partitioned code but is not worker partitioned" "TODO 'serial'" { xfail *-*-* } .-2 }
- ! { dg-bogus "Warning: region contains vector partitioned code but is not vector partitioned" "TODO 'serial'" { xfail *-*-* } .-3 }
+ ! { dg-bogus "\[Ww\]arning: region contains gang partitioned code but is not gang partitioned" "TODO 'serial'" { xfail *-*-* } .-1 }
+ ! { dg-bogus "\[Ww\]arning: region contains worker partitioned code but is not worker partitioned" "TODO 'serial'" { xfail *-*-* } .-2 }
+ ! { dg-bogus "\[Ww\]arning: region contains vector partitioned code but is not vector partitioned" "TODO 'serial'" { xfail *-*-* } .-3 }
if (acc_on_device (acc_device_nvidia)) then
! The GCC nvptx back end enforces vector_length (32).
! It's unclear if that's actually permissible here;
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/parallel-reduction.f90 b/libgomp/testsuite/libgomp.oacc-fortran/parallel-reduction.f90
index af25f0dac9b..99b1441136d 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/parallel-reduction.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/parallel-reduction.f90
@@ -13,7 +13,7 @@ program reduction
s2 = 0
!$acc parallel reduction(+:s1,s2) num_gangs (n) copy(s1)
- ! { dg-bogus "Warning: region is gang partitioned but does not contain gang partitioned code" "TODO 'reduction'" { xfail *-*-* } .-1 }
+ ! { dg-bogus "\[Ww\]arning: region is gang partitioned but does not contain gang partitioned code" "TODO 'reduction'" { xfail *-*-* } .-1 }
s1 = s1 + 1
s2 = s2 + 1
!$acc end parallel
@@ -44,7 +44,7 @@ subroutine redsub(s1, s2, n)
integer :: s1, s2, n
!$acc parallel reduction(+:s1,s2) num_gangs (10) copy(s1)
- ! { dg-bogus "Warning: region is gang partitioned but does not contain gang partitioned code" "TODO 'reduction'" { xfail *-*-* } .-1 }
+ ! { dg-bogus "\[Ww\]arning: region is gang partitioned but does not contain gang partitioned code" "TODO 'reduction'" { xfail *-*-* } .-1 }
s1 = s1 + 1
s2 = s2 + 1
!$acc end parallel
^ permalink raw reply [flat|nested] 7+ messages in thread
end of thread, other threads:[~2021-04-26 21:23 UTC | newest]
Thread overview: 7+ messages (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2021-02-26 12:21 [PATCH] openacc: Warnings for strange partitioning choices for parallel regions Julian Brown
2021-04-26 10:39 ` Thomas Schwinge
2021-04-26 12:32 ` Tobias Burnus
2021-04-26 19:54 ` Thomas Schwinge
2021-04-26 20:51 ` Tobias Burnus
2021-04-26 21:23 ` [Patch, committed] OpenACC: Fix pattern in dg-bogus in Fortran testcases again (Re: [PATCH] openacc: Warnings for strange partitioning choices for parallel regions) Tobias Burnus
2021-04-26 20:00 ` [Patch, committed] OpenACC: Fix pattern in dg-bogus in Fortran testcases (was: " Tobias Burnus
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).