Hi!
On Thu, 21 Jan 2016 22:54:26 +0100, I wrote:
> Committed to gomp-4_0-branch in r232709:
>
> commit 41a76d233e714fd7b79dc1f40823f607c38306ba
> Author: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
> Date: Thu Jan 21 21:52:50 2016 +0000
>
> Un-parallelized OpenACC kernels constructs with nvptx offloading: "avoid
> offloading"
> +! The warning is only triggered for -O2 and higher.
> +! { dg-xfail-if "n/a" { openacc_nvidia_accel_selected } { "-O0" "-O1" } { ""
> } }
> +!$ACC KERNELS /* { dg-warning "OpenACC kernels construct will be executed
> sequentially; will by default avoid offloading to prevent data copy penalty"
> "" { target openacc_nvidia_accel_selected } } */
That can actually be done in a better way, so that we match up exactly
when and where the diagnostic appears, and where not; pushed the attached
to openacc-gcc-8-branch in commit
dfae8e0dab1edfc8d8207eafd1b694c4e1fcd680 'Un-parallelized OpenACC kernels
constructs with nvptx offloading: "avoid offloading": better method for
XFAILing specific cases'.
Grüße
Thomas
>From dfae8e0dab1edfc8d8207eafd1b694c4e1fcd680 Mon Sep 17 00:00:00 2001
From: Thomas Schwinge <tho...@codesourcery.com>
Date: Tue, 29 Jan 2019 20:42:13 +0100
Subject: [PATCH] Un-parallelized OpenACC kernels constructs with nvptx
offloading: "avoid offloading": better method for XFAILing specific cases
gcc/testsuite/
* lib/target-supports.exp
(check_effective_target_opt_levels_2_plus)
(check_effective_target_opt_levels_size): New.
libgomp/
* testsuite/libgomp.oacc-c-c++-common/avoid-offloading-1.c:
Update.
* testsuite/libgomp.oacc-c-c++-common/avoid-offloading-2.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/combined-directives-1.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/kernels-parallel-loop-data-enter-exit.c:
Likewise.
* testsuite/libgomp.oacc-fortran/asyncwait-1.f90: Likewise.
* testsuite/libgomp.oacc-fortran/avoid-offloading-1.f: Likewise.
* testsuite/libgomp.oacc-fortran/avoid-offloading-2.f: Likewise.
* testsuite/libgomp.oacc-fortran/deviceptr-1.f90: Likewise.
* testsuite/libgomp.oacc-fortran/initialize_kernels_loops.f90:
Likewise.
* testsuite/libgomp.oacc-fortran/kernels-loop-2.f95: Likewise.
* testsuite/libgomp.oacc-fortran/kernels-loop-data-2.f95:
Likewise.
* testsuite/libgomp.oacc-fortran/kernels-loop-data-enter-exit-2.f95:
Likewise.
* testsuite/libgomp.oacc-fortran/kernels-loop-data-enter-exit.f95:
Likewise.
* testsuite/libgomp.oacc-fortran/kernels-loop-data-update.f95:
Likewise.
* testsuite/libgomp.oacc-fortran/kernels-loop-data.f95: Likewise.
* testsuite/libgomp.oacc-fortran/kernels-loop.f95: Likewise.
* testsuite/libgomp.oacc-fortran/kernels-parallel-loop-data-enter-exit.f95:
Likewise.
* testsuite/libgomp.oacc-fortran/non-scalar-data.f90: Likewise.
---
gcc/testsuite/ChangeLog.openacc | 6 ++++
gcc/testsuite/lib/target-supports.exp | 10 ++++++
libgomp/ChangeLog.openacc | 31 +++++++++++++++++++
.../avoid-offloading-1.c | 5 +--
.../avoid-offloading-2.c | 5 +--
.../combined-directives-1.c | 10 +++---
.../kernels-parallel-loop-data-enter-exit.c | 8 ++---
.../libgomp.oacc-fortran/asyncwait-1.f90 | 9 ++----
.../libgomp.oacc-fortran/avoid-offloading-1.f | 4 +--
.../libgomp.oacc-fortran/avoid-offloading-2.f | 4 +--
.../libgomp.oacc-fortran/deviceptr-1.f90 | 7 ++---
.../initialize_kernels_loops.f90 | 5 +--
.../libgomp.oacc-fortran/kernels-loop-2.f95 | 9 ++----
.../kernels-loop-data-2.f95 | 9 ++----
.../kernels-loop-data-enter-exit-2.f95 | 9 ++----
.../kernels-loop-data-enter-exit.f95 | 9 ++----
.../kernels-loop-data-update.f95 | 7 ++---
.../kernels-loop-data.f95 | 9 ++----
.../libgomp.oacc-fortran/kernels-loop.f95 | 5 +--
.../kernels-parallel-loop-data-enter-exit.f95 | 7 ++---
.../libgomp.oacc-fortran/non-scalar-data.f90 | 7 ++---
21 files changed, 86 insertions(+), 89 deletions(-)
diff --git a/gcc/testsuite/ChangeLog.openacc b/gcc/testsuite/ChangeLog.openacc
index 41913f7fa02..2479367dce4 100644
--- a/gcc/testsuite/ChangeLog.openacc
+++ b/gcc/testsuite/ChangeLog.openacc
@@ -1,3 +1,9 @@
+2019-01-31 Thomas Schwinge <tho...@codesourcery.com>
+
+ * lib/target-supports.exp
+ (check_effective_target_opt_levels_2_plus)
+ (check_effective_target_opt_levels_size): New.
+
2019-01-31 Julian Brown <jul...@codesourcery.com>
* c-c++-common/goacc/deep-copy-arrayofstruct.c: New test.
diff --git a/gcc/testsuite/lib/target-supports.exp b/gcc/testsuite/lib/target-supports.exp
index cfc22a22975..206bb2a2b11 100644
--- a/gcc/testsuite/lib/target-supports.exp
+++ b/gcc/testsuite/lib/target-supports.exp
@@ -9326,3 +9326,13 @@ proc check_effective_target_cet { } {
}
} "-O2" ]
}
+
+# Return 1 if we have optimization level -O2 or higher.
+proc check_effective_target_opt_levels_2_plus { } {
+ return [check-flags [list "" { *-*-* } { "-O*" } { "-O0" "-Og" "-O" "-O1" }]];
+}
+
+# Return 1 if we have optimization level -Os.
+proc check_effective_target_opt_levels_size { } {
+ return [check-flags [list "" { *-*-* } { "-Os" } { "" }]];
+}
diff --git a/libgomp/ChangeLog.openacc b/libgomp/ChangeLog.openacc
index e638c259cba..d558f0df97b 100644
--- a/libgomp/ChangeLog.openacc
+++ b/libgomp/ChangeLog.openacc
@@ -1,3 +1,34 @@
+2019-01-31 Thomas Schwinge <tho...@codesourcery.com>
+
+ * testsuite/libgomp.oacc-c-c++-common/avoid-offloading-1.c:
+ Update.
+ * testsuite/libgomp.oacc-c-c++-common/avoid-offloading-2.c:
+ Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/combined-directives-1.c:
+ Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/kernels-parallel-loop-data-enter-exit.c:
+ Likewise.
+ * testsuite/libgomp.oacc-fortran/asyncwait-1.f90: Likewise.
+ * testsuite/libgomp.oacc-fortran/avoid-offloading-1.f: Likewise.
+ * testsuite/libgomp.oacc-fortran/avoid-offloading-2.f: Likewise.
+ * testsuite/libgomp.oacc-fortran/deviceptr-1.f90: Likewise.
+ * testsuite/libgomp.oacc-fortran/initialize_kernels_loops.f90:
+ Likewise.
+ * testsuite/libgomp.oacc-fortran/kernels-loop-2.f95: Likewise.
+ * testsuite/libgomp.oacc-fortran/kernels-loop-data-2.f95:
+ Likewise.
+ * testsuite/libgomp.oacc-fortran/kernels-loop-data-enter-exit-2.f95:
+ Likewise.
+ * testsuite/libgomp.oacc-fortran/kernels-loop-data-enter-exit.f95:
+ Likewise.
+ * testsuite/libgomp.oacc-fortran/kernels-loop-data-update.f95:
+ Likewise.
+ * testsuite/libgomp.oacc-fortran/kernels-loop-data.f95: Likewise.
+ * testsuite/libgomp.oacc-fortran/kernels-loop.f95: Likewise.
+ * testsuite/libgomp.oacc-fortran/kernels-parallel-loop-data-enter-exit.f95:
+ Likewise.
+ * testsuite/libgomp.oacc-fortran/non-scalar-data.f90: Likewise.
+
2019-01-31 Julian Brown <jul...@codesourcery.com>
* testsuite/libgomp.oacc-c++/deep-copy-12.C: New test.
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/avoid-offloading-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/avoid-offloading-1.c
index d5fff2d7d8a..72b9ce0ce02 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/avoid-offloading-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/avoid-offloading-1.c
@@ -1,8 +1,5 @@
/* Test that the compiler decides to "avoid offloading". */
-/* The warning is only triggered for -O2 and higher.
- { dg-xfail-if "n/a" { openacc_nvidia_accel_selected } { "-O0" "-O1" } { "" } } */
-
#include <openacc.h>
int main(void)
@@ -10,7 +7,7 @@ int main(void)
int x, y;
#pragma acc data copyout(x, y)
-#pragma acc kernels /* { dg-warning "OpenACC kernels construct will be executed sequentially; will by default avoid offloading to prevent data copy penalty" "" { target openacc_nvidia_accel_selected } } */
+#pragma acc kernels /* { dg-warning "OpenACC kernels construct will be executed sequentially; will by default avoid offloading to prevent data copy penalty" "" { target { openacc_nvidia_accel_selected && opt_levels_2_plus } } } */
*((volatile int *) &x) = 33, y = acc_on_device (acc_device_host);
if (x != 33)
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/avoid-offloading-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/avoid-offloading-2.c
index 41bd6d55f26..9e05d84d792 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/avoid-offloading-2.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/avoid-offloading-2.c
@@ -1,9 +1,6 @@
/* Test that a user can override the compiler's "avoid offloading"
decision. */
-/* The warning is only triggered for -O2 and higher.
- { dg-xfail-if "n/a" { openacc_nvidia_accel_selected } { "-O0" "-O1" } { "" } } */
-
#include <openacc.h>
int main(void)
@@ -22,7 +19,7 @@ int main(void)
int x, y;
#pragma acc data copyout(x, y)
-#pragma acc kernels /* { dg-warning "OpenACC kernels construct will be executed sequentially; will by default avoid offloading to prevent data copy penalty" "" { target openacc_nvidia_accel_selected } } */
+#pragma acc kernels /* { dg-warning "OpenACC kernels construct will be executed sequentially; will by default avoid offloading to prevent data copy penalty" "" { target { openacc_nvidia_accel_selected && opt_levels_2_plus } } } */
*((volatile int *) &x) = 33, y = acc_on_device (acc_device_host);
if (x != 33)
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/combined-directives-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/combined-directives-1.c
index c6abc1d724a..a54ade363a6 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/combined-directives-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/combined-directives-1.c
@@ -1,11 +1,6 @@
/* This test exercises combined directives. */
-/* This test falls back to host execution because struct alias
- analysis is deactivated on OpenACC parallel regions. Consequently,
- parloops can no longer disambiguate arrays a and b. */
-
/* { dg-do run } */
-/* { dg-xfail-if "n/a" { openacc_nvidia_accel_selected } { "-O2" } { "" } } */
#include <stdlib.h>
@@ -38,7 +33,10 @@ main (int argc, char **argv)
abort ();
}
-#pragma acc kernels loop copy (a[0:N]) copy (b[0:N])
+#pragma acc kernels loop copy (a[0:N]) copy (b[0:N]) /* { dg-bogus "OpenACC kernels construct will be executed sequentially; will by default avoid offloading to prevent data copy penalty" "TODO" { xfail { openacc_nvidia_accel_selected && opt_levels_2_plus } } }
+ This runs into "avoid offloading" because struct alias analysis is
+ deactivated on OpenACC parallel regions. Consequently, parloops can no
+ longer disambiguate arrays a and b. */
for (i = 0; i < N; i++)
{
b[i] = 3.0;
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-parallel-loop-data-enter-exit.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-parallel-loop-data-enter-exit.c
index 8cafbc974c9..374014a1e86 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-parallel-loop-data-enter-exit.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-parallel-loop-data-enter-exit.c
@@ -1,7 +1,3 @@
-/* FIXME: OpenACC kernels stopped working with the firstprivate subarray
- changes. */
-/* { dg-prune-output "OpenACC kernels construct will be executed sequentially" } */
-
#include <stdlib.h>
#define N (1024 * 512)
@@ -33,7 +29,9 @@ main (void)
b[i] = i * 4;
}
-#pragma acc kernels present (a[0:N], b[0:N], c[0:N])
+#pragma acc kernels present (a[0:N], b[0:N], c[0:N]) /* { dg-bogus "OpenACC kernels construct will be executed sequentially; will by default avoid offloading to prevent data copy penalty" "TODO" { xfail { openacc_nvidia_accel_selected && opt_levels_2_plus } } }
+ FIXME: OpenACC kernels stopped working with the firstprivate subarray
+ changes. */
{
for (COUNTERTYPE ii = 0; ii < N; ii++)
c[ii] = a[ii] + b[ii];
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/asyncwait-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/asyncwait-1.f90
index f024c1cbe51..ba6809dba7d 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/asyncwait-1.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/asyncwait-1.f90
@@ -1,7 +1,4 @@
! { dg-do run }
-! TODO, <https://gcc.gnu.org/PR80995>.
-! warning: OpenACC kernels construct will be executed sequentially; will by default avoid offloading to prevent data copy penalty
-! { dg-xfail-if "TODO" { openacc_nvidia_accel_selected } { "-Os" } { "" } }
program asyncwait
integer, parameter :: N = 64
@@ -183,13 +180,13 @@ program asyncwait
!$acc data copy (a(1:N)) copy (b(1:N)) copy (c(1:N)) copy (d(1:N))
- !$acc kernels async (1)
+ !$acc kernels async (1) ! { dg-bogus "OpenACC kernels construct will be executed sequentially; will by default avoid offloading to prevent data copy penalty" "PR80995" { xfail { openacc_nvidia_accel_selected && opt_levels_size } } }
do i = 1, N
b(i) = (a(i) * a(i) * a(i)) / a(i)
end do
!$acc end kernels
- !$acc kernels async (1)
+ !$acc kernels async (1) ! { dg-bogus "OpenACC kernels construct will be executed sequentially; will by default avoid offloading to prevent data copy penalty" "PR80995" { xfail { openacc_nvidia_accel_selected && opt_levels_size } } }
do i = 1, N
c(i) = (a(i) * 4) / a(i)
end do
@@ -220,7 +217,7 @@ program asyncwait
!$acc data copy (a(1:N), b(1:N), c(1:N), d(1:N), e(1:N))
- !$acc kernels async (1)
+ !$acc kernels async (1) ! { dg-bogus "OpenACC kernels construct will be executed sequentially; will by default avoid offloading to prevent data copy penalty" "PR80995" { xfail { openacc_nvidia_accel_selected && opt_levels_size } } }
do i = 1, N
b(i) = (a(i) * a(i) * a(i)) / a(i)
end do
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/avoid-offloading-1.f b/libgomp/testsuite/libgomp.oacc-fortran/avoid-offloading-1.f
index da89b93fd54..fb14be19d8d 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/avoid-offloading-1.f
+++ b/libgomp/testsuite/libgomp.oacc-fortran/avoid-offloading-1.f
@@ -2,8 +2,6 @@
! { dg-do run }
! { dg-additional-options "-cpp" }
-! The warning is only triggered for -O2 and higher.
-! { dg-xfail-if "n/a" { openacc_nvidia_accel_selected } { "-O0" "-O1" } { "" } }
! As __OPTIMIZE__ is defined for -O1 and higher, we don't have an (easy) way to
! distinguish -O1 (where we will offload) from -O2 (where we won't offload), so
! for -O1 testing, we expect to abort.
@@ -16,7 +14,7 @@
LOGICAL :: Y
!$ACC DATA COPYOUT(X, Y)
-!$ACC KERNELS /* { dg-warning "OpenACC kernels construct will be executed sequentially; will by default avoid offloading to prevent data copy penalty" "" { target openacc_nvidia_accel_selected } } */
+!$ACC KERNELS ! { dg-warning "OpenACC kernels construct will be executed sequentially; will by default avoid offloading to prevent data copy penalty" "" { target { openacc_nvidia_accel_selected && opt_levels_2_plus } } }
X = 33
Y = ACC_ON_DEVICE (ACC_DEVICE_HOST);
!$ACC END KERNELS
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/avoid-offloading-2.f b/libgomp/testsuite/libgomp.oacc-fortran/avoid-offloading-2.f
index db72602fb1e..5a064618e51 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/avoid-offloading-2.f
+++ b/libgomp/testsuite/libgomp.oacc-fortran/avoid-offloading-2.f
@@ -2,8 +2,6 @@
! { dg-do run }
! { dg-additional-options "-cpp" }
-! The warning is only triggered for -O2 and higher.
-! { dg-xfail-if "n/a" { openacc_nvidia_accel_selected } { "-O0" "-O1" } { "" } }
IMPLICIT NONE
INCLUDE "openacc_lib.h"
@@ -23,7 +21,7 @@
CALL ACC_INIT (D)
!$ACC DATA COPYOUT(X, Y)
-!$ACC KERNELS /* { dg-warning "OpenACC kernels construct will be executed sequentially; will by default avoid offloading to prevent data copy penalty" "" { target openacc_nvidia_accel_selected } } */
+!$ACC KERNELS ! { dg-warning "OpenACC kernels construct will be executed sequentially; will by default avoid offloading to prevent data copy penalty" "" { target { openacc_nvidia_accel_selected && opt_levels_2_plus } } }
X = 33
Y = ACC_ON_DEVICE (ACC_DEVICE_HOST)
!$ACC END KERNELS
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/deviceptr-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/deviceptr-1.f90
index 73838510ea4..418bceac3b9 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/deviceptr-1.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/deviceptr-1.f90
@@ -3,9 +3,6 @@
! the deviceptr variable is implied.
! { dg-do run }
-! TODO, <https://gcc.gnu.org/PR80995>.
-! warning: OpenACC kernels construct will be executed sequentially; will by default avoid offloading to prevent data copy penalty
-! { dg-xfail-if "TODO" { openacc_nvidia_accel_selected } { "-Os" } { "" } }
subroutine subr1 (a, b)
implicit none
@@ -52,7 +49,7 @@ subroutine subr3 (a, b)
integer :: b(N)
integer :: i = 0
- !$acc kernels copy (b)
+ !$acc kernels copy (b) ! { dg-bogus "OpenACC kernels construct will be executed sequentially; will by default avoid offloading to prevent data copy penalty" "PR80995" { xfail { openacc_nvidia_accel_selected && opt_levels_size } } }
do i = 1, N
a(i) = i * 8
b(i) = a(i)
@@ -84,7 +81,7 @@ subroutine subr5 (a, b)
integer :: b(N)
integer :: i = 0
- !$acc kernels deviceptr (a) copy (b)
+ !$acc kernels deviceptr (a) copy (b) ! { dg-bogus "OpenACC kernels construct will be executed sequentially; will by default avoid offloading to prevent data copy penalty" "PR80995" { xfail { openacc_nvidia_accel_selected && opt_levels_size } } }
do i = 1, N
a(i) = i * 32
b(i) = a(i)
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/initialize_kernels_loops.f90 b/libgomp/testsuite/libgomp.oacc-fortran/initialize_kernels_loops.f90
index 6d1713157b7..8eb02b88d25 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/initialize_kernels_loops.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/initialize_kernels_loops.f90
@@ -1,14 +1,11 @@
! { dg-do run }
-!TODO
-! warning: OpenACC kernels construct will be executed sequentially; will by default avoid offloading to prevent data copy penalty
-! { dg-xfail-if "TODO" { openacc_nvidia_accel_selected } { "*" } { "-O0" "-O1" } }
subroutine kernel(lo, hi, a, b, c)
implicit none
integer :: lo, hi, i
real, dimension(lo:hi) :: a, b, c
-!$acc kernels
+!$acc kernels ! { dg-bogus "OpenACC kernels construct will be executed sequentially; will by default avoid offloading to prevent data copy penalty" "TODO" { xfail { openacc_nvidia_accel_selected && opt_levels_2_plus } } }
!$acc loop independent
do i = lo, hi
b(i) = a(i)
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/kernels-loop-2.f95 b/libgomp/testsuite/libgomp.oacc-fortran/kernels-loop-2.f95
index af491d053e6..8becc159dd1 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/kernels-loop-2.f95
+++ b/libgomp/testsuite/libgomp.oacc-fortran/kernels-loop-2.f95
@@ -1,7 +1,4 @@
! { dg-do run }
-! TODO, <https://gcc.gnu.org/PR80995>.
-! warning: OpenACC kernels construct will be executed sequentially; will by default avoid offloading to prevent data copy penalty
-! { dg-xfail-if "TODO" { openacc_nvidia_accel_selected } { "-Os" } { "" } }
program main
implicit none
@@ -11,7 +8,7 @@ program main
! Parallelism dimensions: compiler/runtime decides.
!$acc kernels copyout (a(0:n-1))
- do i = 0, n - 1
+ do i = 0, n - 1 ! { dg-bogus "OpenACC kernels construct will be executed sequentially; will by default avoid offloading to prevent data copy penalty" "PR80995" { xfail { openacc_nvidia_accel_selected && opt_levels_size } } }
a(i) = i * 2
end do
!$acc end kernels
@@ -20,7 +17,7 @@ program main
!$acc kernels copyout (b(0:n-1)) &
!$acc num_gangs (3 + a(3)) num_workers (5 + a(5)) vector_length (7 + a(7))
! { dg-prune-output "using vector_length \\(32\\), ignoring runtime setting" }
- do i = 0, n -1
+ do i = 0, n -1 ! { dg-bogus "OpenACC kernels construct will be executed sequentially; will by default avoid offloading to prevent data copy penalty" "PR80995" { xfail { openacc_nvidia_accel_selected && opt_levels_size } } }
b(i) = i * 4
end do
!$acc end kernels
@@ -29,7 +26,7 @@ program main
!$acc kernels copyin (a(0:n-1), b(0:n-1)) copyout (c(0:n-1)) &
!$acc num_gangs (3) num_workers (5) vector_length (7)
! { dg-prune-output "using vector_length \\(32\\), ignoring 7" }
- do ii = 0, n - 1
+ do ii = 0, n - 1 ! { dg-bogus "OpenACC kernels construct will be executed sequentially; will by default avoid offloading to prevent data copy penalty" "PR80995" { xfail { openacc_nvidia_accel_selected && opt_levels_size } } }
c(ii) = a(ii) + b(ii)
end do
!$acc end kernels
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/kernels-loop-data-2.f95 b/libgomp/testsuite/libgomp.oacc-fortran/kernels-loop-data-2.f95
index ca1ac70cd9e..2191ebedee3 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/kernels-loop-data-2.f95
+++ b/libgomp/testsuite/libgomp.oacc-fortran/kernels-loop-data-2.f95
@@ -1,7 +1,4 @@
! { dg-do run }
-! TODO, <https://gcc.gnu.org/PR80995>.
-! warning: OpenACC kernels construct will be executed sequentially; will by default avoid offloading to prevent data copy penalty
-! { dg-xfail-if "TODO" { openacc_nvidia_accel_selected } { "-Os" } { "" } }
program main
implicit none
@@ -11,7 +8,7 @@ program main
!$acc data copyout (a(0:n-1))
!$acc kernels present (a(0:n-1))
- do i = 0, n - 1
+ do i = 0, n - 1 ! { dg-bogus "OpenACC kernels construct will be executed sequentially; will by default avoid offloading to prevent data copy penalty" "PR80995" { xfail { openacc_nvidia_accel_selected && opt_levels_size } } }
a(i) = i * 2
end do
!$acc end kernels
@@ -19,7 +16,7 @@ program main
!$acc data copyout (b(0:n-1))
!$acc kernels present (b(0:n-1))
- do i = 0, n -1
+ do i = 0, n -1 ! { dg-bogus "OpenACC kernels construct will be executed sequentially; will by default avoid offloading to prevent data copy penalty" "PR80995" { xfail { openacc_nvidia_accel_selected && opt_levels_size } } }
b(i) = i * 4
end do
!$acc end kernels
@@ -27,7 +24,7 @@ program main
!$acc data copyin (a(0:n-1), b(0:n-1)) copyout (c(0:n-1))
!$acc kernels present (a(0:n-1), b(0:n-1), c(0:n-1))
- do ii = 0, n - 1
+ do ii = 0, n - 1 ! { dg-bogus "OpenACC kernels construct will be executed sequentially; will by default avoid offloading to prevent data copy penalty" "PR80995" { xfail { openacc_nvidia_accel_selected && opt_levels_size } } }
c(ii) = a(ii) + b(ii)
end do
!$acc end kernels
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/kernels-loop-data-enter-exit-2.f95 b/libgomp/testsuite/libgomp.oacc-fortran/kernels-loop-data-enter-exit-2.f95
index 5103f3ede04..75fb8a32beb 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/kernels-loop-data-enter-exit-2.f95
+++ b/libgomp/testsuite/libgomp.oacc-fortran/kernels-loop-data-enter-exit-2.f95
@@ -1,7 +1,4 @@
! { dg-do run }
-! TODO, <https://gcc.gnu.org/PR80995>.
-! warning: OpenACC kernels construct will be executed sequentially; will by default avoid offloading to prevent data copy penalty
-! { dg-xfail-if "TODO" { openacc_nvidia_accel_selected } { "-Os" } { "" } }
program main
implicit none
@@ -11,7 +8,7 @@ program main
!$acc enter data create (a(0:n-1))
!$acc kernels present (a(0:n-1))
- do i = 0, n - 1
+ do i = 0, n - 1 ! { dg-bogus "OpenACC kernels construct will be executed sequentially; will by default avoid offloading to prevent data copy penalty" "PR80995" { xfail { openacc_nvidia_accel_selected && opt_levels_size } } }
a(i) = i * 2
end do
!$acc end kernels
@@ -19,7 +16,7 @@ program main
!$acc enter data create (b(0:n-1))
!$acc kernels present (b(0:n-1))
- do i = 0, n -1
+ do i = 0, n -1 ! { dg-bogus "OpenACC kernels construct will be executed sequentially; will by default avoid offloading to prevent data copy penalty" "PR80995" { xfail { openacc_nvidia_accel_selected && opt_levels_size } } }
b(i) = i * 4
end do
!$acc end kernels
@@ -27,7 +24,7 @@ program main
!$acc enter data copyin (a(0:n-1), b(0:n-1)) create (c(0:n-1))
!$acc kernels present (a(0:n-1), b(0:n-1), c(0:n-1))
- do ii = 0, n - 1
+ do ii = 0, n - 1 ! { dg-bogus "OpenACC kernels construct will be executed sequentially; will by default avoid offloading to prevent data copy penalty" "PR80995" { xfail { openacc_nvidia_accel_selected && opt_levels_size } } }
c(ii) = a(ii) + b(ii)
end do
!$acc end kernels
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/kernels-loop-data-enter-exit.f95 b/libgomp/testsuite/libgomp.oacc-fortran/kernels-loop-data-enter-exit.f95
index 5c1fd52c956..8ea34bf7bf8 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/kernels-loop-data-enter-exit.f95
+++ b/libgomp/testsuite/libgomp.oacc-fortran/kernels-loop-data-enter-exit.f95
@@ -1,7 +1,4 @@
! { dg-do run }
-! TODO, <https://gcc.gnu.org/PR80995>.
-! warning: OpenACC kernels construct will be executed sequentially; will by default avoid offloading to prevent data copy penalty
-! { dg-xfail-if "TODO" { openacc_nvidia_accel_selected } { "-Os" } { "" } }
program main
implicit none
@@ -12,19 +9,19 @@ program main
!$acc enter data create (a(0:n-1), b(0:n-1), c(0:n-1))
!$acc kernels present (a(0:n-1))
- do i = 0, n - 1
+ do i = 0, n - 1 ! { dg-bogus "OpenACC kernels construct will be executed sequentially; will by default avoid offloading to prevent data copy penalty" "PR80995" { xfail { openacc_nvidia_accel_selected && opt_levels_size } } }
a(i) = i * 2
end do
!$acc end kernels
!$acc kernels present (b(0:n-1))
- do i = 0, n -1
+ do i = 0, n -1 ! { dg-bogus "OpenACC kernels construct will be executed sequentially; will by default avoid offloading to prevent data copy penalty" "PR80995" { xfail { openacc_nvidia_accel_selected && opt_levels_size } } }
b(i) = i * 4
end do
!$acc end kernels
!$acc kernels present (a(0:n-1), b(0:n-1), c(0:n-1))
- do ii = 0, n - 1
+ do ii = 0, n - 1 ! { dg-bogus "OpenACC kernels construct will be executed sequentially; will by default avoid offloading to prevent data copy penalty" "PR80995" { xfail { openacc_nvidia_accel_selected && opt_levels_size } } }
c(ii) = a(ii) + b(ii)
end do
!$acc end kernels
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/kernels-loop-data-update.f95 b/libgomp/testsuite/libgomp.oacc-fortran/kernels-loop-data-update.f95
index 3f889a41a0f..710068a707a 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/kernels-loop-data-update.f95
+++ b/libgomp/testsuite/libgomp.oacc-fortran/kernels-loop-data-update.f95
@@ -1,7 +1,4 @@
! { dg-do run }
-! TODO, <https://gcc.gnu.org/PR80995>.
-! warning: OpenACC kernels construct will be executed sequentially; will by default avoid offloading to prevent data copy penalty
-! { dg-xfail-if "TODO" { openacc_nvidia_accel_selected } { "-Os" } { "" } }
program main
implicit none
@@ -12,7 +9,7 @@ program main
!$acc enter data create (a(0:n-1), b(0:n-1), c(0:n-1))
!$acc kernels present (a(0:n-1))
- do i = 0, n - 1
+ do i = 0, n - 1 ! { dg-bogus "OpenACC kernels construct will be executed sequentially; will by default avoid offloading to prevent data copy penalty" "PR80995" { xfail { openacc_nvidia_accel_selected && opt_levels_size } } }
a(i) = i * 2
end do
!$acc end kernels
@@ -24,7 +21,7 @@ program main
!$acc update device (b(0:n-1))
!$acc kernels present (a(0:n-1), b(0:n-1), c(0:n-1))
- do ii = 0, n - 1
+ do ii = 0, n - 1 ! { dg-bogus "OpenACC kernels construct will be executed sequentially; will by default avoid offloading to prevent data copy penalty" "PR80995" { xfail { openacc_nvidia_accel_selected && opt_levels_size } } }
c(ii) = a(ii) + b(ii)
end do
!$acc end kernels
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/kernels-loop-data.f95 b/libgomp/testsuite/libgomp.oacc-fortran/kernels-loop-data.f95
index e495663a72e..c1dec2c89c5 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/kernels-loop-data.f95
+++ b/libgomp/testsuite/libgomp.oacc-fortran/kernels-loop-data.f95
@@ -1,7 +1,4 @@
! { dg-do run }
-! TODO, <https://gcc.gnu.org/PR80995>.
-! warning: OpenACC kernels construct will be executed sequentially; will by default avoid offloading to prevent data copy penalty
-! { dg-xfail-if "TODO" { openacc_nvidia_accel_selected } { "-Os" } { "" } }
program main
implicit none
@@ -12,19 +9,19 @@ program main
!$acc data copyout (a(0:n-1), b(0:n-1), c(0:n-1))
!$acc kernels present (a(0:n-1))
- do i = 0, n - 1
+ do i = 0, n - 1 ! { dg-bogus "OpenACC kernels construct will be executed sequentially; will by default avoid offloading to prevent data copy penalty" "PR80995" { xfail { openacc_nvidia_accel_selected && opt_levels_size } } }
a(i) = i * 2
end do
!$acc end kernels
!$acc kernels present (b(0:n-1))
- do i = 0, n -1
+ do i = 0, n -1 ! { dg-bogus "OpenACC kernels construct will be executed sequentially; will by default avoid offloading to prevent data copy penalty" "PR80995" { xfail { openacc_nvidia_accel_selected && opt_levels_size } } }
b(i) = i * 4
end do
!$acc end kernels
!$acc kernels present (a(0:n-1), b(0:n-1), c(0:n-1))
- do ii = 0, n - 1
+ do ii = 0, n - 1 ! { dg-bogus "OpenACC kernels construct will be executed sequentially; will by default avoid offloading to prevent data copy penalty" "PR80995" { xfail { openacc_nvidia_accel_selected && opt_levels_size } } }
c(ii) = a(ii) + b(ii)
end do
!$acc end kernels
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/kernels-loop.f95 b/libgomp/testsuite/libgomp.oacc-fortran/kernels-loop.f95
index 7377ed8780a..c9d3c4adc95 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/kernels-loop.f95
+++ b/libgomp/testsuite/libgomp.oacc-fortran/kernels-loop.f95
@@ -1,7 +1,4 @@
! { dg-do run }
-! TODO, <https://gcc.gnu.org/PR80995>.
-! warning: OpenACC kernels construct will be executed sequentially; will by default avoid offloading to prevent data copy penalty
-! { dg-xfail-if "TODO" { openacc_nvidia_accel_selected } { "-Os" } { "" } }
program main
implicit none
@@ -18,7 +15,7 @@ program main
end do
!$acc kernels copyin (a(0:n-1), b(0:n-1)) copyout (c(0:n-1))
- do ii = 0, n - 1
+ do ii = 0, n - 1 ! { dg-bogus "OpenACC kernels construct will be executed sequentially; will by default avoid offloading to prevent data copy penalty" "PR80995" { xfail { openacc_nvidia_accel_selected && opt_levels_size } } }
c(ii) = a(ii) + b(ii)
end do
!$acc end kernels
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/kernels-parallel-loop-data-enter-exit.f95 b/libgomp/testsuite/libgomp.oacc-fortran/kernels-parallel-loop-data-enter-exit.f95
index 8685275046b..99300ec88b1 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/kernels-parallel-loop-data-enter-exit.f95
+++ b/libgomp/testsuite/libgomp.oacc-fortran/kernels-parallel-loop-data-enter-exit.f95
@@ -1,7 +1,4 @@
! { dg-do run }
-! TODO, <https://gcc.gnu.org/PR80995>.
-! warning: OpenACC kernels construct will be executed sequentially; will by default avoid offloading to prevent data copy penalty
-! { dg-xfail-if "TODO" { openacc_nvidia_accel_selected } { "-Os" } { "" } }
program main
implicit none
@@ -12,7 +9,7 @@ program main
!$acc enter data create (a(0:n-1), b(0:n-1), c(0:n-1))
!$acc kernels present (a(0:n-1))
- do i = 0, n - 1
+ do i = 0, n - 1 ! { dg-bogus "OpenACC kernels construct will be executed sequentially; will by default avoid offloading to prevent data copy penalty" "PR80995" { xfail { openacc_nvidia_accel_selected && opt_levels_size } } }
a(i) = i * 2
end do
!$acc end kernels
@@ -25,7 +22,7 @@ program main
!$acc end parallel
!$acc kernels present (a(0:n-1), b(0:n-1), c(0:n-1))
- do ii = 0, n - 1
+ do ii = 0, n - 1 ! { dg-bogus "OpenACC kernels construct will be executed sequentially; will by default avoid offloading to prevent data copy penalty" "PR80995" { xfail { openacc_nvidia_accel_selected && opt_levels_size } } }
c(ii) = a(ii) + b(ii)
end do
!$acc end kernels
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/non-scalar-data.f90 b/libgomp/testsuite/libgomp.oacc-fortran/non-scalar-data.f90
index 99bd69207b6..f037b9ac9fd 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/non-scalar-data.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/non-scalar-data.f90
@@ -3,9 +3,6 @@
! present.
! { dg-do run }
-! TODO, <https://gcc.gnu.org/PR80995>.
-! warning: OpenACC kernels construct will be executed sequentially; will by default avoid offloading to prevent data copy penalty
-! { dg-xfail-if "TODO" { openacc_nvidia_accel_selected } { "-Os" } { "" } }
program main
implicit none
@@ -54,7 +51,7 @@ subroutine kernels (array, n)
integer, dimension (n) :: array
integer :: n, i
- !$acc kernels
+ !$acc kernels ! { dg-bogus "OpenACC kernels construct will be executed sequentially; will by default avoid offloading to prevent data copy penalty" "PR80995" { xfail { openacc_nvidia_accel_selected && opt_levels_size } } }
do i = 1, n
array(i) = i
end do
@@ -65,7 +62,7 @@ subroutine kernels_default_present (array, n)
integer, dimension (n) :: array
integer :: n, i
- !$acc kernels default(present)
+ !$acc kernels default(present) ! { dg-bogus "OpenACC kernels construct will be executed sequentially; will by default avoid offloading to prevent data copy penalty" "PR80995" { xfail { openacc_nvidia_accel_selected && opt_levels_size } } }
do i = 1, n
array(i) = i+1
end do
--
2.17.1
