mirror of
https://forge.sourceware.org/marek/gcc.git
synced 2026-02-22 03:47:02 -05:00
This patch changes the mapping node arrangement used for array components of derived types in order to accommodate for changes made in the previous patch, particularly the use of "GOMP_MAP_ATTACH_DETACH" for pointer-typed derived-type members instead of "GOMP_MAP_ALWAYS_POINTER". We change the mapping nodes used for a derived-type mapping like this: type T integer, pointer, dimension(:) :: arrptr end type T type(T) :: tvar [...] !$omp target map(tofrom: tvar%arrptr) So that the nodes used look like this: 1) map(to: tvar%arrptr) --> GOMP_MAP_TO [implicit] *tvar%arrptr%data (the array data) GOMP_MAP_TO_PSET tvar%arrptr (the descriptor) GOMP_MAP_ATTACH_DETACH tvar%arrptr%data 2) map(tofrom: tvar%arrptr(3:8) --> GOMP_MAP_TOFROM *tvar%arrptr%data(3) (size 8-3+1, etc.) GOMP_MAP_TO_PSET tvar%arrptr GOMP_MAP_ATTACH_DETACH tvar%arrptr%data (bias 3, etc.) In this case, we can determine in the front-end that the whole-array/pointer mapping (1) is only needed to map the pointer -- so we drop it entirely. (Note also that we set -- early -- the OMP_CLAUSE_MAP_RUNTIME_IMPLICIT_P flag for whole-array-via-pointer mappings. See below.) In the middle end, we process mappings using the struct sibling-list handling machinery by moving the "GOMP_MAP_TO_PSET" node from the middle of the group of three mapping nodes to the proper sorted position after the GOMP_MAP_STRUCT mapping: GOMP_MAP_STRUCT tvar (len: 1) GOMP_MAP_TO_PSET tvar%arr (size: 64, etc.) <--. moved here [...] | GOMP_MAP_TOFROM *tvar%arrptr%data(3) ___| GOMP_MAP_ATTACH_DETACH tvar%arrptr%data In another case, if we have an array of derived-type values "dtarr", and mappings like: i = 1 j = 1 map(to: dtarr(i)%arrptr) map(tofrom: dtarr(j)%arrptr(3:8)) We still map the same way, but this time we cannot prove that the base expressions "dtarr(i) and "dtarr(j)" are the same in the front-end. So we keep both mappings, but we move the "[implicit]" mapping of the full-array reference to the end of the clause list in gimplify.cc (by adjusting the topological sorting algorithm): GOMP_MAP_STRUCT dtvar (len: 2) GOMP_MAP_TO_PSET dtvar(i)%arrptr GOMP_MAP_TO_PSET dtvar(j)%arrptr [...] GOMP_MAP_TOFROM *dtvar(j)%arrptr%data(3) (size: 8-3+1) GOMP_MAP_ATTACH_DETACH dtvar(j)%arrptr%data GOMP_MAP_TO [implicit] *dtvar(i)%arrptr%data(1) (size: whole array) GOMP_MAP_ATTACH_DETACH dtvar(i)%arrptr%data Always moving "[implicit]" full-array mappings after array-section mappings (without that bit set) means that we'll avoid copying the whole array unnecessarily -- even in cases where we can't prove that the arrays are the same. The patch also fixes some bugs with "enter data" and "exit data" directives with this new mapping arrangement. Also now if you have mappings like this: #pragma omp target enter data map(to: dv, dv%arr(1:20)) The whole of the derived-type variable "dv" is mapped, so the GOMP_MAP_TO_PSET for the array-section mapping can be dropped: GOMP_MAP_TO dv GOMP_MAP_TO *dv%arr%data GOMP_MAP_TO_PSET dv%arr <-- deleted (array section mapping) GOMP_MAP_ATTACH_DETACH dv%arr%data To accommodate for recent changes to mapping nodes made by Tobias, this version of the patch avoids using GOMP_MAP_TO_PSET for "exit data" directives, in favour of using the "correct" GOMP_MAP_RELEASE/GOMP_MAP_DELETE kinds during early expansion. A new flag is introduced so the middle-end knows when the latter two kinds are being used specifically for an array descriptor. This version of the patch fixes "omp target exit data" handling for GOMP_MAP_DELETE, and adds pretty-printing dump output for the OMP_CLAUSE_RELEASE_DESCRIPTOR flag (for a little extra clarity). Also I noticed the handling of descriptors on *OpenACC* exit-data directives was inconsistent, so I've made those use GOMP_MAP_RELEASE/GOMP_MAP_DELETE with the new flag in the same way as OpenMP too. In the end it doesn't actually matter to the runtime, which handles GOMP_MAP_RELEASE/GOMP_MAP_DELETE/GOMP_MAP_TO_PSET for array descriptors on OpenACC "exit data" directives the same, anyway, and doing it this way in the FE avoids needless divergence. I've added a couple of new tests (gomp/target-enter-exit-data.f90 and goacc/enter-exit-data-2.f90). 2023-12-07 Julian Brown <julian@codesourcery.com> gcc/fortran/ * dependency.cc (gfc_omp_expr_prefix_same): New function. * dependency.h (gfc_omp_expr_prefix_same): Add prototype. * gfortran.h (gfc_omp_namelist): Add "duplicate_of" field to "u2" union. * trans-openmp.cc (dependency.h): Include. (gfc_trans_omp_array_section): Adjust mapping node arrangement for array descriptors. Use GOMP_MAP_TO_PSET or GOMP_MAP_RELEASE/GOMP_MAP_DELETE with the OMP_CLAUSE_RELEASE_DESCRIPTOR flag set. (gfc_symbol_rooted_namelist): New function. (gfc_trans_omp_clauses): Check subcomponent and subarray/element accesses elsewhere in the clause list for pointers to derived types or array descriptors, and adjust or drop mapping nodes appropriately. Adjust for changes to mapping node arrangement. (gfc_trans_oacc_executable_directive): Pass code op through. gcc/ * gimplify.cc (omp_map_clause_descriptor_p): New function. (build_omp_struct_comp_nodes, omp_get_attachment, omp_group_base): Use above function. (omp_tsort_mapping_groups): Process nodes that have OMP_CLAUSE_MAP_RUNTIME_IMPLICIT_P set after those that don't. Add enter_exit_data parameter. (omp_resolve_clause_dependencies): Remove GOMP_MAP_TO_PSET mappings if we're mapping the whole containing derived-type variable. (omp_accumulate_sibling_list): Adjust GOMP_MAP_TO_PSET handling. Remove GOMP_MAP_ALWAYS_POINTER handling. (gimplify_scan_omp_clauses): Pass enter_exit argument to omp_tsort_mapping_groups. Don't adjust/remove GOMP_MAP_TO_PSET mappings for derived-type components here. * tree.h (OMP_CLAUSE_RELEASE_DESCRIPTOR): New macro. * tree-pretty-print.cc (dump_omp_clause): Show OMP_CLAUSE_RELEASE_DESCRIPTOR in dump output (with GOMP_MAP_TO_PSET-like syntax). gcc/testsuite/ * gfortran.dg/goacc/enter-exit-data-2.f90: New test. * gfortran.dg/goacc/finalize-1.f: Adjust scan output. * gfortran.dg/gomp/map-9.f90: Adjust scan output. * gfortran.dg/gomp/map-subarray-2.f90: New test. * gfortran.dg/gomp/map-subarray.f90: New test. * gfortran.dg/gomp/target-enter-exit-data.f90: New test. libgomp/ * testsuite/libgomp.fortran/map-subarray.f90: New test. * testsuite/libgomp.fortran/map-subarray-2.f90: New test. * testsuite/libgomp.fortran/map-subarray-3.f90: New test. * testsuite/libgomp.fortran/map-subarray-4.f90: New test. * testsuite/libgomp.fortran/map-subarray-6.f90: New test. * testsuite/libgomp.fortran/map-subarray-7.f90: New test. * testsuite/libgomp.fortran/map-subarray-8.f90: New test. * testsuite/libgomp.fortran/map-subcomponents.f90: New test. * testsuite/libgomp.fortran/struct-elem-map-1.f90: Adjust for descriptor-mapping changes. Remove XFAIL.
592 lines
21 KiB
Fortran
592 lines
21 KiB
Fortran
! { dg-do run }
|
||
!
|
||
! Test OpenMP 4.5 structure-element mapping
|
||
|
||
! TODO: ...%str4 + %uni4 should be tested but that currently fails due to
|
||
! PR fortran/95868 (see commented lined)
|
||
! TODO: Test also 'var' as array and/or pointer; nested derived types,
|
||
! type-extended types.
|
||
|
||
program main
|
||
implicit none
|
||
|
||
type t2
|
||
integer :: a, b
|
||
! For complex, assume small integers are exactly representable
|
||
complex(kind=8) :: c
|
||
integer :: d(10)
|
||
integer, pointer :: e => null(), f(:) => null()
|
||
character(len=5) :: str1
|
||
character(len=5) :: str2(4)
|
||
character(len=:), pointer :: str3 => null()
|
||
character(len=:), pointer :: str4(:) => null()
|
||
character(kind=4, len=5) :: uni1
|
||
character(kind=4, len=5) :: uni2(4)
|
||
character(kind=4, len=:), pointer :: uni3 => null()
|
||
character(kind=4, len=:), pointer :: uni4(:) => null()
|
||
end type t2
|
||
|
||
integer :: i
|
||
|
||
call one ()
|
||
call two ()
|
||
call three ()
|
||
call four ()
|
||
call five ()
|
||
call six ()
|
||
call seven ()
|
||
call eight ()
|
||
call nine ()
|
||
call ten ()
|
||
call eleven ()
|
||
call twelve ()
|
||
|
||
contains
|
||
! Implicitly mapped – but no pointers are mapped
|
||
subroutine one()
|
||
type(t2) :: var
|
||
|
||
print '(g0)', '==== TESTCASE "one" ===='
|
||
|
||
var = t2(a = 1, &
|
||
b = 2, c = cmplx(-1.0_8, 2.0_8,kind=8), &
|
||
d = [(-3*i, i = 1, 10)], &
|
||
str1 = "abcde", &
|
||
str2 = ["12345", "67890", "ABCDE", "FGHIJ"], &
|
||
uni1 = 4_"abcde", &
|
||
uni2 = [4_"12345", 4_"67890", 4_"ABCDE", 4_"FGHIJ"])
|
||
allocate (var%e, source=99)
|
||
allocate (var%f, source=[22, 33, 44, 55])
|
||
allocate (var%str3, source="HelloWorld")
|
||
allocate (var%str4, source=["Let's", "Go!!!"])
|
||
allocate (var%uni3, source=4_"HelloWorld")
|
||
allocate (var%uni4, source=[4_"Let's", 4_"Go!!!"])
|
||
|
||
!$omp target map(tofrom:var)
|
||
if (var%a /= 1) stop 1
|
||
if (var%b /= 2) stop 2
|
||
if (var%c%re /= -1.0_8 .or. var%c%im /= 2.0_8) stop 3
|
||
if (any (var%d /= [(-3*i, i = 1, 10)])) stop 4
|
||
if (var%str1 /= "abcde") stop 5
|
||
if (any (var%str2 /= ["12345", "67890", "ABCDE", "FGHIJ"])) stop 6
|
||
if (var%uni1 /= 4_"abcde") stop 7
|
||
if (any (var%uni2 /= [4_"12345", 4_"67890", 4_"ABCDE", 4_"FGHIJ"])) stop 8
|
||
!$omp end target
|
||
|
||
deallocate(var%e, var%f, var%str3, var%str4, var%uni3, var%uni4)
|
||
end subroutine one
|
||
|
||
! Explicitly mapped – all and full arrays
|
||
subroutine two()
|
||
type(t2) :: var
|
||
|
||
print '(g0)', '==== TESTCASE "two" ===='
|
||
|
||
var = t2(a = 1, &
|
||
b = 2, c = cmplx(-1.0_8, 2.0_8,kind=8), &
|
||
d = [(-3*i, i = 1, 10)], &
|
||
str1 = "abcde", &
|
||
str2 = ["12345", "67890", "ABCDE", "FGHIJ"], &
|
||
uni1 = 4_"abcde", &
|
||
uni2 = [4_"12345", 4_"67890", 4_"ABCDE", 4_"FGHIJ"])
|
||
allocate (var%e, source=99)
|
||
allocate (var%f, source=[22, 33, 44, 55])
|
||
allocate (var%str3, source="HelloWorld")
|
||
allocate (var%str4, source=["Let's", "Go!!!"])
|
||
allocate (var%uni3, source=4_"HelloWorld")
|
||
allocate (var%uni4, source=[4_"Let's", 4_"Go!!!"])
|
||
|
||
!$omp target map(tofrom: var%a, var%b, var%c, var%d, var%e, var%f, &
|
||
!$omp& var%str1, var%str2, var%str3, var%str4, &
|
||
!$omp& var%uni1, var%uni2, var%uni3, var%uni4)
|
||
if (var%a /= 1) stop 1
|
||
if (var%b /= 2) stop 2
|
||
if (var%c%re /= -1.0_8 .or. var%c%im /= 2.0_8) stop 3
|
||
if (any (var%d /= [(-3*i, i = 1, 10)])) stop 4
|
||
if (var%str1 /= "abcde") stop 5
|
||
if (any (var%str2 /= ["12345", "67890", "ABCDE", "FGHIJ"])) stop 6
|
||
|
||
if (.not. associated (var%e)) stop 7
|
||
if (var%e /= 99) stop 8
|
||
if (.not. associated (var%f)) stop 9
|
||
if (size (var%f) /= 4) stop 10
|
||
if (any (var%f /= [22, 33, 44, 55])) stop 11
|
||
if (.not. associated (var%str3)) stop 12
|
||
if (len (var%str3) /= len ("HelloWorld")) stop 13
|
||
if (var%str3 /= "HelloWorld") stop 14
|
||
if (.not. associated (var%str4)) stop 15
|
||
if (len (var%str4) /= 5) stop 16
|
||
if (size (var%str4) /= 2) stop 17
|
||
if (any (var%str4 /= ["Let's", "Go!!!"])) stop 18
|
||
|
||
if (var%uni1 /= 4_"abcde") stop 19
|
||
if (any (var%uni2 /= [4_"12345", 4_"67890", 4_"ABCDE", 4_"FGHIJ"])) stop 20
|
||
if (.not. associated (var%uni3)) stop 21
|
||
if (len (var%uni3) /= len (4_"HelloWorld")) stop 22
|
||
if (var%uni3 /= 4_"HelloWorld") stop 23
|
||
if (.not. associated (var%uni4)) stop 24
|
||
if (len (var%uni4) /= 5) stop 25
|
||
if (size (var%uni4) /= 2) stop 26
|
||
if (any (var%uni4 /= [4_"Let's", 4_"Go!!!"])) stop 27
|
||
!$omp end target
|
||
|
||
deallocate(var%e, var%f, var%str3, var%str4, var%uni3, var%uni4)
|
||
end subroutine two
|
||
|
||
! Explicitly mapped – one by one but full arrays
|
||
subroutine three()
|
||
type(t2) :: var
|
||
|
||
print '(g0)', '==== TESTCASE "three" ===='
|
||
|
||
var = t2(a = 1, &
|
||
b = 2, c = cmplx(-1.0_8, 2.0_8,kind=8), &
|
||
d = [(-3*i, i = 1, 10)], &
|
||
str1 = "abcde", &
|
||
str2 = ["12345", "67890", "ABCDE", "FGHIJ"], &
|
||
uni1 = 4_"abcde", &
|
||
uni2 = [4_"12345", 4_"67890", 4_"ABCDE", 4_"FGHIJ"])
|
||
allocate (var%e, source=99)
|
||
allocate (var%f, source=[22, 33, 44, 55])
|
||
allocate (var%str3, source="HelloWorld")
|
||
allocate (var%str4, source=["Let's", "Go!!!"])
|
||
allocate (var%uni3, source=4_"HelloWorld")
|
||
allocate (var%uni4, source=[4_"Let's", 4_"Go!!!"])
|
||
|
||
!$omp target map(tofrom: var%a)
|
||
if (var%a /= 1) stop 1
|
||
!$omp end target
|
||
!$omp target map(tofrom: var%b)
|
||
if (var%b /= 2) stop 2
|
||
!$omp end target
|
||
!$omp target map(tofrom: var%c)
|
||
if (var%c%re /= -1.0_8 .or. var%c%im /= 2.0_8) stop 3
|
||
!$omp end target
|
||
!$omp target map(tofrom: var%d)
|
||
if (any (var%d /= [(-3*i, i = 1, 10)])) stop 4
|
||
!$omp end target
|
||
!$omp target map(tofrom: var%str1)
|
||
if (var%str1 /= "abcde") stop 5
|
||
!$omp end target
|
||
!$omp target map(tofrom: var%str2)
|
||
if (any (var%str2 /= ["12345", "67890", "ABCDE", "FGHIJ"])) stop 6
|
||
!$omp end target
|
||
|
||
!$omp target map(tofrom: var%e)
|
||
if (.not. associated (var%e)) stop 7
|
||
if (var%e /= 99) stop 8
|
||
!$omp end target
|
||
!$omp target map(tofrom: var%f)
|
||
if (.not. associated (var%f)) stop 9
|
||
if (size (var%f) /= 4) stop 10
|
||
if (any (var%f /= [22, 33, 44, 55])) stop 11
|
||
!$omp end target
|
||
!$omp target map(tofrom: var%str3)
|
||
if (.not. associated (var%str3)) stop 12
|
||
if (len (var%str3) /= len ("HelloWorld")) stop 13
|
||
if (var%str3 /= "HelloWorld") stop 14
|
||
!$omp end target
|
||
!$omp target map(tofrom: var%str4)
|
||
if (.not. associated (var%str4)) stop 15
|
||
if (len (var%str4) /= 5) stop 16
|
||
if (size (var%str4) /= 2) stop 17
|
||
if (any (var%str4 /= ["Let's", "Go!!!"])) stop 18
|
||
!$omp end target
|
||
|
||
!$omp target map(tofrom: var%uni1)
|
||
if (var%uni1 /= 4_"abcde") stop 19
|
||
!$omp end target
|
||
!$omp target map(tofrom: var%uni2)
|
||
if (any (var%uni2 /= [4_"12345", 4_"67890", 4_"ABCDE", 4_"FGHIJ"])) stop 20
|
||
!$omp end target
|
||
!$omp target map(tofrom: var%uni3)
|
||
if (.not. associated (var%uni3)) stop 21
|
||
if (len (var%uni3) /= len (4_"HelloWorld")) stop 22
|
||
if (var%uni3 /= 4_"HelloWorld") stop 23
|
||
!$omp end target
|
||
!$omp target map(tofrom: var%uni4)
|
||
if (.not. associated (var%uni4)) stop 24
|
||
if (len (var%uni4) /= 5) stop 25
|
||
if (size (var%uni4) /= 2) stop 26
|
||
if (any (var%uni4 /= [4_"Let's", 4_"Go!!!"])) stop 27
|
||
!$omp end target
|
||
|
||
deallocate(var%e, var%f, var%str3, var%str4, var%uni3, var%uni4)
|
||
end subroutine three
|
||
|
||
! Explicitly mapped – all but only subarrays
|
||
subroutine four()
|
||
type(t2) :: var
|
||
|
||
print '(g0)', '==== TESTCASE "four" ===='
|
||
|
||
var = t2(a = 1, &
|
||
b = 2, c = cmplx(-1.0_8, 2.0_8,kind=8), &
|
||
d = [(-3*i, i = 1, 10)], &
|
||
str1 = "abcde", &
|
||
str2 = ["12345", "67890", "ABCDE", "FGHIJ"], &
|
||
uni1 = 4_"abcde", &
|
||
uni2 = [4_"12345", 4_"67890", 4_"ABCDE", 4_"FGHIJ"])
|
||
allocate (var%f, source=[22, 33, 44, 55])
|
||
allocate (var%str4, source=["Let's", "Go!!!"])
|
||
allocate (var%uni4, source=[4_"Let's", 4_"Go!!!"])
|
||
|
||
! !$omp target map(tofrom: var%d(4:7), var%f(2:3), var%str2(2:3)) &
|
||
! !$omp& map(tofrom: var%str4(2:2), var%uni2(2:3), var%uni4(2:2))
|
||
!$omp target map(tofrom: var%d(4:7), var%f(2:3), var%str2(2:3), var%uni2(2:3))
|
||
if (any (var%d(4:7) /= [(-3*i, i = 4, 7)])) stop 4
|
||
if (any (var%str2(2:3) /= ["67890", "ABCDE"])) stop 6
|
||
|
||
if (.not. associated (var%f)) stop 9
|
||
if (size (var%f) /= 4) stop 10
|
||
if (any (var%f(2:3) /= [33, 44])) stop 11
|
||
! if (.not. associated (var%str4)) stop 15
|
||
! if (len (var%str4) /= 5) stop 16
|
||
! if (size (var%str4) /= 2) stop 17
|
||
! if (var%str4(2) /= "Go!!!") stop 18
|
||
|
||
if (any (var%uni2(2:3) /= [4_"67890", 4_"ABCDE"])) stop 19
|
||
! if (.not. associated (var%uni4)) stop 20
|
||
! if (len (var%uni4) /= 5) stop 21
|
||
! if (size (var%uni4) /= 2) stop 22
|
||
! if (var%uni4(2) /= "Go!!!") stop 23
|
||
!$omp end target
|
||
|
||
deallocate(var%f, var%str4)
|
||
end subroutine four
|
||
|
||
! Explicitly mapped – all but only subarrays and one by one
|
||
subroutine five()
|
||
type(t2) :: var
|
||
|
||
print '(g0)', '==== TESTCASE "five" ===='
|
||
|
||
var = t2(a = 1, &
|
||
b = 2, c = cmplx(-1.0_8, 2.0_8,kind=8), &
|
||
d = [(-3*i, i = 1, 10)], &
|
||
str1 = "abcde", &
|
||
str2 = ["12345", "67890", "ABCDE", "FGHIJ"], &
|
||
uni1 = 4_"abcde", &
|
||
uni2 = [4_"12345", 4_"67890", 4_"ABCDE", 4_"FGHIJ"])
|
||
allocate (var%f, source=[22, 33, 44, 55])
|
||
allocate (var%str4, source=["Let's", "Go!!!"])
|
||
|
||
!$omp target map(tofrom: var%d(4:7))
|
||
if (any (var%d(4:7) /= [(-3*i, i = 4, 7)])) stop 4
|
||
!$omp end target
|
||
!$omp target map(tofrom: var%str2(2:3))
|
||
if (any (var%str2(2:3) /= ["67890", "ABCDE"])) stop 6
|
||
!$omp end target
|
||
|
||
!$omp target map(tofrom: var%f(2:3))
|
||
if (.not. associated (var%f)) stop 9
|
||
if (size (var%f) /= 4) stop 10
|
||
if (any (var%f(2:3) /= [33, 44])) stop 11
|
||
!$omp end target
|
||
! !$omp target map(tofrom: var%str4(2:2))
|
||
! if (.not. associated (var%str4)) stop 15
|
||
! if (len (var%str4) /= 5) stop 16
|
||
! if (size (var%str4) /= 2) stop 17
|
||
! if (var%str4(2) /= "Go!!!") stop 18
|
||
! !$omp end target
|
||
! !$omp target map(tofrom: var%uni4(2:2))
|
||
! if (.not. associated (var%uni4)) stop 15
|
||
! if (len (var%uni4) /= 5) stop 16
|
||
! if (size (var%uni4) /= 2) stop 17
|
||
! if (var%uni4(2) /= 4_"Go!!!") stop 18
|
||
! !$omp end target
|
||
|
||
deallocate(var%f, var%str4)
|
||
end subroutine five
|
||
|
||
! Explicitly mapped – all but only array elements
|
||
subroutine six()
|
||
type(t2) :: var
|
||
|
||
print '(g0)', '==== TESTCASE "six" ===='
|
||
|
||
var = t2(a = 1, &
|
||
b = 2, c = cmplx(-1.0_8, 2.0_8,kind=8), &
|
||
d = [(-3*i, i = 1, 10)], &
|
||
str1 = "abcde", &
|
||
str2 = ["12345", "67890", "ABCDE", "FGHIJ"], &
|
||
uni1 = 4_"abcde", &
|
||
uni2 = [4_"12345", 4_"67890", 4_"ABCDE", 4_"FGHIJ"])
|
||
allocate (var%f, source=[22, 33, 44, 55])
|
||
allocate (var%str4, source=["Let's", "Go!!!"])
|
||
allocate (var%uni4, source=[4_"Let's", 4_"Go!!!"])
|
||
|
||
! !$omp target map(tofrom: var%d(5), var%f(3), var%str2(3), &
|
||
! !$omp var%str4(2), var%uni2(3), var%uni4(2))
|
||
!$omp target map(tofrom: var%d(5), var%f(3), var%str2(3), var%uni2(3))
|
||
if (var%d(5) /= -3*5) stop 4
|
||
if (var%str2(3) /= "ABCDE") stop 6
|
||
if (var%uni2(3) /= 4_"ABCDE") stop 7
|
||
|
||
if (.not. associated (var%f)) stop 9
|
||
if (size (var%f) /= 4) stop 10
|
||
if (var%f(3) /= 44) stop 11
|
||
! if (.not. associated (var%str4)) stop 15
|
||
! if (len (var%str4) /= 5) stop 16
|
||
! if (size (var%str4) /= 2) stop 17
|
||
! if (var%str4(2) /= "Go!!!") stop 18
|
||
! if (.not. associated (var%uni4)) stop 19
|
||
! if (len (var%uni4) /= 5) stop 20
|
||
! if (size (var%uni4) /= 2) stop 21
|
||
! if (var%uni4(2) /= 4_"Go!!!") stop 22
|
||
!$omp end target
|
||
|
||
deallocate(var%f, var%str4, var%uni4)
|
||
end subroutine six
|
||
|
||
! Explicitly mapped – all but only array elements and one by one
|
||
subroutine seven()
|
||
type(t2) :: var
|
||
|
||
print '(g0)', '==== TESTCASE "seven" ===='
|
||
|
||
var = t2(a = 1, &
|
||
b = 2, c = cmplx(-1.0_8, 2.0_8,kind=8), &
|
||
d = [(-3*i, i = 1, 10)], &
|
||
str1 = "abcde", &
|
||
str2 = ["12345", "67890", "ABCDE", "FGHIJ"], &
|
||
uni1 = 4_"abcde", &
|
||
uni2 = [4_"12345", 4_"67890", 4_"ABCDE", 4_"FGHIJ"])
|
||
allocate (var%f, source=[22, 33, 44, 55])
|
||
allocate (var%str4, source=["Let's", "Go!!!"])
|
||
allocate (var%uni4, source=[4_"Let's", 4_"Go!!!"])
|
||
|
||
!$omp target map(tofrom: var%d(5))
|
||
if (var%d(5) /= (-3*5)) stop 4
|
||
!$omp end target
|
||
!$omp target map(tofrom: var%str2(2:3))
|
||
if (any (var%str2(2:3) /= ["67890", "ABCDE"])) stop 6
|
||
!$omp end target
|
||
!$omp target map(tofrom: var%uni2(2:3))
|
||
if (any (var%uni2(2:3) /= [4_"67890", 4_"ABCDE"])) stop 7
|
||
!$omp end target
|
||
|
||
!$omp target map(tofrom: var%f(2:3))
|
||
if (.not. associated (var%f)) stop 9
|
||
if (size (var%f) /= 4) stop 10
|
||
if (any (var%f(2:3) /= [33, 44])) stop 11
|
||
!$omp end target
|
||
! !$omp target map(tofrom: var%str4(2:2))
|
||
! if (.not. associated (var%str4)) stop 15
|
||
! if (len (var%str4) /= 5) stop 16
|
||
! if (size (var%str4) /= 2) stop 17
|
||
! if (var%str4(2) /= "Go!!!") stop 18
|
||
! !$omp end target
|
||
! !$omp target map(tofrom: var%uni4(2:2))
|
||
! if (.not. associated (var%uni4)) stop 15
|
||
! if (len (var%uni4) /= 5) stop 16
|
||
! if (size (var%uni4) /= 2) stop 17
|
||
! if (var%uni4(2) /= 4_"Go!!!") stop 18
|
||
! !$omp end target
|
||
|
||
deallocate(var%f, var%str4, var%uni4)
|
||
end subroutine seven
|
||
|
||
! Check mapping of NULL pointers
|
||
subroutine eight()
|
||
type(t2) :: var
|
||
|
||
print '(g0)', '==== TESTCASE "eight" ===='
|
||
|
||
var = t2(a = 1, &
|
||
b = 2, c = cmplx(-1.0_8, 2.0_8,kind=8), &
|
||
d = [(-3*i, i = 1, 10)], &
|
||
str1 = "abcde", &
|
||
str2 = ["12345", "67890", "ABCDE", "FGHIJ"], &
|
||
uni1 = 4_"abcde", &
|
||
uni2 = [4_"12345", 4_"67890", 4_"ABCDE", 4_"FGHIJ"])
|
||
|
||
! !$omp target map(tofrom: var%e, var%f, var%str3, var%str4, var%uni3, var%uni4)
|
||
!$omp target map(tofrom: var%e, var%str3, var%uni3)
|
||
if (associated (var%e)) stop 1
|
||
! if (associated (var%f)) stop 2
|
||
if (associated (var%str3)) stop 3
|
||
! if (associated (var%str4)) stop 4
|
||
if (associated (var%uni3)) stop 5
|
||
! if (associated (var%uni4)) stop 6
|
||
!$omp end target
|
||
end subroutine eight
|
||
|
||
! This is "subroutine four" but with explicit base-pointer mappings
|
||
! (var%f, etc.).
|
||
subroutine nine()
|
||
type(t2) :: var
|
||
|
||
print '(g0)', '==== TESTCASE "nine" ===='
|
||
|
||
var = t2(a = 1, &
|
||
b = 2, c = cmplx(-1.0_8, 2.0_8,kind=8), &
|
||
d = [(-3*i, i = 1, 10)], &
|
||
str1 = "abcde", &
|
||
str2 = ["12345", "67890", "ABCDE", "FGHIJ"], &
|
||
uni1 = 4_"abcde", &
|
||
uni2 = [4_"12345", 4_"67890", 4_"ABCDE", 4_"FGHIJ"])
|
||
allocate (var%f, source=[22, 33, 44, 55])
|
||
allocate (var%str4, source=["Let's", "Go!!!"])
|
||
allocate (var%uni4, source=[4_"Let's", 4_"Go!!!"])
|
||
|
||
! !$omp target map(tofrom: var%d(4:7), var%f(2:3), var%str2(2:3)) &
|
||
! !$omp& map(tofrom: var%str4(2:2), var%uni2(2:3), var%uni4(2:2))
|
||
!$omp target map(to: var%f) map(tofrom: var%d(4:7), var%f(2:3), &
|
||
!$omp& var%str2(2:3), var%uni2(2:3))
|
||
if (any (var%d(4:7) /= [(-3*i, i = 4, 7)])) stop 4
|
||
if (any (var%str2(2:3) /= ["67890", "ABCDE"])) stop 6
|
||
|
||
if (.not. associated (var%f)) stop 9
|
||
if (size (var%f) /= 4) stop 10
|
||
if (any (var%f(2:3) /= [33, 44])) stop 11
|
||
! if (.not. associated (var%str4)) stop 15
|
||
! if (len (var%str4) /= 5) stop 16
|
||
! if (size (var%str4) /= 2) stop 17
|
||
! if (var%str4(2) /= "Go!!!") stop 18
|
||
|
||
if (any (var%uni2(2:3) /= [4_"67890", 4_"ABCDE"])) stop 19
|
||
! if (.not. associated (var%uni4)) stop 20
|
||
! if (len (var%uni4) /= 5) stop 21
|
||
! if (size (var%uni4) /= 2) stop 22
|
||
! if (var%uni4(2) /= "Go!!!") stop 23
|
||
!$omp end target
|
||
|
||
deallocate(var%f, var%str4)
|
||
end subroutine nine
|
||
|
||
! This is "subroutine five" but with explicit base-pointer mappings.
|
||
subroutine ten()
|
||
type(t2) :: var
|
||
|
||
print '(g0)', '==== TESTCASE "ten" ===='
|
||
|
||
var = t2(a = 1, &
|
||
b = 2, c = cmplx(-1.0_8, 2.0_8,kind=8), &
|
||
d = [(-3*i, i = 1, 10)], &
|
||
str1 = "abcde", &
|
||
str2 = ["12345", "67890", "ABCDE", "FGHIJ"], &
|
||
uni1 = 4_"abcde", &
|
||
uni2 = [4_"12345", 4_"67890", 4_"ABCDE", 4_"FGHIJ"])
|
||
allocate (var%f, source=[22, 33, 44, 55])
|
||
allocate (var%str4, source=["Let's", "Go!!!"])
|
||
|
||
!$omp target map(tofrom: var%d(4:7))
|
||
if (any (var%d(4:7) /= [(-3*i, i = 4, 7)])) stop 4
|
||
!$omp end target
|
||
!$omp target map(tofrom: var%str2(2:3))
|
||
if (any (var%str2(2:3) /= ["67890", "ABCDE"])) stop 6
|
||
!$omp end target
|
||
|
||
!$omp target map(to: var%f) map(tofrom: var%f(2:3))
|
||
if (.not. associated (var%f)) stop 9
|
||
if (size (var%f) /= 4) stop 10
|
||
if (any (var%f(2:3) /= [33, 44])) stop 11
|
||
!$omp end target
|
||
! !$omp target map(tofrom: var%str4(2:2))
|
||
! if (.not. associated (var%str4)) stop 15
|
||
! if (len (var%str4) /= 5) stop 16
|
||
! if (size (var%str4) /= 2) stop 17
|
||
! if (var%str4(2) /= "Go!!!") stop 18
|
||
! !$omp end target
|
||
! !$omp target map(tofrom: var%uni4(2:2))
|
||
! if (.not. associated (var%uni4)) stop 15
|
||
! if (len (var%uni4) /= 5) stop 16
|
||
! if (size (var%uni4) /= 2) stop 17
|
||
! if (var%uni4(2) /= 4_"Go!!!") stop 18
|
||
! !$omp end target
|
||
|
||
deallocate(var%f, var%str4)
|
||
end subroutine ten
|
||
|
||
! This is "subroutine six" but with explicit base pointer mappings.
|
||
subroutine eleven()
|
||
type(t2) :: var
|
||
|
||
print '(g0)', '==== TESTCASE "eleven" ===='
|
||
|
||
var = t2(a = 1, &
|
||
b = 2, c = cmplx(-1.0_8, 2.0_8,kind=8), &
|
||
d = [(-3*i, i = 1, 10)], &
|
||
str1 = "abcde", &
|
||
str2 = ["12345", "67890", "ABCDE", "FGHIJ"], &
|
||
uni1 = 4_"abcde", &
|
||
uni2 = [4_"12345", 4_"67890", 4_"ABCDE", 4_"FGHIJ"])
|
||
allocate (var%f, source=[22, 33, 44, 55])
|
||
allocate (var%str4, source=["Let's", "Go!!!"])
|
||
allocate (var%uni4, source=[4_"Let's", 4_"Go!!!"])
|
||
|
||
! !$omp target map(tofrom: var%d(5), var%f(3), var%str2(3), &
|
||
! !$omp var%str4(2), var%uni2(3), var%uni4(2))
|
||
!$omp target map(to: var%f) map(tofrom: var%d(5), var%f(3), &
|
||
!$omp& var%str2(3), var%uni2(3))
|
||
if (var%d(5) /= -3*5) stop 4
|
||
if (var%str2(3) /= "ABCDE") stop 6
|
||
if (var%uni2(3) /= 4_"ABCDE") stop 7
|
||
|
||
if (.not. associated (var%f)) stop 9
|
||
if (size (var%f) /= 4) stop 10
|
||
if (var%f(3) /= 44) stop 11
|
||
! if (.not. associated (var%str4)) stop 15
|
||
! if (len (var%str4) /= 5) stop 16
|
||
! if (size (var%str4) /= 2) stop 17
|
||
! if (var%str4(2) /= "Go!!!") stop 18
|
||
! if (.not. associated (var%uni4)) stop 19
|
||
! if (len (var%uni4) /= 5) stop 20
|
||
! if (size (var%uni4) /= 2) stop 21
|
||
! if (var%uni4(2) /= 4_"Go!!!") stop 22
|
||
!$omp end target
|
||
|
||
deallocate(var%f, var%str4, var%uni4)
|
||
end subroutine eleven
|
||
|
||
! This is "subroutine seven" but with explicit base-pointer mappings.
|
||
subroutine twelve()
|
||
type(t2) :: var
|
||
|
||
print '(g0)', '==== TESTCASE "twelve" ===='
|
||
|
||
var = t2(a = 1, &
|
||
b = 2, c = cmplx(-1.0_8, 2.0_8,kind=8), &
|
||
d = [(-3*i, i = 1, 10)], &
|
||
str1 = "abcde", &
|
||
str2 = ["12345", "67890", "ABCDE", "FGHIJ"], &
|
||
uni1 = 4_"abcde", &
|
||
uni2 = [4_"12345", 4_"67890", 4_"ABCDE", 4_"FGHIJ"])
|
||
allocate (var%f, source=[22, 33, 44, 55])
|
||
allocate (var%str4, source=["Let's", "Go!!!"])
|
||
allocate (var%uni4, source=[4_"Let's", 4_"Go!!!"])
|
||
|
||
!$omp target map(tofrom: var%d(5))
|
||
if (var%d(5) /= (-3*5)) stop 4
|
||
!$omp end target
|
||
!$omp target map(tofrom: var%str2(2:3))
|
||
if (any (var%str2(2:3) /= ["67890", "ABCDE"])) stop 6
|
||
!$omp end target
|
||
!$omp target map(tofrom: var%uni2(2:3))
|
||
if (any (var%uni2(2:3) /= [4_"67890", 4_"ABCDE"])) stop 7
|
||
!$omp end target
|
||
|
||
!$omp target map(to: var%f) map(tofrom: var%f(2:3))
|
||
if (.not. associated (var%f)) stop 9
|
||
if (size (var%f) /= 4) stop 10
|
||
if (any (var%f(2:3) /= [33, 44])) stop 11
|
||
!$omp end target
|
||
! !$omp target map(tofrom: var%str4(2:2))
|
||
! if (.not. associated (var%str4)) stop 15
|
||
! if (len (var%str4) /= 5) stop 16
|
||
! if (size (var%str4) /= 2) stop 17
|
||
! if (var%str4(2) /= "Go!!!") stop 18
|
||
! !$omp end target
|
||
! !$omp target map(tofrom: var%uni4(2:2))
|
||
! if (.not. associated (var%uni4)) stop 15
|
||
! if (len (var%uni4) /= 5) stop 16
|
||
! if (size (var%uni4) /= 2) stop 17
|
||
! if (var%uni4(2) /= 4_"Go!!!") stop 18
|
||
! !$omp end target
|
||
|
||
deallocate(var%f, var%str4, var%uni4)
|
||
end subroutine twelve
|
||
|
||
end program main
|