[PATCH] openacc: Warnings for strange partitioning choices for parallel regions
Julian Brown
julian@codesourcery.com
Fri Feb 26 12:21:54 GMT 2021
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
More information about the Gcc-patches
mailing list