public inbox for gcc-cvs@sourceware.org
help / color / mirror / Atom feed
* [gcc r14-4910] Extend test suite coverage for OpenACC 'self' clause for compute constructs
@ 2023-10-25 9:25 Thomas Schwinge
0 siblings, 0 replies; only message in thread
From: Thomas Schwinge @ 2023-10-25 9:25 UTC (permalink / raw)
To: gcc-cvs
https://gcc.gnu.org/g:047841a68ebf5f991e842961f9e54f3c10b94f2c
commit r14-4910-g047841a68ebf5f991e842961f9e54f3c10b94f2c
Author: Thomas Schwinge <thomas@codesourcery.com>
Date: Mon Oct 23 14:53:29 2023 +0200
Extend test suite coverage for OpenACC 'self' clause for compute constructs
... on top of what was provided in recent
commit 3a3596389c2e539cb8fd5dc5784a4e2afe193a2a
"OpenACC 2.7: Implement self clause for compute constructs".
gcc/testsuite/
* c-c++-common/goacc/if-clause-2.c: Enhance.
* c-c++-common/goacc/self-clause-1.c: Likewise.
* c-c++-common/goacc/self-clause-2.c: Likewise.
* gfortran.dg/goacc/if.f95: Likewise.
* gfortran.dg/goacc/kernels-tree.f95: Likewise.
* gfortran.dg/goacc/parallel-tree.f95: Likewise.
* gfortran.dg/goacc/self.f95: Likewise.
libgomp/
* testsuite/libgomp.oacc-c-c++-common/if-1.c: Enhance.
* testsuite/libgomp.oacc-c-c++-common/self-1.c: Likewise.
* testsuite/libgomp.oacc-fortran/if-1.f90: Likewise.
* testsuite/libgomp.oacc-c-c++-common/if-self-1.c: New.
* testsuite/libgomp.oacc-fortran/self-1.f90: Likewise.
Diff:
---
gcc/testsuite/c-c++-common/goacc/if-clause-2.c | 2 +
gcc/testsuite/c-c++-common/goacc/self-clause-1.c | 6 +
gcc/testsuite/c-c++-common/goacc/self-clause-2.c | 20 +
gcc/testsuite/gfortran.dg/goacc/if.f95 | 10 +-
gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95 | 5 +-
gcc/testsuite/gfortran.dg/goacc/parallel-tree.f95 | 3 +-
gcc/testsuite/gfortran.dg/goacc/self.f95 | 8 +
libgomp/testsuite/libgomp.oacc-c-c++-common/if-1.c | 4 +
.../libgomp.oacc-c-c++-common/if-self-1.c | 36 +
.../testsuite/libgomp.oacc-c-c++-common/self-1.c | 5 +
libgomp/testsuite/libgomp.oacc-fortran/if-1.f90 | 4 +
libgomp/testsuite/libgomp.oacc-fortran/self-1.f90 | 996 +++++++++++++++++++++
12 files changed, 1094 insertions(+), 5 deletions(-)
diff --git a/gcc/testsuite/c-c++-common/goacc/if-clause-2.c b/gcc/testsuite/c-c++-common/goacc/if-clause-2.c
index a48072509e1..71475521758 100644
--- a/gcc/testsuite/c-c++-common/goacc/if-clause-2.c
+++ b/gcc/testsuite/c-c++-common/goacc/if-clause-2.c
@@ -1,3 +1,5 @@
+/* See also 'self-clause-2.c'. */
+
/* { dg-additional-options "-fdump-tree-gimple" } */
/* { dg-additional-options "--param=openacc-kernels=decompose" }
{ dg-additional-options "-fdump-tree-omp_oacc_kernels_decompose" } */
diff --git a/gcc/testsuite/c-c++-common/goacc/self-clause-1.c b/gcc/testsuite/c-c++-common/goacc/self-clause-1.c
index fe892bea210..28de3dc0584 100644
--- a/gcc/testsuite/c-c++-common/goacc/self-clause-1.c
+++ b/gcc/testsuite/c-c++-common/goacc/self-clause-1.c
@@ -5,6 +5,8 @@ f (int b)
{
struct { int i; } *p;
+#pragma acc parallel self(0) self(b) /* { dg-error "too many 'self' clauses" } */
+ ;
#pragma acc parallel self self(b) /* { dg-error "too many 'self' clauses" } */
;
#pragma acc parallel self(*p)
@@ -12,6 +14,8 @@ f (int b)
{ dg-error {could not convert '\* p' from 'f\(int\)::<unnamed struct>' to 'bool'} {} { target c++ } .-2 } */
;
+#pragma acc kernels self(0) self(b) /* { dg-error "too many 'self' clauses" } */
+ ;
#pragma acc kernels self self(b) /* { dg-error "too many 'self' clauses" } */
;
#pragma acc kernels self(*p)
@@ -19,6 +23,8 @@ f (int b)
{ dg-error {could not convert '\* p' from 'f\(int\)::<unnamed struct>' to 'bool'} {} { target c++ } .-2 } */
;
+#pragma acc serial self(0) self(b) /* { dg-error "too many 'self' clauses" } */
+ ;
#pragma acc serial self self(b) /* { dg-error "too many 'self' clauses" } */
;
#pragma acc serial self(*p)
diff --git a/gcc/testsuite/c-c++-common/goacc/self-clause-2.c b/gcc/testsuite/c-c++-common/goacc/self-clause-2.c
index d932ac9a4a6..769694baec9 100644
--- a/gcc/testsuite/c-c++-common/goacc/self-clause-2.c
+++ b/gcc/testsuite/c-c++-common/goacc/self-clause-2.c
@@ -1,3 +1,5 @@
+/* See also 'if-clause-2.c'. */
+
/* { dg-additional-options "-fdump-tree-gimple" } */
void
@@ -15,3 +17,21 @@ f (short c)
/* { dg-final { scan-tree-dump-times {(?n)#pragma omp target oacc_serial map\(tofrom:c \[len: [0-9]+\]\) self\(_[0-9]+\)$} 1 "gimple" } } */
++c;
}
+
+/* The same, but with implicit 'true' condition-argument. */
+
+void
+g (short d)
+{
+#pragma acc parallel self copy(d)
+ /* { dg-final { scan-tree-dump-times {(?n)#pragma omp target oacc_parallel map\(tofrom:d \[len: [0-9]+\]\) self\(1\)$} 1 "gimple" } } */
+ ++d;
+
+#pragma acc kernels self copy(d)
+ /* { dg-final { scan-tree-dump-times {(?n)#pragma omp target oacc_kernels map\(tofrom:d \[len: [0-9]+\]\) self\(1\)$} 1 "gimple" } } */
+ ++d;
+
+#pragma acc serial self copy(d)
+ /* { dg-final { scan-tree-dump-times {(?n)#pragma omp target oacc_serial map\(tofrom:d \[len: [0-9]+\]\) self\(1\)$} 1 "gimple" } } */
+ ++d;
+}
diff --git a/gcc/testsuite/gfortran.dg/goacc/if.f95 b/gcc/testsuite/gfortran.dg/goacc/if.f95
index 56f3711f320..753ef8251c2 100644
--- a/gcc/testsuite/gfortran.dg/goacc/if.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/if.f95
@@ -1,3 +1,5 @@
+! See also 'self.f95'.
+
! { dg-do compile }
program test
@@ -12,12 +14,14 @@ program test
!$acc end parallel
!$acc parallel if (1) ! { dg-error "scalar LOGICAL expression" }
!$acc end parallel
- !$acc kernels if (i) ! { dg-error "scalar LOGICAL expression" }
- !$acc end kernels
+
!$acc kernels if ! { dg-error "Expected '\\(' after 'if'" }
!$acc kernels if () ! { dg-error "Invalid character" }
+ !$acc kernels if (i) ! { dg-error "scalar LOGICAL expression" }
+ !$acc end kernels
!$acc kernels if (1) ! { dg-error "scalar LOGICAL expression" }
!$acc end kernels
+
!$acc data if ! { dg-error "Expected '\\(' after 'if'" }
!$acc data if () ! { dg-error "Invalid character" }
!$acc data if (i) ! { dg-error "scalar LOGICAL expression" }
@@ -36,12 +40,14 @@ program test
!$acc end parallel
!$acc parallel if (i.gt.1)
!$acc end parallel
+
!$acc kernels if (x)
!$acc end kernels
!$acc kernels if (.true.)
!$acc end kernels
!$acc kernels if (i.gt.1)
!$acc end kernels
+
!$acc data if (x)
!$acc end data
!$acc data if (.true.)
diff --git a/gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95
index ceb07fbb9e9..1ba04a84e12 100644
--- a/gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95
@@ -12,6 +12,7 @@ program test
logical :: l = .true.
!$acc kernels if(l) async num_gangs(i) num_workers(i) vector_length(i) &
+ !$acc self &
!$acc copy(i), copyin(j), copyout(k), create(m) &
!$acc no_create(n) &
!$acc present(o), pcopy(p), pcopyin(r), pcopyout(s), pcreate(t) &
@@ -27,7 +28,7 @@ end program test
! { dg-final { scan-tree-dump-times "num_gangs" 1 "original" } }
! { dg-final { scan-tree-dump-times "num_workers" 1 "original" } }
! { dg-final { scan-tree-dump-times "vector_length" 1 "original" } }
-
+! { dg-final { scan-tree-dump-times "self\\(1\\)" 1 "original" } }
! { dg-final { scan-tree-dump-times "map\\(tofrom:i\\)" 1 "original" } }
! { dg-final { scan-tree-dump-times "map\\(to:j\\)" 1 "original" } }
! { dg-final { scan-tree-dump-times "map\\(from:k\\)" 1 "original" } }
@@ -42,4 +43,4 @@ end program test
! { dg-final { scan-tree-dump-times "map\\(force_deviceptr:u\\)" 1 "original" } }
! { dg-final { scan-tree-dump-times {(?n)#pragma omp target oacc_data_kernels if\((?:D\.|_)[0-9]+\)$} 1 "omp_oacc_kernels_decompose" } }
-! { dg-final { scan-tree-dump-times {(?n)#pragma omp target oacc_parallel_kernels_gang_single num_gangs\(1\) if\((?:D\.|_)[0-9]+\) async\(-1\)$} 1 "omp_oacc_kernels_decompose" } }
+! { dg-final { scan-tree-dump-times {(?n)#pragma omp target oacc_parallel_kernels_gang_single num_gangs\(1\) if\((?:D\.|_)[0-9]+\) self\(1\) async\(-1\)$} 1 "omp_oacc_kernels_decompose" } }
diff --git a/gcc/testsuite/gfortran.dg/goacc/parallel-tree.f95 b/gcc/testsuite/gfortran.dg/goacc/parallel-tree.f95
index 6110d93b91e..0d4ec1133af 100644
--- a/gcc/testsuite/gfortran.dg/goacc/parallel-tree.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/parallel-tree.f95
@@ -14,6 +14,7 @@ program test
logical :: l = .true.
!$acc parallel if(l) async num_gangs(i) num_workers(i) vector_length(i) &
+ !$acc self &
!$acc reduction(max:q), copy(i), copyin(j), copyout(k), create(m) &
!$acc no_create(n) &
!$acc present(o), pcopy(p), pcopyin(r), pcopyout(s), pcreate(t) &
@@ -33,7 +34,7 @@ end program test
! { dg-final { scan-tree-dump-times "num_gangs" 1 "original" } }
! { dg-final { scan-tree-dump-times "num_workers" 1 "original" } }
! { dg-final { scan-tree-dump-times "vector_length" 1 "original" } }
-
+! { dg-final { scan-tree-dump-times "self\\(1\\)" 1 "original" } }
! { dg-final { scan-tree-dump-times "reduction\\(max:q\\)" 1 "original" } }
! { dg-final { scan-tree-dump-times "map\\(tofrom:i\\)" 1 "original" } }
! { dg-final { scan-tree-dump-times "map\\(to:j\\)" 1 "original" } }
diff --git a/gcc/testsuite/gfortran.dg/goacc/self.f95 b/gcc/testsuite/gfortran.dg/goacc/self.f95
index 4817f16be56..aa0f6fe88c5 100644
--- a/gcc/testsuite/gfortran.dg/goacc/self.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/self.f95
@@ -1,3 +1,5 @@
+! See also 'if.f95'.
+
! { dg-do compile }
program test
@@ -29,6 +31,8 @@ program test
!$acc kernels self (.false.) self (.false.) { dg-error "Duplicated 'self' clause" }
!$acc serial self (.false.) self (.false.) { dg-error "Duplicated 'self' clause" }
+ !$acc parallel self
+ !$acc end parallel
!$acc parallel self (x)
!$acc end parallel
!$acc parallel self (.true.)
@@ -36,6 +40,8 @@ program test
!$acc parallel self (i.gt.1)
!$acc end parallel
+ !$acc kernels self
+ !$acc end kernels
!$acc kernels self (x)
!$acc end kernels
!$acc kernels self (.true.)
@@ -43,6 +49,8 @@ program test
!$acc kernels self (i.gt.1)
!$acc end kernels
+ !$acc serial self
+ !$acc end serial
!$acc serial self (x)
!$acc end serial
!$acc serial self (.true.)
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/if-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/if-1.c
index 5398905f411..f04c02fdb89 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/if-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/if-1.c
@@ -1,3 +1,7 @@
+/* OpenACC 'if' clause. */
+
+/* See also 'self-1.c'. */
+
#include <openacc.h>
#include <stdlib.h>
#include <stdbool.h>
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/if-self-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/if-self-1.c
new file mode 100644
index 00000000000..f77edda83e3
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/if-self-1.c
@@ -0,0 +1,36 @@
+/* Test 'if' and 'self' clause appearing together. */
+
+#include <openacc.h>
+
+static int test(float i, long double s)
+{
+ int ret;
+#pragma acc serial copyout(ret) if(i) self(s)
+ /* { dg-bogus {using 'vector_length \(32\)', ignoring 1} {} { xfail openacc_nvidia_accel_selected } .-1 } */
+ {
+ ret = acc_on_device(acc_device_host);
+ }
+ return ret;
+}
+
+int main()
+{
+ if (!test(0, 0))
+ __builtin_abort();
+
+ if (!test(0, 1))
+ __builtin_abort();
+
+#if ACC_MEM_SHARED
+ if (!test(1, 0))
+ __builtin_abort();
+#else
+ if (test(1, 0))
+ __builtin_abort();
+#endif
+
+ if (!test(1, 1))
+ __builtin_abort();
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/self-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/self-1.c
index 752e16e8545..d17e27c8af3 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/self-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/self-1.c
@@ -1,3 +1,8 @@
+/* OpenACC 'self' clause. */
+
+/* This is 'if-1.c' with 'self(!cond)' instead of 'if(cond)' on compute
+ constructs. */
+
#include <openacc.h>
#include <stdlib.h>
#include <stdbool.h>
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/if-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/if-1.f90
index e0cfd912d0f..5ba6509749e 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/if-1.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/if-1.f90
@@ -1,3 +1,7 @@
+! OpenACC 'if' clause.
+
+! See also 'self-1.f90'.
+
! { dg-do run }
! { dg-additional-options "-cpp" }
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/self-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/self-1.f90
new file mode 100644
index 00000000000..b9ec9de08d9
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/self-1.f90
@@ -0,0 +1,996 @@
+! OpenACC 'self' clause.
+
+! This is 'if-1.f90' with 'self(!cond)' instead of 'if(cond)' on compute
+! constructs.
+! ..., which the exception of certain 'kernels' constructs.
+
+! { dg-do run }
+! { dg-additional-options "-cpp" }
+
+! { dg-additional-options "--param=openacc-kernels=decompose" }
+
+! { dg-additional-options "-fopt-info-note-omp" }
+! { dg-additional-options "-foffload=-fopt-info-note-omp" }
+
+! { dg-additional-options "--param=openacc-privatization=noisy" }
+! { dg-additional-options "-foffload=--param=openacc-privatization=noisy" }
+! Prune a few: uninteresting, and potentially varying depending on GCC configuration (data types):
+! { dg-prune-output {note: variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} }
+
+! It's only with Tcl 8.5 (released in 2007) that "the variable 'varName'
+! passed to 'incr' may be unset, and in that case, it will be set to [...]",
+! so to maintain compatibility with earlier Tcl releases, we manually
+! initialize counter variables:
+! { dg-line l_dummy[variable c_compute 0] }
+! { dg-message "dummy" "" { target iN-VAl-Id } l_dummy } to avoid
+! "WARNING: dg-line var l_dummy defined, but not used". */
+
+program main
+ use openacc
+ implicit none
+
+ integer, parameter :: N = 8
+ integer, parameter :: one = 1
+ integer, parameter :: zero = 0
+ integer i, nn
+ real, allocatable :: a(:), b(:)
+ real exp, exp2
+
+ i = 0
+
+ allocate (a(N))
+ allocate (b(N))
+
+ a(:) = 4.0
+
+ !$acc parallel copyin (a(1:N)) copyout (b(1:N)) self (1 /= 1) ! { dg-line l_compute[incr c_compute] }
+ do i = 1, N
+ ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute }
+ !TODO Unhandled 'CONST_DECL' instances for constant argument in 'acc_on_device' call.
+ if (acc_on_device (acc_device_host) .eqv. .TRUE.) then
+ b(i) = a(i) + 1
+ else
+ b(i) = a(i)
+ end if
+ end do
+ !$acc end parallel
+
+#if ACC_MEM_SHARED
+ exp = 5.0
+#else
+ exp = 4.0
+#endif
+
+ do i = 1, N
+ if (b(i) .ne. exp) STOP 1
+ end do
+
+ a(:) = 16.0
+
+ !$acc parallel self (0 /= 1) ! { dg-line l_compute[incr c_compute] }
+ do i = 1, N
+ ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute }
+ if (acc_on_device (acc_device_host) .eqv. .TRUE.) then
+ b(i) = a(i) + 1
+ else
+ b(i) = a(i)
+ end if
+ end do
+ !$acc end parallel
+
+ do i = 1, N
+ if (b(i) .ne. 17.0) STOP 2
+ end do
+
+ a(:) = 8.0
+
+ !$acc parallel copyin (a(1:N)) copyout (b(1:N)) self (one /= 1) ! { dg-line l_compute[incr c_compute] }
+ do i = 1, N
+ ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute }
+ if (acc_on_device (acc_device_host) .eqv. .TRUE.) then
+ b(i) = a(i) + 1
+ else
+ b(i) = a(i)
+ end if
+ end do
+ !$acc end parallel
+
+#if ACC_MEM_SHARED
+ exp = 9.0
+#else
+ exp = 8.0
+#endif
+
+ do i = 1, N
+ if (b(i) .ne. exp) STOP 3
+ end do
+
+ a(:) = 22.0
+
+ !$acc parallel self (zero /= 1) ! { dg-line l_compute[incr c_compute] }
+ do i = 1, N
+ ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute }
+ if (acc_on_device (acc_device_host) .eqv. .TRUE.) then
+ b(i) = a(i) + 1
+ else
+ b(i) = a(i)
+ end if
+ end do
+ !$acc end parallel
+
+ do i = 1, N
+ if (b(i) .ne. 23.0) STOP 4
+ end do
+
+ a(:) = 16.0
+
+ !$acc parallel copyin (a(1:N)) copyout (b(1:N)) self (.FALSE.) ! { dg-line l_compute[incr c_compute] }
+ do i = 1, N
+ ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute }
+ if (acc_on_device (acc_device_host) .eqv. .TRUE.) then
+ b(i) = a(i) + 1
+ else
+ b(i) = a(i)
+ end if
+ end do
+ !$acc end parallel
+
+#if ACC_MEM_SHARED
+ exp = 17.0;
+#else
+ exp = 16.0;
+#endif
+
+ do i = 1, N
+ if (b(i) .ne. exp) STOP 5
+ end do
+
+ a(:) = 76.0
+
+ !$acc parallel self (.TRUE.) ! { dg-line l_compute[incr c_compute] }
+ do i = 1, N
+ ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute }
+ if (acc_on_device (acc_device_host) .eqv. .TRUE.) then
+ b(i) = a(i) + 1
+ else
+ b(i) = a(i)
+ end if
+ end do
+ !$acc end parallel
+
+ do i = 1, N
+ if (b(i) .ne. 77.0) STOP 6
+ end do
+
+ a(:) = 22.0
+
+ nn = 1
+
+ !$acc parallel copyin (a(1:N)) copyout (b(1:N)) self (nn /= 1) ! { dg-line l_compute[incr c_compute] }
+ do i = 1, N
+ ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute }
+ if (acc_on_device (acc_device_host) .eqv. .TRUE.) then
+ b(i) = a(i) + 1
+ else
+ b(i) = a(i)
+ end if
+ end do
+ !$acc end parallel
+
+#if ACC_MEM_SHARED
+ exp = 23.0;
+#else
+ exp = 22.0;
+#endif
+
+ do i = 1, N
+ if (b(i) .ne. exp) STOP 7
+ end do
+
+ a(:) = 18.0
+
+ nn = 0
+
+ !$acc parallel self (nn /= 1) ! { dg-line l_compute[incr c_compute] }
+ do i = 1, N
+ ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute }
+ if (acc_on_device (acc_device_host) .eqv. .TRUE.) then
+ b(i) = a(i) + 1
+ else
+ b(i) = a(i)
+ end if
+ end do
+ !$acc end parallel
+
+ do i = 1, N
+ if (b(i) .ne. 19.0) STOP 8
+ end do
+
+ a(:) = 49.0
+
+ nn = 1
+
+ !$acc parallel copyin (a(1:N)) copyout (b(1:N)) self (.NOT. ((nn + nn) > 0)) ! { dg-line l_compute[incr c_compute] }
+ do i = 1, N
+ ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute }
+ if (acc_on_device (acc_device_host) .eqv. .TRUE.) then
+ b(i) = a(i) + 1
+ else
+ b(i) = a(i)
+ end if
+ end do
+ !$acc end parallel
+
+#if ACC_MEM_SHARED
+ exp = 50.0
+#else
+ exp = 49.0
+#endif
+
+ do i = 1, N
+ if (b(i) .ne. exp) STOP 9
+ end do
+
+ a(:) = 38.0
+
+ nn = 0;
+
+ !$acc parallel copyin (a(1:N)) copyout (b(1:N)) self (.NOT. ((nn + nn) > 0)) ! { dg-line l_compute[incr c_compute] }
+ do i = 1, N
+ ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute }
+ if (acc_on_device (acc_device_host) .eqv. .TRUE.) then
+ b(i) = a(i) + 1
+ else
+ b(i) = a(i)
+ end if
+ end do
+ !$acc end parallel
+
+ do i = 1, N
+ if (b(i) .ne. 39.0) STOP 10
+ end do
+
+ a(:) = 91.0
+
+ !$acc parallel copyin (a(1:N)) copyout (b(1:N)) self (.NOT. (-2 > 0)) ! { dg-line l_compute[incr c_compute] }
+ do i = 1, N
+ ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute }
+ if (acc_on_device (acc_device_host) .eqv. .TRUE.) then
+ b(i) = a(i) + 1
+ else
+ b(i) = a(i)
+ end if
+ end do
+ !$acc end parallel
+
+ do i = 1, N
+ if (b(i) .ne. 92.0) STOP 11
+ end do
+
+ a(:) = 43.0
+
+ !$acc parallel copyin (a(1:N)) copyout (b(1:N)) self (one /= 1) ! { dg-line l_compute[incr c_compute] }
+ do i = 1, N
+ ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute }
+ if (acc_on_device (acc_device_host) .eqv. .TRUE.) then
+ b(i) = a(i) + 1
+ else
+ b(i) = a(i)
+ end if
+ end do
+ !$acc end parallel
+
+#if ACC_MEM_SHARED
+ exp = 44.0
+#else
+ exp = 43.0
+#endif
+
+ do i = 1, N
+ if (b(i) .ne. exp) STOP 12
+ end do
+
+ a(:) = 87.0
+
+ !$acc parallel self (one /= 0) ! { dg-line l_compute[incr c_compute] }
+ do i = 1, N
+ ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute }
+ if (acc_on_device (acc_device_host) .eqv. .TRUE.) then
+ b(i) = a(i) + 1
+ else
+ b(i) = a(i)
+ end if
+ end do
+ !$acc end parallel
+
+ do i = 1, N
+ if (b(i) .ne. 88.0) STOP 13
+ end do
+
+ a(:) = 3.0
+ b(:) = 9.0
+
+#if ACC_MEM_SHARED
+ exp = 0.0
+ exp2 = 0.0
+#else
+ call acc_copyin (a, sizeof (a))
+ call acc_copyin (b, sizeof (b))
+ exp = 3.0;
+ exp2 = 9.0;
+#endif
+
+ !$acc update device (a(1:N), b(1:N)) if (1 == 1)
+
+ a(:) = 0.0
+ b(:) = 0.0
+
+ !$acc update host (a(1:N), b(1:N)) if (1 == 1)
+
+ do i = 1, N
+ if (a(i) .ne. exp) STOP 14
+ if (b(i) .ne. exp2) STOP 15
+ end do
+
+ a(:) = 6.0
+ b(:) = 12.0
+
+ !$acc update device (a(1:N), b(1:N)) if (0 == 1)
+
+ a(:) = 0.0
+ b(:) = 0.0
+
+ !$acc update host (a(1:N), b(1:N)) if (1 == 1)
+
+ do i = 1, N
+ if (a(i) .ne. exp) STOP 16
+ if (b(i) .ne. exp2) STOP 17
+ end do
+
+ a(:) = 26.0
+ b(:) = 21.0
+
+ !$acc update device (a(1:N), b(1:N)) if (1 == 1)
+
+ a(:) = 0.0
+ b(:) = 0.0
+
+ !$acc update host (a(1:N), b(1:N)) if (0 == 1)
+
+ do i = 1, N
+ if (a(i) .ne. 0.0) STOP 18
+ if (b(i) .ne. 0.0) STOP 19
+ end do
+
+#if !ACC_MEM_SHARED
+ call acc_copyout (a, sizeof (a))
+ call acc_copyout (b, sizeof (b))
+#endif
+
+ a(:) = 4.0
+ b(:) = 0.0
+
+ !$acc data copyin (a(1:N)) copyout (b(1:N)) if (1 == 1)
+ ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-1 }
+
+ !$acc parallel present (a(1:N))
+ do i = 1, N
+ b(i) = a(i)
+ end do
+ !$acc end parallel
+ !$acc end data
+
+ do i = 1, N
+ if (b(i) .ne. 4.0) STOP 20
+ end do
+
+ a(:) = 8.0
+ b(:) = 1.0
+
+ !$acc data copyin (a(1:N)) copyout (b(1:N)) if (0 == 1)
+ ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target { ! openacc_host_selected } } .-1 }
+
+#if !ACC_MEM_SHARED
+ if (acc_is_present (a) .eqv. .TRUE.) STOP 21
+ if (acc_is_present (b) .eqv. .TRUE.) STOP 22
+#endif
+
+ !$acc end data
+
+ a(:) = 18.0
+ b(:) = 21.0
+
+ !$acc data copyin (a(1:N)) if (1 == 1)
+ ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-1 }
+ ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target { ! openacc_host_selected } } .-2 }
+
+#if !ACC_MEM_SHARED
+ if (acc_is_present (a) .eqv. .FALSE.) STOP 23
+#endif
+
+ !$acc data copyout (b(1:N)) if (0 == 1)
+ ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-1 }
+ ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target { ! openacc_host_selected } } .-2 }
+#if !ACC_MEM_SHARED
+ if (acc_is_present (b) .eqv. .TRUE.) STOP 24
+#endif
+ !$acc data copyout (b(1:N)) if (1 == 1)
+ ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-1 }
+
+ !$acc parallel present (a(1:N)) present (b(1:N))
+ do i = 1, N
+ b(i) = a(i)
+ end do
+ !$acc end parallel
+
+ !$acc end data
+
+#if !ACC_MEM_SHARED
+ if (acc_is_present (b) .eqv. .TRUE.) STOP 25
+#endif
+ !$acc end data
+ !$acc end data
+
+ do i = 1, N
+ if (b(1) .ne. 18.0) STOP 26
+ end do
+
+ !$acc enter data copyin (b(1:N)) if (0 == 1)
+
+#if !ACC_MEM_SHARED
+ if (acc_is_present (b) .eqv. .TRUE.) STOP 27
+#endif
+
+ !$acc exit data delete (b(1:N)) if (0 == 1)
+
+ !$acc enter data copyin (b(1:N)) if (1 == 1)
+
+#if !ACC_MEM_SHARED
+ if (acc_is_present (b) .eqv. .FALSE.) STOP 28
+#endif
+
+ !$acc exit data delete (b(1:N)) if (1 == 1)
+
+#if !ACC_MEM_SHARED
+ if (acc_is_present (b) .eqv. .TRUE.) STOP 29
+#endif
+
+ !$acc enter data copyin (b(1:N)) if (zero == 1)
+
+#if !ACC_MEM_SHARED
+ if (acc_is_present (b) .eqv. .TRUE.) STOP 30
+#endif
+
+ !$acc exit data delete (b(1:N)) if (zero == 1)
+
+ !$acc enter data copyin (b(1:N)) if (one == 1)
+
+#if !ACC_MEM_SHARED
+ if (acc_is_present (b) .eqv. .FALSE.) STOP 31
+#endif
+
+ !$acc exit data delete (b(1:N)) if (one == 1)
+
+#if !ACC_MEM_SHARED
+ if (acc_is_present (b) .eqv. .TRUE.) STOP 32
+#endif
+
+ !$acc enter data copyin (b(1:N)) if (one == 0)
+
+#if !ACC_MEM_SHARED
+ if (acc_is_present (b) .eqv. .TRUE.) STOP 33
+#endif
+
+ !$acc exit data delete (b(1:N)) if (one == 0)
+
+ !$acc enter data copyin (b(1:N)) if (one == 1)
+
+#if !ACC_MEM_SHARED
+ if (acc_is_present (b) .eqv. .FALSE.) STOP 34
+#endif
+
+ !$acc exit data delete (b(1:N)) if (one == 1)
+
+#if !ACC_MEM_SHARED
+ if (acc_is_present (b) .eqv. .TRUE.) STOP 35
+#endif
+
+ a(:) = 4.0
+
+ !$acc kernels copyin (a(1:N)) copyout (b(1:N)) self (1 /= 1) ! { dg-line l_compute[incr c_compute] }
+ ! { dg-note {OpenACC 'kernels' decomposition: variable 'i' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute }
+ ! { dg-note {variable 'i' made addressable} {} { target *-*-* } l_compute$c_compute } */
+ ! { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} "" { target *-*-* } .+1 }
+ do i = 1, N
+ ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute }
+ if (acc_on_device (acc_device_host) .eqv. .TRUE.) then
+ b(i) = a(i) + 1
+ else
+ b(i) = a(i)
+ end if
+ end do
+ !$acc end kernels
+
+#if ACC_MEM_SHARED
+ exp = 5.0
+#else
+ exp = 4.0
+#endif
+
+ do i = 1, N
+ if (b(i) .ne. exp) STOP 36
+ end do
+
+ a(:) = 16.0
+
+ !$acc kernels if (0 == 1) ! { dg-line l_compute[incr c_compute] }
+ ! { dg-note {OpenACC 'kernels' decomposition: variable 'i' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute }
+ ! { dg-note {variable 'i' already made addressable} {} { target *-*-* } l_compute$c_compute } */
+ ! { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} "" { target *-*-* } .+1 }
+ do i = 1, N
+ ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute }
+ if (acc_on_device (acc_device_host) .eqv. .TRUE.) then
+ b(i) = a(i) + 1
+ else
+ b(i) = a(i)
+ end if
+ end do
+ !$acc end kernels
+
+ do i = 1, N
+ if (b(i) .ne. 17.0) STOP 37
+ end do
+
+ a(:) = 8.0
+
+ !$acc kernels copyin (a(1:N)) copyout (b(1:N)) self (one /= 1) ! { dg-line l_compute[incr c_compute] }
+ ! { dg-note {OpenACC 'kernels' decomposition: variable 'i' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute }
+ ! { dg-note {variable 'i' already made addressable} {} { target *-*-* } l_compute$c_compute } */
+ ! { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} "" { target *-*-* } .+1 }
+ do i = 1, N
+ ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute }
+ if (acc_on_device (acc_device_host) .eqv. .TRUE.) then
+ b(i) = a(i) + 1
+ else
+ b(i) = a(i)
+ end if
+ end do
+ !$acc end kernels
+
+#if ACC_MEM_SHARED
+ exp = 9.0
+#else
+ exp = 8.0
+#endif
+
+ do i = 1, N
+ if (b(i) .ne. exp) STOP 38
+ end do
+
+ a(:) = 22.0
+
+ !$acc kernels if (zero == 1) ! { dg-line l_compute[incr c_compute] }
+ ! { dg-note {OpenACC 'kernels' decomposition: variable 'i' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute }
+ ! { dg-note {variable 'i' already made addressable} {} { target *-*-* } l_compute$c_compute } */
+ ! { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} "" { target *-*-* } .+1 }
+ do i = 1, N
+ ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute }
+ if (acc_on_device (acc_device_host) .eqv. .TRUE.) then
+ b(i) = a(i) + 1
+ else
+ b(i) = a(i)
+ end if
+ end do
+ !$acc end kernels
+
+ do i = 1, N
+ if (b(i) .ne. 23.0) STOP 39
+ end do
+
+ a(:) = 16.0
+
+ !$acc kernels copyin (a(1:N)) copyout (b(1:N)) self (.FALSE.) ! { dg-line l_compute[incr c_compute] }
+ ! { dg-note {OpenACC 'kernels' decomposition: variable 'i' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute }
+ ! { dg-note {variable 'i' already made addressable} {} { target *-*-* } l_compute$c_compute } */
+ ! { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} "" { target *-*-* } .+1 }
+ do i = 1, N
+ ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute }
+ if (acc_on_device (acc_device_host) .eqv. .TRUE.) then
+ b(i) = a(i) + 1
+ else
+ b(i) = a(i)
+ end if
+ end do
+ !$acc end kernels
+
+#if ACC_MEM_SHARED
+ exp = 17.0;
+#else
+ exp = 16.0;
+#endif
+
+ do i = 1, N
+ if (b(i) .ne. exp) STOP 40
+ end do
+
+ a(:) = 76.0
+
+ !$acc kernels if (.FALSE.) ! { dg-line l_compute[incr c_compute] }
+ ! { dg-note {OpenACC 'kernels' decomposition: variable 'i' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute }
+ ! { dg-note {variable 'i' already made addressable} {} { target *-*-* } l_compute$c_compute } */
+ ! { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} "" { target *-*-* } .+1 }
+ do i = 1, N
+ ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute }
+ if (acc_on_device (acc_device_host) .eqv. .TRUE.) then
+ b(i) = a(i) + 1
+ else
+ b(i) = a(i)
+ end if
+ end do
+ !$acc end kernels
+
+ do i = 1, N
+ if (b(i) .ne. 77.0) STOP 41
+ end do
+
+ a(:) = 22.0
+
+ nn = 1
+
+ !$acc kernels copyin (a(1:N)) copyout (b(1:N)) self (nn /= 1) ! { dg-line l_compute[incr c_compute] }
+ ! { dg-note {OpenACC 'kernels' decomposition: variable 'i' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute }
+ ! { dg-note {variable 'i' already made addressable} {} { target *-*-* } l_compute$c_compute } */
+ ! { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} "" { target *-*-* } .+1 }
+ do i = 1, N
+ ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute }
+ if (acc_on_device (acc_device_host) .eqv. .TRUE.) then
+ b(i) = a(i) + 1
+ else
+ b(i) = a(i)
+ end if
+ end do
+ !$acc end kernels
+
+#if ACC_MEM_SHARED
+ exp = 23.0;
+#else
+ exp = 22.0;
+#endif
+
+ do i = 1, N
+ if (b(i) .ne. exp) STOP 42
+ end do
+
+ a(:) = 18.0
+
+ nn = 0
+
+ !$acc kernels if (nn == 1) ! { dg-line l_compute[incr c_compute] }
+ ! { dg-note {OpenACC 'kernels' decomposition: variable 'i' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute }
+ ! { dg-note {variable 'i' already made addressable} {} { target *-*-* } l_compute$c_compute } */
+ ! { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} "" { target *-*-* } .+1 }
+ do i = 1, N
+ ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute }
+ if (acc_on_device (acc_device_host) .eqv. .TRUE.) then
+ b(i) = a(i) + 1
+ else
+ b(i) = a(i)
+ end if
+ end do
+ !$acc end kernels
+
+ do i = 1, N
+ if (b(i) .ne. 19.0) STOP 43
+ end do
+
+ a(:) = 49.0
+
+ nn = 1
+
+ !$acc kernels copyin (a(1:N)) copyout (b(1:N)) self (.NOT. ((nn + nn) > 0)) ! { dg-line l_compute[incr c_compute] }
+ ! { dg-note {OpenACC 'kernels' decomposition: variable 'i' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute }
+ ! { dg-note {variable 'i' already made addressable} {} { target *-*-* } l_compute$c_compute } */
+ ! { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} "" { target *-*-* } .+1 }
+ do i = 1, N
+ ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute }
+ if (acc_on_device (acc_device_host) .eqv. .TRUE.) then
+ b(i) = a(i) + 1
+ else
+ b(i) = a(i)
+ end if
+ end do
+ !$acc end kernels
+
+#if ACC_MEM_SHARED
+ exp = 50.0
+#else
+ exp = 49.0
+#endif
+
+ do i = 1, N
+ if (b(i) .ne. exp) STOP 44
+ end do
+
+ a(:) = 38.0
+
+ nn = 0;
+
+ !$acc kernels copyin (a(1:N)) copyout (b(1:N)) if ((nn + nn) > 0) ! { dg-line l_compute[incr c_compute] }
+ ! { dg-note {OpenACC 'kernels' decomposition: variable 'i' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute }
+ ! { dg-note {variable 'i' already made addressable} {} { target *-*-* } l_compute$c_compute } */
+ ! { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} "" { target *-*-* } .+1 }
+ do i = 1, N
+ ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute }
+ if (acc_on_device (acc_device_host) .eqv. .TRUE.) then
+ b(i) = a(i) + 1
+ else
+ b(i) = a(i)
+ end if
+ end do
+ !$acc end kernels
+
+ do i = 1, N
+ if (b(i) .ne. 39.0) STOP 45
+ end do
+
+ a(:) = 91.0
+
+ !$acc kernels copyin (a(1:N)) copyout (b(1:N)) if (-2 > 0) ! { dg-line l_compute[incr c_compute] }
+ ! { dg-note {OpenACC 'kernels' decomposition: variable 'i' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute }
+ ! { dg-note {variable 'i' already made addressable} {} { target *-*-* } l_compute$c_compute } */
+ ! { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} "" { target *-*-* } .+1 }
+ do i = 1, N
+ ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute }
+ if (acc_on_device (acc_device_host) .eqv. .TRUE.) then
+ b(i) = a(i) + 1
+ else
+ b(i) = a(i)
+ end if
+ end do
+ !$acc end kernels
+
+ do i = 1, N
+ if (b(i) .ne. 92.0) STOP 46
+ end do
+
+ a(:) = 43.0
+
+ !$acc kernels copyin (a(1:N)) copyout (b(1:N)) self (one /= 1) ! { dg-line l_compute[incr c_compute] }
+ ! { dg-note {OpenACC 'kernels' decomposition: variable 'i' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute }
+ ! { dg-note {variable 'i' already made addressable} {} { target *-*-* } l_compute$c_compute } */
+ ! { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} "" { target *-*-* } .+1 }
+ do i = 1, N
+ ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute }
+ if (acc_on_device (acc_device_host) .eqv. .TRUE.) then
+ b(i) = a(i) + 1
+ else
+ b(i) = a(i)
+ end if
+ end do
+ !$acc end kernels
+
+#if ACC_MEM_SHARED
+ exp = 44.0
+#else
+ exp = 43.0
+#endif
+
+ do i = 1, N
+ if (b(i) .ne. exp) STOP 47
+ end do
+
+ a(:) = 87.0
+
+ !$acc kernels if (one == 0) ! { dg-line l_compute[incr c_compute] }
+ ! { dg-note {OpenACC 'kernels' decomposition: variable 'i' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute }
+ ! { dg-note {variable 'i' already made addressable} {} { target *-*-* } l_compute$c_compute } */
+ ! { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} "" { target *-*-* } .+1 }
+ do i = 1, N
+ ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute }
+ if (acc_on_device (acc_device_host) .eqv. .TRUE.) then
+ b(i) = a(i) + 1
+ else
+ b(i) = a(i)
+ end if
+ end do
+ !$acc end kernels
+
+ do i = 1, N
+ if (b(i) .ne. 88.0) STOP 48
+ end do
+
+ a(:) = 3.0
+ b(:) = 9.0
+
+#if ACC_MEM_SHARED
+ exp = 0.0
+ exp2 = 0.0
+#else
+ call acc_copyin (a, sizeof (a))
+ call acc_copyin (b, sizeof (b))
+ exp = 3.0;
+ exp2 = 9.0;
+#endif
+
+ !$acc update device (a(1:N), b(1:N)) if (1 == 1)
+
+ a(:) = 0.0
+ b(:) = 0.0
+
+ !$acc update host (a(1:N), b(1:N)) if (1 == 1)
+
+ do i = 1, N
+ if (a(i) .ne. exp) STOP 49
+ if (b(i) .ne. exp2) STOP 50
+ end do
+
+ a(:) = 6.0
+ b(:) = 12.0
+
+ !$acc update device (a(1:N), b(1:N)) if (0 == 1)
+
+ a(:) = 0.0
+ b(:) = 0.0
+
+ !$acc update host (a(1:N), b(1:N)) if (1 == 1)
+
+ do i = 1, N
+ if (a(i) .ne. exp) STOP 51
+ if (b(i) .ne. exp2) STOP 52
+ end do
+
+ a(:) = 26.0
+ b(:) = 21.0
+
+ !$acc update device (a(1:N), b(1:N)) if (1 == 1)
+
+ a(:) = 0.0
+ b(:) = 0.0
+
+ !$acc update host (a(1:N), b(1:N)) if (0 == 1)
+
+ do i = 1, N
+ if (a(i) .ne. 0.0) STOP 53
+ if (b(i) .ne. 0.0) STOP 54
+ end do
+
+#if !ACC_MEM_SHARED
+ call acc_copyout (a, sizeof (a))
+ call acc_copyout (b, sizeof (b))
+#endif
+
+ a(:) = 4.0
+ b(:) = 0.0
+
+ !$acc data copyin (a(1:N)) copyout (b(1:N)) if (1 == 1)
+ ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-1 }
+
+ !$acc kernels present (a(1:N)) ! { dg-line l_compute[incr c_compute] }
+ ! { dg-note {OpenACC 'kernels' decomposition: variable 'i' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute }
+ ! { dg-note {variable 'i' already made addressable} {} { target *-*-* } l_compute$c_compute } */
+ ! { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} "" { target *-*-* } .+1 }
+ do i = 1, N
+ b(i) = a(i)
+ end do
+ !$acc end kernels
+ !$acc end data
+
+ do i = 1, N
+ if (b(i) .ne. 4.0) STOP 55
+ end do
+
+ a(:) = 8.0
+ b(:) = 1.0
+
+ !$acc data copyin (a(1:N)) copyout (b(1:N)) if (0 == 1)
+ ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target { ! openacc_host_selected } } .-1 }
+
+#if !ACC_MEM_SHARED
+ if (acc_is_present (a) .eqv. .TRUE.) STOP 56
+ if (acc_is_present (b) .eqv. .TRUE.) STOP 57
+#endif
+
+ !$acc end data
+
+ a(:) = 18.0
+ b(:) = 21.0
+
+ !$acc data copyin (a(1:N)) if (1 == 1)
+ ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-1 }
+ ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target { ! openacc_host_selected } } .-2 }
+
+#if !ACC_MEM_SHARED
+ if (acc_is_present (a) .eqv. .FALSE.) STOP 58
+#endif
+
+ !$acc data copyout (b(1:N)) if (0 == 1)
+ ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-1 }
+ ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target { ! openacc_host_selected } } .-2 }
+#if !ACC_MEM_SHARED
+ if (acc_is_present (b) .eqv. .TRUE.) STOP 59
+#endif
+ !$acc data copyout (b(1:N)) if (1 == 1)
+ ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-1 }
+
+ !$acc kernels present (a(1:N)) present (b(1:N)) ! { dg-line l_compute[incr c_compute] }
+ ! { dg-note {OpenACC 'kernels' decomposition: variable 'i' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute }
+ ! { dg-note {variable 'i' already made addressable} {} { target *-*-* } l_compute$c_compute } */
+ ! { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} "" { target *-*-* } .+1 }
+ do i = 1, N
+ b(i) = a(i)
+ end do
+ !$acc end kernels
+
+ !$acc end data
+
+#if !ACC_MEM_SHARED
+ if (acc_is_present (b) .eqv. .TRUE.) STOP 60
+#endif
+ !$acc end data
+ !$acc end data
+
+ do i = 1, N
+ if (b(1) .ne. 18.0) STOP 61
+ end do
+
+ !$acc enter data copyin (b(1:N)) if (0 == 1)
+
+#if !ACC_MEM_SHARED
+ if (acc_is_present (b) .eqv. .TRUE.) STOP 62
+#endif
+
+ !$acc exit data delete (b(1:N)) if (0 == 1)
+
+ !$acc enter data copyin (b(1:N)) if (1 == 1)
+
+#if !ACC_MEM_SHARED
+ if (acc_is_present (b) .eqv. .FALSE.) STOP 63
+#endif
+
+ !$acc exit data delete (b(1:N)) if (1 == 1)
+
+#if !ACC_MEM_SHARED
+ if (acc_is_present (b) .eqv. .TRUE.) STOP 64
+#endif
+
+ !$acc enter data copyin (b(1:N)) if (zero == 1)
+
+#if !ACC_MEM_SHARED
+ if (acc_is_present (b) .eqv. .TRUE.) STOP 65
+#endif
+
+ !$acc exit data delete (b(1:N)) if (zero == 1)
+
+ !$acc enter data copyin (b(1:N)) if (one == 1)
+
+#if !ACC_MEM_SHARED
+ if (acc_is_present (b) .eqv. .FALSE.) STOP 66
+#endif
+
+ !$acc exit data delete (b(1:N)) if (one == 1)
+
+#if !ACC_MEM_SHARED
+ if (acc_is_present (b) .eqv. .TRUE.) STOP 67
+#endif
+
+ !$acc enter data copyin (b(1:N)) if (one == 0)
+
+#if !ACC_MEM_SHARED
+ if (acc_is_present (b) .eqv. .TRUE.) STOP 68
+#endif
+
+ !$acc exit data delete (b(1:N)) if (one == 0)
+
+ !$acc enter data copyin (b(1:N)) if (one == 1)
+
+#if !ACC_MEM_SHARED
+ if (acc_is_present (b) .eqv. .FALSE.) STOP 69
+#endif
+
+ !$acc exit data delete (b(1:N)) if (one == 1)
+
+#if !ACC_MEM_SHARED
+ if (acc_is_present (b) .eqv. .TRUE.) STOP 70
+#endif
+
+end program main
^ permalink raw reply [flat|nested] only message in thread
only message in thread, other threads:[~2023-10-25 9:25 UTC | newest]
Thread overview: (only message) (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2023-10-25 9:25 [gcc r14-4910] Extend test suite coverage for OpenACC 'self' clause for compute constructs Thomas Schwinge
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).