public inbox for gcc-cvs@sourceware.org
help / color / mirror / Atom feed
* [gcc/devel/omp/gcc-12] Fortran "declare create"/allocate support for OpenACC: adjust 'libgomp.oacc-fortran/declare-allocata
@ 2022-11-02 21:08 Thomas Schwinge
0 siblings, 0 replies; 2+ messages in thread
From: Thomas Schwinge @ 2022-11-02 21:08 UTC (permalink / raw)
To: gcc-cvs
https://gcc.gnu.org/g:5f1635e2e69fedaace1b7c77f2a02137b9c38ba4
commit 5f1635e2e69fedaace1b7c77f2a02137b9c38ba4
Author: Thomas Schwinge <thomas@codesourcery.com>
Date: Wed Nov 2 16:49:13 2022 +0100
Fortran "declare create"/allocate support for OpenACC: adjust 'libgomp.oacc-fortran/declare-allocatable-1*.f90'
libgomp/
* testsuite/libgomp.oacc-fortran/declare-allocatable-1-directive.f90:
Adjust.
* testsuite/libgomp.oacc-fortran/declare-allocatable-1-runtime.f90:
Likewise.
* testsuite/libgomp.oacc-fortran/declare-allocatable-1.f90:
Likewise.
Diff:
---
libgomp/ChangeLog.omp | 7 +++++++
.../libgomp.oacc-fortran/declare-allocatable-1-directive.f90 | 7 +++----
.../libgomp.oacc-fortran/declare-allocatable-1-runtime.f90 | 7 +++----
libgomp/testsuite/libgomp.oacc-fortran/declare-allocatable-1.f90 | 5 +----
4 files changed, 14 insertions(+), 12 deletions(-)
diff --git a/libgomp/ChangeLog.omp b/libgomp/ChangeLog.omp
index 38048785174..03867008500 100644
--- a/libgomp/ChangeLog.omp
+++ b/libgomp/ChangeLog.omp
@@ -1,5 +1,12 @@
2022-11-02 Thomas Schwinge <thomas@codesourcery.com>
+ * testsuite/libgomp.oacc-fortran/declare-allocatable-1-directive.f90:
+ Adjust.
+ * testsuite/libgomp.oacc-fortran/declare-allocatable-1-runtime.f90:
+ Likewise.
+ * testsuite/libgomp.oacc-fortran/declare-allocatable-1.f90:
+ Likewise.
+
* testsuite/libgomp.oacc-fortran/declare-allocatable-1-directive.f90:
XFAIL some OpenACC 'kernels' confusion.
* testsuite/libgomp.oacc-fortran/declare-allocatable-1-runtime.f90:
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/declare-allocatable-1-directive.f90 b/libgomp/testsuite/libgomp.oacc-fortran/declare-allocatable-1-directive.f90
index 6dd6f480519..6c5d6d87355 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/declare-allocatable-1-directive.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/declare-allocatable-1-directive.f90
@@ -2,11 +2,10 @@
! { dg-do run }
-!TODO-OpenACC-declare-allocate
-! Missing support for OpenACC "Changes from Version 2.0 to 2.5":
+! We've got support for OpenACC "Changes from Version 2.0 to 2.5":
! "The 'declare create' directive with a Fortran 'allocatable' has new behavior".
-! Thus, after 'allocate'/before 'deallocate', do
-! '!$acc enter data create'/'!$acc exit data delete' manually.
+! Yet, after 'allocate'/before 'deallocate', do
+! '!$acc enter data create'/'!$acc exit data delete' manually, too.
!TODO { dg-additional-options -fno-inline } for stable results regarding OpenACC 'routine'.
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/declare-allocatable-1-runtime.f90 b/libgomp/testsuite/libgomp.oacc-fortran/declare-allocatable-1-runtime.f90
index f86413b6b11..66a87e8850f 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/declare-allocatable-1-runtime.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/declare-allocatable-1-runtime.f90
@@ -2,11 +2,10 @@
! { dg-do run }
-!TODO-OpenACC-declare-allocate
-! Missing support for OpenACC "Changes from Version 2.0 to 2.5":
+! We've got support for OpenACC "Changes from Version 2.0 to 2.5":
! "The 'declare create' directive with a Fortran 'allocatable' has new behavior".
-! Thus, after 'allocate'/before 'deallocate', call 'acc_create'/'acc_delete'
-! manually.
+! Yet, after 'allocate'/before 'deallocate', call 'acc_create'/'acc_delete'
+! manually, too.
!TODO { dg-additional-options -fno-inline } for stable results regarding OpenACC 'routine'.
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/declare-allocatable-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/declare-allocatable-1.f90
index fb2ede13465..01275c69346 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/declare-allocatable-1.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/declare-allocatable-1.f90
@@ -2,11 +2,8 @@
! { dg-do run }
-!TODO-OpenACC-declare-allocate
-! Not currently implementing correct '-DACC_MEM_SHARED=0' behavior:
-! Missing support for OpenACC "Changes from Version 2.0 to 2.5":
+! We've got support for OpenACC "Changes from Version 2.0 to 2.5":
! "The 'declare create' directive with a Fortran 'allocatable' has new behavior".
-! { dg-xfail-run-if TODO { *-*-* } { -DACC_MEM_SHARED=0 } }
!TODO { dg-additional-options -fno-inline } for stable results regarding OpenACC 'routine'.
^ permalink raw reply [flat|nested] 2+ messages in thread
* [gcc/devel/omp/gcc-12] Fortran "declare create"/allocate support for OpenACC: adjust 'libgomp.oacc-fortran/declare-allocata
@ 2022-11-02 21:08 Thomas Schwinge
0 siblings, 0 replies; 2+ messages in thread
From: Thomas Schwinge @ 2022-11-02 21:08 UTC (permalink / raw)
To: gcc-cvs
https://gcc.gnu.org/g:f8a1c619be162d4446048f3d5678002859b561c7
commit f8a1c619be162d4446048f3d5678002859b561c7
Author: Thomas Schwinge <thomas@codesourcery.com>
Date: Wed Nov 2 16:50:33 2022 +0100
Fortran "declare create"/allocate support for OpenACC: adjust 'libgomp.oacc-fortran/declare-allocatable-array_descriptor-1*.f90'
libgomp/
* testsuite/libgomp.oacc-fortran/declare-allocatable-array_descriptor-1-directive.f90:
Adjust.
* testsuite/libgomp.oacc-fortran/declare-allocatable-array_descriptor-1-runtime.f90:
Likewise.
* testsuite/libgomp.oacc-fortran/declare-allocatable-array_descriptor-1.f90:
New.
Diff:
---
libgomp/ChangeLog.omp | 7 +
...re-allocatable-array_descriptor-1-directive.f90 | 41 +--
...lare-allocatable-array_descriptor-1-runtime.f90 | 107 ++++--
.../declare-allocatable-array_descriptor-1.f90 | 405 +++++++++++++++++++++
4 files changed, 499 insertions(+), 61 deletions(-)
diff --git a/libgomp/ChangeLog.omp b/libgomp/ChangeLog.omp
index 03867008500..72b74493685 100644
--- a/libgomp/ChangeLog.omp
+++ b/libgomp/ChangeLog.omp
@@ -1,5 +1,12 @@
2022-11-02 Thomas Schwinge <thomas@codesourcery.com>
+ * testsuite/libgomp.oacc-fortran/declare-allocatable-array_descriptor-1-directive.f90:
+ Adjust.
+ * testsuite/libgomp.oacc-fortran/declare-allocatable-array_descriptor-1-runtime.f90:
+ Likewise.
+ * testsuite/libgomp.oacc-fortran/declare-allocatable-array_descriptor-1.f90:
+ New.
+
* testsuite/libgomp.oacc-fortran/declare-allocatable-1-directive.f90:
Adjust.
* testsuite/libgomp.oacc-fortran/declare-allocatable-1-runtime.f90:
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/declare-allocatable-array_descriptor-1-directive.f90 b/libgomp/testsuite/libgomp.oacc-fortran/declare-allocatable-array_descriptor-1-directive.f90
index 6604f72c5c1..0f4d21a138b 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/declare-allocatable-array_descriptor-1-directive.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/declare-allocatable-array_descriptor-1-directive.f90
@@ -7,11 +7,10 @@
! host/device array descriptors.
! { dg-skip-if n/a { *-*-* } { -DACC_MEM_SHARED=1 } }
-!TODO-OpenACC-declare-allocate
-! Missing support for OpenACC "Changes from Version 2.0 to 2.5":
+! We've got support for OpenACC "Changes from Version 2.0 to 2.5":
! "The 'declare create' directive with a Fortran 'allocatable' has new behavior".
-! Thus, after 'allocate'/before 'deallocate', do
-! '!$acc enter data create'/'!$acc exit data delete' manually.
+! Yet, after 'allocate'/before 'deallocate', do
+! '!$acc enter data create'/'!$acc exit data delete' manually, too.
!TODO { dg-additional-options -fno-inline } for stable results regarding OpenACC 'routine'.
@@ -101,8 +100,6 @@ program test
allocate (b(n1_lb:n1_ub))
call verify_n1_allocated
- if (acc_is_present (b)) error stop
- !$acc enter data create (b)
! This is now OpenACC "present":
if (.not.acc_is_present (b)) error stop
! ..., and got the actual array descriptor installed:
@@ -110,15 +107,16 @@ program test
call verify_n1_allocated
!$acc end serial
+ !$acc enter data create (b)
+ if (.not.acc_is_present (b)) error stop
+ !$acc serial
+ call verify_n1_allocated
+ !$acc end serial
+
do i = n1_lb, n1_ub
b(i) = i - 1
end do
- ! In 'declare-allocatable-array_descriptor-1-runtime.f90', this does "verify
- ! that host-to-device copy doesn't touch the device-side (still initial)
- ! array descriptor (but it does copy the array data"). This is here not
- ! applicable anymore, as we've already gotten the actual array descriptor
- ! installed. Thus now verify that it does copy the array data.
call acc_update_device (b)
!$acc serial
call verify_n1_allocated
@@ -143,12 +141,6 @@ program test
!TODO 'GOMP_MAP_TO_PSET':
! { dg-final { scan-tree-dump-times {(?n)^ *#pragma omp target oacc_parallel map\(tofrom:MEM <integer\(kind=[0-9]+\)\[0:\]> \[\(integer\(kind=[0-9]+\)\[0:\] \*\)[^\]]+\] \[len: [^\]]+\]\) map\(alloc:b\.data \[pointer assign, bias: 0\]\) map\(from:id1_2 \[len: [0-9]+\]\)$} 1 gimple } }
- ! In 'declare-allocatable-array_descriptor-1-runtime.f90', this does "verify
- ! that device-to-host copy doesn't touch the host-side array descriptor,
- ! doesn't copy out the device-side (still initial) array descriptor (but it
- ! does copy the array data)". This is here not applicable anymore, as we've
- ! already gotten the actual array descriptor installed. Thus now verify that
- ! it does copy the array data.
call acc_update_self (b)
call verify_n1_allocated
@@ -223,14 +215,13 @@ program test
!$acc exit data delete (b)
if (.not.allocated (b)) error stop
- if (acc_is_present (b)) error stop
- ! The device-side array descriptor doesn't get updated, so 'b' still appears
- ! as "allocated":
+ if (.not.acc_is_present (b)) error stop
!$acc serial
call verify_n1_allocated
!$acc end serial
deallocate (b)
+ !if (acc_is_present (b)) error stop
call verify_n1_deallocated (.false.)
! The device-side array descriptor doesn't get updated, so 'b' still appears
! as "allocated":
@@ -260,10 +251,13 @@ program test
allocate (b(n2_lb:n2_ub))
call verify_n2_allocated
- if (acc_is_present (b)) error stop
+ if (.not.acc_is_present (b)) error stop
+ !$acc serial
+ call verify_n2_allocated
+ !$acc end serial
+
!$acc enter data create (b)
if (.not.acc_is_present (b)) error stop
- ! ..., and got the actual array descriptor installed:
!$acc serial
call verify_n2_allocated
!$acc end serial
@@ -337,12 +331,13 @@ program test
!$acc exit data delete (b)
if (.not.allocated (b)) error stop
- if (acc_is_present (b)) error stop
+ if (.not.acc_is_present (b)) error stop
!$acc serial
call verify_n2_allocated
!$acc end serial
deallocate (b)
+ !if (acc_is_present (b)) error stop
call verify_n2_deallocated (.false.)
!$acc serial
call verify_n2_allocated
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/declare-allocatable-array_descriptor-1-runtime.f90 b/libgomp/testsuite/libgomp.oacc-fortran/declare-allocatable-array_descriptor-1-runtime.f90
index b27f312631d..0682256dd91 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/declare-allocatable-array_descriptor-1-runtime.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/declare-allocatable-array_descriptor-1-runtime.f90
@@ -7,11 +7,10 @@
! host/device array descriptors.
! { dg-skip-if n/a { *-*-* } { -DACC_MEM_SHARED=1 } }
-!TODO-OpenACC-declare-allocate
-! Missing support for OpenACC "Changes from Version 2.0 to 2.5":
+! We've got support for OpenACC "Changes from Version 2.0 to 2.5":
! "The 'declare create' directive with a Fortran 'allocatable' has new behavior".
-! Thus, after 'allocate'/before 'deallocate', call 'acc_create'/'acc_delete'
-! manually.
+! Yet, after 'allocate'/before 'deallocate', call 'acc_create'/'acc_delete'
+! manually, too.
!TODO { dg-additional-options -fno-inline } for stable results regarding OpenACC 'routine'.
@@ -101,31 +100,47 @@ program test
allocate (b(n1_lb:n1_ub))
call verify_n1_allocated
- if (acc_is_present (b)) error stop
- call acc_create (b)
! This is now OpenACC "present":
if (.not.acc_is_present (b)) error stop
- ! This still has the initial array descriptor:
+ ! ..., and got the actual array descriptor installed:
!$acc serial
- call verify_initial
+ call verify_n1_allocated
+ !$acc end serial
+
+ call acc_create (b)
+ if (.not.acc_is_present (b)) error stop
+ !$acc serial
+ call verify_n1_allocated
!$acc end serial
do i = n1_lb, n1_ub
b(i) = i - 1
end do
- ! Verify that host-to-device copy doesn't touch the device-side (still
- ! initial) array descriptor (but it does copy the array data).
call acc_update_device (b)
!$acc serial
- call verify_initial
+ call verify_n1_allocated
!$acc end serial
b = 40
- ! Verify that device-to-host copy doesn't touch the host-side array
- ! descriptor, doesn't copy out the device-side (still initial) array
- ! descriptor (but it does copy the array data).
+ !$acc parallel copyout (id1_1) ! No data clause for 'b' (explicit or implicit): no 'GOMP_MAP_TO_PSET'.
+ call verify_n1_values (-1)
+ id1_1 = 0
+ !$acc end parallel
+ ! { dg-final { scan-tree-dump-times {(?n)^ *#pragma acc parallel map\(from:id1_1\)$} 1 original } }
+ ! { dg-final { scan-tree-dump-times {(?n)^ *#pragma omp target oacc_parallel map\(from:id1_1 \[len: [0-9]+\]\)$} 1 gimple } }
+
+ !$acc parallel copy (b) copyout (id1_2)
+ ! As already present, 'copy (b)' doesn't copy; addend is still '-1'.
+ call verify_n1_values (-1)
+ id1_2 = 0
+ !$acc end parallel
+ ! { dg-final { scan-tree-dump-times {(?n)^ *#pragma acc parallel map\(tofrom:\*\(integer\(kind=[0-9]+\)\[0:\] \* restrict\) b\.data \[len: [^\]]+\]\) map\(to:b \[pointer set, len: [0-9]+\]\) map\(alloc:\(integer\(kind=[0-9]+\)\[0:\] \* restrict\) b\.data \[pointer assign, bias: 0\]\) map\(from:id1_2\)$} 1 original } }
+ !TODO ..., but without an actual use of 'b', the gimplifier removes the
+ !TODO 'GOMP_MAP_TO_PSET':
+ ! { dg-final { scan-tree-dump-times {(?n)^ *#pragma omp target oacc_parallel map\(tofrom:MEM <integer\(kind=[0-9]+\)\[0:\]> \[\(integer\(kind=[0-9]+\)\[0:\] \*\)[^\]]+\] \[len: [^\]]+\]\) map\(alloc:b\.data \[pointer assign, bias: 0\]\) map\(from:id1_2 \[len: [0-9]+\]\)$} 1 gimple } }
+
call acc_update_self (b)
call verify_n1_allocated
@@ -142,11 +157,19 @@ program test
! { dg-final { scan-tree-dump-times {(?n)^ *#pragma omp target oacc_update map\(force_to:MEM <integer\(kind=[0-9]+\)\[0:\]> \[\(integer\(kind=[0-9]+\)\[0:\] \*\)[^\]]+\] \[len: [^\]]+\]\) map\(to:b \[pointer set, len: [0-9]+\]\) map\(alloc:b\.data \[pointer assign, bias: 0\]\) map\(force_from:id1_1 \[len: [0-9]+\]\)$} 1 gimple } }
! ..., but it's silently skipped in 'GOACC_update'.
!$acc serial
- call verify_initial
+ call verify_n1_allocated
!$acc end serial
b = 41
+ !$acc parallel
+ call verify_n1_values (1)
+ !$acc end parallel
+
+ !$acc parallel copy (b)
+ call verify_n1_values (1)
+ !$acc end parallel
+
!$acc update self (b) self (id1_2)
! We do have 'GOMP_MAP_TO_PSET' here:
! { dg-final { scan-tree-dump-times {(?n)^ *#pragma acc update map\(force_from:\*\(integer\(kind=[0-9]+\)\[0:\] \* restrict\) b\.data \[len: [^\]]+\]\) map\(to:b \[pointer set, len: [0-9]+\]\) map\(alloc:\(integer\(kind=[0-9]+\)\[0:\] \* restrict\) b\.data \[pointer assign, bias: 0\]\) map\(force_from:id1_2\);$} 1 original } }
@@ -159,20 +182,9 @@ program test
b(i) = b(i) + 2
end do
- ! Now install the actual array descriptor, via a data clause for 'b'
- ! (explicit or implicit): must get a 'GOMP_MAP_TO_PSET', which then in
- ! 'gomp_map_vars_internal' is handled as 'declare target', and because of
- ! '*(void **) hostaddrs[i] != NULL', we've got 'has_always_ptrset == true',
- ! 'always_to_cnt == 1', and therefore 'gomp_map_vars_existing' does update
- ! the 'GOMP_MAP_TO_PSET'.
- !$acc serial present (b) copyin (id1_1)
- call verify_initial
- id1_1 = 0
- !$acc end serial
- ! { dg-final { scan-tree-dump-times {(?n)^ *#pragma acc serial map\(force_present:\*\(integer\(kind=[0-9]+\)\[0:\] \* restrict\) b\.data \[len: [^\]]+\]\) map\(to:b \[pointer set, len: [0-9]+\]\) map\(alloc:\(integer\(kind=[0-9]+\)\[0:\] \* restrict\) b\.data \[pointer assign, bias: 0\]\) map\(to:id1_1\)$} 1 original } }
- !TODO ..., but without an actual use of 'b', the gimplifier removes the
- !TODO 'GOMP_MAP_TO_PSET':
- ! { dg-final { scan-tree-dump-times {(?n)^ *#pragma omp target oacc_serial map\(force_present:MEM <integer\(kind=[0-9]+\)\[0:\]> \[\(integer\(kind=[0-9]+\)\[0:\] \*\)[^\]]+\] \[len: [^\]]+\]\) map\(alloc:b\.data \[pointer assign, bias: 0\]\) map\(to:id1_1 \[len: [0-9]+\]\)$} 1 gimple } }
+ ! Now test that (potentially re-)installing the actual array descriptor is a
+ ! no-op, via a data clause for 'b' (explicit or implicit): must get a
+ ! 'GOMP_MAP_TO_PSET'.
!$acc serial present (b) copyin (id1_2)
call verify_n1_allocated
!TODO Use of 'b':
@@ -203,14 +215,13 @@ program test
call acc_delete (b)
if (.not.allocated (b)) error stop
- if (acc_is_present (b)) error stop
- ! The device-side array descriptor doesn't get updated, so 'b' still appears
- ! as "allocated":
+ if (.not.acc_is_present (b)) error stop
!$acc serial
call verify_n1_allocated
!$acc end serial
deallocate (b)
+ !if (acc_is_present (b)) error stop
call verify_n1_deallocated (.false.)
! The device-side array descriptor doesn't get updated, so 'b' still appears
! as "allocated":
@@ -240,12 +251,15 @@ program test
allocate (b(n2_lb:n2_ub))
call verify_n2_allocated
- if (acc_is_present (b)) error stop
+ if (.not.acc_is_present (b)) error stop
+ !$acc serial
+ call verify_n2_allocated
+ !$acc end serial
+
call acc_create (b)
if (.not.acc_is_present (b)) error stop
- ! This still has the previous (n1) array descriptor:
!$acc serial
- call verify_n1_deallocated (.true.)
+ call verify_n2_allocated
!$acc end serial
do i = n2_lb, n2_ub
@@ -254,11 +268,19 @@ program test
call acc_update_device (b)
!$acc serial
- call verify_n1_deallocated (.true.)
+ call verify_n2_allocated
!$acc end serial
b = -40
+ !$acc parallel
+ call verify_n2_values (20)
+ !$acc end parallel
+
+ !$acc parallel copy (b)
+ call verify_n2_values (20)
+ !$acc end parallel
+
call acc_update_self (b)
call verify_n2_allocated
@@ -269,11 +291,19 @@ program test
!$acc update device (b)
!$acc serial
- call verify_n1_deallocated (.true.)
+ call verify_n2_allocated
!$acc end serial
b = -41
+ !$acc parallel
+ call verify_n2_values (-20)
+ !$acc end parallel
+
+ !$acc parallel copy (b)
+ call verify_n2_values (-20)
+ !$acc end parallel
+
!$acc update self (b)
call verify_n2_allocated
@@ -301,12 +331,13 @@ program test
call acc_delete (b)
if (.not.allocated (b)) error stop
- if (acc_is_present (b)) error stop
+ if (.not.acc_is_present (b)) error stop
!$acc serial
call verify_n2_allocated
!$acc end serial
deallocate (b)
+ !if (acc_is_present (b)) error stop
call verify_n2_deallocated (.false.)
!$acc serial
call verify_n2_allocated
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/declare-allocatable-array_descriptor-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/declare-allocatable-array_descriptor-1.f90
new file mode 100644
index 00000000000..1105a578618
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/declare-allocatable-array_descriptor-1.f90
@@ -0,0 +1,405 @@
+! Test OpenACC 'declare create' with allocatable arrays.
+
+! { dg-do run }
+
+! Note that we're not testing OpenACC semantics here, but rather documenting
+! current GCC behavior, specifically, behavior concerning updating of
+! host/device array descriptors.
+! { dg-skip-if n/a { *-*-* } { -DACC_MEM_SHARED=1 } }
+
+! We've got support for OpenACC "Changes from Version 2.0 to 2.5":
+! "The 'declare create' directive with a Fortran 'allocatable' has new behavior".
+
+
+!TODO { dg-additional-options -fno-inline } for stable results regarding OpenACC 'routine'.
+
+
+!TODO OpenACC 'serial' vs. GCC/nvptx:
+!TODO { dg-prune-output {using 'vector_length \(32\)', ignoring 1} }
+
+
+! { dg-additional-options -fdump-tree-original }
+! { dg-additional-options -fdump-tree-gimple }
+
+
+module vars
+ implicit none
+ integer, parameter :: n1_lb = -3
+ integer, parameter :: n1_ub = 6
+ integer, parameter :: n2_lb = -9999
+ integer, parameter :: n2_ub = 22222
+
+ integer, allocatable :: b(:)
+ !$acc declare create (b)
+
+end module vars
+
+program test
+ use vars
+ use openacc
+ implicit none
+ integer :: i
+
+ ! Identifiers for purposes of reliable '-fdump-tree-[...]' scanning.
+ integer :: id1_1, id1_2
+
+ interface
+
+ subroutine verify_initial
+ implicit none
+ !$acc routine seq
+ end subroutine verify_initial
+
+ subroutine verify_n1_allocated
+ implicit none
+ !$acc routine seq
+ end subroutine verify_n1_allocated
+
+ subroutine verify_n1_values (addend)
+ implicit none
+ !$acc routine gang
+ integer, value :: addend
+ end subroutine verify_n1_values
+
+ subroutine verify_n1_deallocated (expect_allocated)
+ implicit none
+ !$acc routine seq
+ logical, value :: expect_allocated
+ end subroutine verify_n1_deallocated
+
+ subroutine verify_n2_allocated
+ implicit none
+ !$acc routine seq
+ end subroutine verify_n2_allocated
+
+ subroutine verify_n2_values (addend)
+ implicit none
+ !$acc routine gang
+ integer, value :: addend
+ end subroutine verify_n2_values
+
+ subroutine verify_n2_deallocated (expect_allocated)
+ implicit none
+ !$acc routine seq
+ logical, value :: expect_allocated
+ end subroutine verify_n2_deallocated
+
+ end interface
+
+ call acc_create (id1_1)
+ call acc_create (id1_2)
+
+ call verify_initial
+ ! It is important here (and similarly, following) that there is no data
+ ! clause for 'b' (explicit or implicit): no 'GOMP_MAP_TO_PSET'.
+ !$acc serial
+ call verify_initial
+ !$acc end serial
+
+ allocate (b(n1_lb:n1_ub))
+ call verify_n1_allocated
+ ! This is now OpenACC "present":
+ if (.not.acc_is_present (b)) error stop
+ ! ..., and got the actual array descriptor installed:
+ !$acc serial
+ call verify_n1_allocated
+ !$acc end serial
+
+ do i = n1_lb, n1_ub
+ b(i) = i - 1
+ end do
+
+ call acc_update_device (b)
+ !$acc serial
+ call verify_n1_allocated
+ !$acc end serial
+
+ b = 40
+
+ !$acc parallel copyout (id1_1) ! No data clause for 'b' (explicit or implicit): no 'GOMP_MAP_TO_PSET'.
+ call verify_n1_values (-1)
+ id1_1 = 0
+ !$acc end parallel
+ ! { dg-final { scan-tree-dump-times {(?n)^ *#pragma acc parallel map\(from:id1_1\)$} 1 original } }
+ ! { dg-final { scan-tree-dump-times {(?n)^ *#pragma omp target oacc_parallel map\(from:id1_1 \[len: [0-9]+\]\)$} 1 gimple } }
+
+ !$acc parallel copy (b) copyout (id1_2)
+ ! As already present, 'copy (b)' doesn't copy; addend is still '-1'.
+ call verify_n1_values (-1)
+ id1_2 = 0
+ !$acc end parallel
+ ! { dg-final { scan-tree-dump-times {(?n)^ *#pragma acc parallel map\(tofrom:\*\(integer\(kind=[0-9]+\)\[0:\] \* restrict\) b\.data \[len: [^\]]+\]\) map\(to:b \[pointer set, len: [0-9]+\]\) map\(alloc:\(integer\(kind=[0-9]+\)\[0:\] \* restrict\) b\.data \[pointer assign, bias: 0\]\) map\(from:id1_2\)$} 1 original } }
+ !TODO ..., but without an actual use of 'b', the gimplifier removes the
+ !TODO 'GOMP_MAP_TO_PSET':
+ ! { dg-final { scan-tree-dump-times {(?n)^ *#pragma omp target oacc_parallel map\(tofrom:MEM <integer\(kind=[0-9]+\)\[0:\]> \[\(integer\(kind=[0-9]+\)\[0:\] \*\)[^\]]+\] \[len: [^\]]+\]\) map\(alloc:b\.data \[pointer assign, bias: 0\]\) map\(from:id1_2 \[len: [0-9]+\]\)$} 1 gimple } }
+
+ call acc_update_self (b)
+ call verify_n1_allocated
+
+ do i = n1_lb, n1_ub
+ if (b(i) /= i - 1) error stop
+ b(i) = b(i) + 2
+ end do
+
+ ! The same using the OpenACC 'update' directive.
+
+ !$acc update device (b) self (id1_1)
+ ! We do have 'GOMP_MAP_TO_PSET' here:
+ ! { dg-final { scan-tree-dump-times {(?n)^ *#pragma acc update map\(force_to:\*\(integer\(kind=[0-9]+\)\[0:\] \* restrict\) b\.data \[len: [^\]]+\]\) map\(to:b \[pointer set, len: [0-9]+\]\) map\(alloc:\(integer\(kind=[0-9]+\)\[0:\] \* restrict\) b\.data \[pointer assign, bias: 0\]\) map\(force_from:id1_1\);$} 1 original } }
+ ! { dg-final { scan-tree-dump-times {(?n)^ *#pragma omp target oacc_update map\(force_to:MEM <integer\(kind=[0-9]+\)\[0:\]> \[\(integer\(kind=[0-9]+\)\[0:\] \*\)[^\]]+\] \[len: [^\]]+\]\) map\(to:b \[pointer set, len: [0-9]+\]\) map\(alloc:b\.data \[pointer assign, bias: 0\]\) map\(force_from:id1_1 \[len: [0-9]+\]\)$} 1 gimple } }
+ ! ..., but it's silently skipped in 'GOACC_update'.
+ !$acc serial
+ call verify_n1_allocated
+ !$acc end serial
+
+ b = 41
+
+ !$acc parallel
+ call verify_n1_values (1)
+ !$acc end parallel
+
+ !$acc parallel copy (b)
+ call verify_n1_values (1)
+ !$acc end parallel
+
+ !$acc update self (b) self (id1_2)
+ ! We do have 'GOMP_MAP_TO_PSET' here:
+ ! { dg-final { scan-tree-dump-times {(?n)^ *#pragma acc update map\(force_from:\*\(integer\(kind=[0-9]+\)\[0:\] \* restrict\) b\.data \[len: [^\]]+\]\) map\(to:b \[pointer set, len: [0-9]+\]\) map\(alloc:\(integer\(kind=[0-9]+\)\[0:\] \* restrict\) b\.data \[pointer assign, bias: 0\]\) map\(force_from:id1_2\);$} 1 original } }
+ ! { dg-final { scan-tree-dump-times {(?n)^ *#pragma omp target oacc_update map\(force_from:MEM <integer\(kind=[0-9]+\)\[0:\]> \[\(integer\(kind=[0-9]+\)\[0:\] \*\)[^\]]+\] \[len: [^\]]+\]\) map\(to:b \[pointer set, len: [0-9]+\]\) map\(alloc:b\.data \[pointer assign, bias: 0\]\) map\(force_from:id1_2 \[len: [0-9]+\]\)$} 1 gimple } }
+ ! ..., but it's silently skipped in 'GOACC_update'.
+ call verify_n1_allocated
+
+ do i = n1_lb, n1_ub
+ if (b(i) /= i + 1) error stop
+ b(i) = b(i) + 2
+ end do
+
+ ! Now test that (potentially re-)installing the actual array descriptor is a
+ ! no-op, via a data clause for 'b' (explicit or implicit): must get a
+ ! 'GOMP_MAP_TO_PSET'.
+ !$acc serial present (b) copyin (id1_2)
+ call verify_n1_allocated
+ !TODO Use of 'b':
+ id1_2 = ubound (b, 1)
+ !$acc end serial
+ ! { dg-final { scan-tree-dump-times {(?n)^ *#pragma acc serial map\(force_present:\*\(integer\(kind=[0-9]+\)\[0:\] \* restrict\) b\.data \[len: [^\]]+\]\) map\(to:b \[pointer set, len: [0-9]+\]\) map\(alloc:\(integer\(kind=[0-9]+\)\[0:\] \* restrict\) b\.data \[pointer assign, bias: 0\]\) map\(to:id1_2\)$} 1 original } }
+ ! { dg-final { scan-tree-dump-times {(?n)^ *#pragma omp target oacc_serial map\(force_present:MEM <integer\(kind=[0-9]+\)\[0:\]> \[\(integer\(kind=[0-9]+\)\[0:\] \*\)[^\]]+\] \[len: [^\]]+\]\) map\(to:b \[pointer set, len: [0-9]+\]\) map\(alloc:b\.data \[pointer assign, bias: 0\]\) map\(to:id1_2 \[len: [0-9]+\]\)$} 1 gimple } }
+
+ !$acc parallel copyin (id1_1) ! No data clause for 'b' (explicit or implicit): no 'GOMP_MAP_TO_PSET'.
+ call verify_n1_values (1)
+ id1_1 = 0
+ !$acc end parallel
+ ! { dg-final { scan-tree-dump-times {(?n)^ *#pragma acc parallel map\(to:id1_1\)$} 1 original } }
+ ! { dg-final { scan-tree-dump-times {(?n)^ *#pragma omp target oacc_parallel map\(to:id1_1 \[len: [0-9]+\]\)$} 1 gimple } }
+
+ !$acc parallel copy (b) copyin (id1_2)
+ ! As already present, 'copy (b)' doesn't copy; addend is still '1'.
+ call verify_n1_values (1)
+ id1_2 = 0
+ !$acc end parallel
+ ! { dg-final { scan-tree-dump-times {(?n)^ *#pragma acc parallel map\(tofrom:\*\(integer\(kind=[0-9]+\)\[0:\] \* restrict\) b\.data \[len: [^\]]+\]\) map\(to:b \[pointer set, len: [0-9]+\]\) map\(alloc:\(integer\(kind=[0-9]+\)\[0:\] \* restrict\) b\.data \[pointer assign, bias: 0\]\) map\(to:id1_2\)$} 1 original } }
+ !TODO ..., but without an actual use of 'b', the gimplifier removes the
+ !TODO 'GOMP_MAP_TO_PSET':
+ ! { dg-final { scan-tree-dump-times {(?n)^ *#pragma omp target oacc_parallel map\(tofrom:MEM <integer\(kind=[0-9]+\)\[0:\]> \[\(integer\(kind=[0-9]+\)\[0:\] \*\)[^\]]+\] \[len: [^\]]+\]\) map\(alloc:b\.data \[pointer assign, bias: 0\]\) map\(to:id1_2 \[len: [0-9]+\]\)$} 1 gimple } }
+
+ call verify_n1_allocated
+ if (.not.acc_is_present (b)) error stop
+
+ deallocate (b)
+ !if (acc_is_present (b)) error stop
+ call verify_n1_deallocated (.false.)
+ ! The device-side array descriptor doesn't get updated, so 'b' still appears
+ ! as "allocated":
+ !$acc serial
+ call verify_n1_allocated
+ !$acc end serial
+
+ ! Now try to install the actual array descriptor, via a data clause for 'b'
+ ! (explicit or implicit): must get a 'GOMP_MAP_TO_PSET', which then in
+ ! 'gomp_map_vars_internal' is handled as 'declare target', but because of
+ ! '*(void **) hostaddrs[i] == NULL', we've got 'has_always_ptrset == false',
+ ! 'always_to_cnt == 0', and therefore 'gomp_map_vars_existing' doesn't update
+ ! the 'GOMP_MAP_TO_PSET'.
+ ! The device-side array descriptor doesn't get updated, so 'b' still appears
+ ! as "allocated":
+ !TODO Why does 'present (b)' still work here?
+ !$acc serial present (b) copyout (id1_2)
+ call verify_n1_deallocated (.true.)
+ !TODO Use of 'b'.
+ id1_2 = ubound (b, 1)
+ !$acc end serial
+ ! { dg-final { scan-tree-dump-times {(?n)^ *#pragma acc serial map\(force_present:\*\(integer\(kind=[0-9]+\)\[0:\] \* restrict\) b\.data \[len: [^\]]+\]\) map\(to:b \[pointer set, len: [0-9]+\]\) map\(alloc:\(integer\(kind=[0-9]+\)\[0:\] \* restrict\) b\.data \[pointer assign, bias: 0\]\) map\(from:id1_2\)$} 1 original } }
+ ! { dg-final { scan-tree-dump-times {(?n)^ *#pragma omp target oacc_serial map\(force_present:MEM <integer\(kind=[0-9]+\)\[0:\]> \[\(integer\(kind=[0-9]+\)\[0:\] \*\)[^\]]+\] \[len: [^\]]+\]\) map\(to:b \[pointer set, len: [0-9]+\]\) map\(alloc:b\.data \[pointer assign, bias: 0\]\) map\(from:id1_2 \[len: [0-9]+\]\)$} 1 gimple } }
+
+
+ ! Restart the procedure, with different array dimensions.
+
+ allocate (b(n2_lb:n2_ub))
+ call verify_n2_allocated
+ if (.not.acc_is_present (b)) error stop
+ !$acc serial
+ call verify_n2_allocated
+ !$acc end serial
+
+ do i = n2_lb, n2_ub
+ b(i) = i + 20
+ end do
+
+ call acc_update_device (b)
+ !$acc serial
+ call verify_n2_allocated
+ !$acc end serial
+
+ b = -40
+
+ !$acc parallel
+ call verify_n2_values (20)
+ !$acc end parallel
+
+ !$acc parallel copy (b)
+ call verify_n2_values (20)
+ !$acc end parallel
+
+ call acc_update_self (b)
+ call verify_n2_allocated
+
+ do i = n2_lb, n2_ub
+ if (b(i) /= i + 20) error stop
+ b(i) = b(i) - 40
+ end do
+
+ !$acc update device (b)
+ !$acc serial
+ call verify_n2_allocated
+ !$acc end serial
+
+ b = -41
+
+ !$acc parallel
+ call verify_n2_values (-20)
+ !$acc end parallel
+
+ !$acc parallel copy (b)
+ call verify_n2_values (-20)
+ !$acc end parallel
+
+ !$acc update self (b)
+ call verify_n2_allocated
+
+ do i = n2_lb, n2_ub
+ if (b(i) /= i - 20) error stop
+ b(i) = b(i) + 10
+ end do
+
+ !$acc serial present (b) copy (id1_2)
+ call verify_n2_allocated
+ !TODO Use of 'b':
+ id1_2 = ubound (b, 1)
+ !$acc end serial
+
+ !$acc parallel
+ call verify_n2_values (-20)
+ !$acc end parallel
+
+ !$acc parallel copy (b)
+ call verify_n2_values (-20)
+ !$acc end parallel
+
+ call verify_n2_allocated
+ if (.not.acc_is_present (b)) error stop
+
+ deallocate (b)
+ !if (acc_is_present (b)) error stop
+ call verify_n2_deallocated (.false.)
+ !$acc serial
+ call verify_n2_allocated
+ !$acc end serial
+
+ !$acc serial present (b) copy (id1_2)
+ call verify_n2_deallocated (.true.)
+ !TODO Use of 'b':
+ id1_2 = ubound (b, 1)
+ !$acc end serial
+
+end program test
+
+
+subroutine verify_initial
+ use vars
+ implicit none
+ !$acc routine seq
+
+ if (allocated (b)) error stop "verify_initial allocated"
+ if (any (lbound (b) /= [0])) error stop "verify_initial lbound"
+ if (any (ubound (b) /= [0])) error stop "verify_initial ubound"
+end subroutine verify_initial
+
+subroutine verify_n1_allocated
+ use vars
+ implicit none
+ !$acc routine seq
+
+ if (.not.allocated (b)) error stop "verify_n1_allocated allocated"
+ if (any (lbound (b) /= [n1_lb])) error stop "verify_n1_allocated lbound"
+ if (any (ubound (b) /= [n1_ub])) error stop "verify_n1_allocated ubound"
+end subroutine verify_n1_allocated
+
+subroutine verify_n1_values (addend)
+ use vars
+ implicit none
+ !$acc routine gang
+ integer, value :: addend
+ integer :: i
+
+ !$acc loop
+ do i = n1_lb, n1_ub
+ if (b(i) /= i + addend) error stop
+ end do
+end subroutine verify_n1_values
+
+subroutine verify_n1_deallocated (expect_allocated)
+ use vars
+ implicit none
+ !$acc routine seq
+ logical, value :: expect_allocated
+
+ if (allocated(b) .neqv. expect_allocated) error stop "verify_n1_deallocated allocated"
+ ! Apparently 'deallocate'ing doesn't unset the bounds.
+ if (any (lbound (b) /= [n1_lb])) error stop "verify_n1_deallocated lbound"
+ if (any (ubound (b) /= [n1_ub])) error stop "verify_n1_deallocated ubound"
+end subroutine verify_n1_deallocated
+
+subroutine verify_n2_allocated
+ use vars
+ implicit none
+ !$acc routine seq
+
+ if (.not.allocated(b)) error stop "verify_n2_allocated allocated"
+ if (any (lbound (b) /= [n2_lb])) error stop "verify_n2_allocated lbound"
+ if (any (ubound (b) /= [n2_ub])) error stop "verify_n2_allocated ubound"
+end subroutine verify_n2_allocated
+
+subroutine verify_n2_values (addend)
+ use vars
+ implicit none
+ !$acc routine gang
+ integer, value :: addend
+ integer :: i
+
+ !$acc loop
+ do i = n2_lb, n2_ub
+ if (b(i) /= i + addend) error stop
+ end do
+end subroutine verify_n2_values
+
+subroutine verify_n2_deallocated (expect_allocated)
+ use vars
+ implicit none
+ !$acc routine seq
+ logical, value :: expect_allocated
+
+ if (allocated(b) .neqv. expect_allocated) error stop "verify_n2_deallocated allocated"
+ ! Apparently 'deallocate'ing doesn't unset the bounds.
+ if (any (lbound (b) /= [n2_lb])) error stop "verify_n2_deallocated lbound"
+ if (any (ubound (b) /= [n2_ub])) error stop "verify_n2_deallocated ubound"
+end subroutine verify_n2_deallocated
^ permalink raw reply [flat|nested] 2+ messages in thread
end of thread, other threads:[~2022-11-02 21:08 UTC | newest]
Thread overview: 2+ messages (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2022-11-02 21:08 [gcc/devel/omp/gcc-12] Fortran "declare create"/allocate support for OpenACC: adjust 'libgomp.oacc-fortran/declare-allocata Thomas Schwinge
2022-11-02 21:08 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).