mirror of
git://gcc.gnu.org/git/gcc.git
synced 2025-04-18 10:00:35 +08:00
openacc: Don't strip TO_PSET/POINTER for enter/exit data
OpenACC 2.6 specifies that the array descriptor (when present) must be copied to the target before attaching pointers in Fortran. This patch reverses the stripping of GOMP_MAP_TO_PSET and GOMP_MAP_POINTER that was introduced by the "OpenACC reference count overhaul" patch. 2020-07-10 Julian Brown <julian@codesourcery.com> Thomas Schwinge <thomas@codesourcery.com> gcc/ * gimplify.c (gimplify_scan_omp_clauses): Do not strip GOMP_MAP_TO_PSET/GOMP_MAP_POINTER for OpenACC enter/exit data directives (see also PR92929). gcc/testsuite/ * gfortran.dg/goacc/finalize-1.f: Update expected dump output. libgomp/ * testsuite/libgomp.oacc-fortran/dynamic-pointer-1.f90: New test. Co-Authored-By: Thomas Schwinge <thomas@codesourcery.com>
This commit is contained in:
Julian Brown
parent
7a4770f039
commit
b20097c65d
@ -8768,6 +8768,8 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
|
||||
case OMP_TARGET_DATA:
|
||||
case OMP_TARGET_ENTER_DATA:
|
||||
case OMP_TARGET_EXIT_DATA:
|
||||
case OACC_ENTER_DATA:
|
||||
case OACC_EXIT_DATA:
|
||||
case OACC_HOST_DATA:
|
||||
if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER
|
||||
|| (OMP_CLAUSE_MAP_KIND (c)
|
||||
@ -8776,15 +8778,6 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
|
||||
mapped, but not the pointer to it. */
|
||||
remove = true;
|
||||
break;
|
||||
case OACC_ENTER_DATA:
|
||||
case OACC_EXIT_DATA:
|
||||
if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER
|
||||
|| OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_TO_PSET
|
||||
|| OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER
|
||||
|| (OMP_CLAUSE_MAP_KIND (c)
|
||||
== GOMP_MAP_FIRSTPRIVATE_REFERENCE))
|
||||
remove = true;
|
||||
break;
|
||||
default:
|
||||
break;
|
||||
}
|
||||
@ -8794,7 +8787,15 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
|
||||
does not make sense. Likewise, for 'update' only transferring the
|
||||
data itself is needed as the rest has been handled in previous
|
||||
directives. However, for 'exit data', the array descriptor needs
|
||||
to be delete; hence, we turn the MAP_TO_PSET into a MAP_DELETE. */
|
||||
to be delete; hence, we turn the MAP_TO_PSET into a MAP_DELETE.
|
||||
|
||||
NOTE: Generally, it is not safe to perform "enter data" operations
|
||||
on arrays where the data *or the descriptor* may go out of scope
|
||||
before a corresponding "exit data" operation -- and such a
|
||||
descriptor may be synthesized temporarily, e.g. to pass an
|
||||
explicit-shape array to a function expecting an assumed-shape
|
||||
argument. Performing "enter data" inside the called function
|
||||
would thus be problematic. */
|
||||
if (code == OMP_TARGET_EXIT_DATA
|
||||
&& OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_TO_PSET)
|
||||
OMP_CLAUSE_SET_MAP_KIND (c, OMP_CLAUSE_MAP_KIND (*prev_list_p)
|
||||
|
||||
@ -21,7 +21,7 @@
|
||||
|
||||
!$ACC EXIT DATA FINALIZE DELETE (del_f_p(2:5))
|
||||
! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(release:\\*\\(c_char \\*\\) parm\\.0\\.data \\\[len: \[^\\\]\]+\\\]\\) map\\(to:del_f_p \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(alloc:\\(integer\\(kind=1\\)\\\[0:\\\] \\* restrict\\) del_f_p\\.data \\\[pointer assign, bias: \\(sizetype\\) parm\\.0\\.data - \\(sizetype\\) del_f_p\\.data\\\]\\) finalize;$" 1 "original" } }
|
||||
! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(delete:MEM\\\[\\(c_char \\*\\)\[^\\\]\]+\\\] \\\[len: \[^\\\]\]+\\\]\\) finalize$" 1 "gimple" } }
|
||||
! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(delete:MEM\\\[\\(c_char \\*\\)\[^\\\]\]+\\\] \\\[len: \[^\\\]\]+\\\]\\) map\\(to:del_f_p \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(alloc:del_f_p\\.data \\\[pointer assign, bias: \[^\\\]\]+\\\]\\) finalize$" 1 "gimple" } }
|
||||
|
||||
!$ACC EXIT DATA COPYOUT (cpo_r)
|
||||
! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(from:cpo_r\\);$" 1 "original" } }
|
||||
@ -33,5 +33,5 @@
|
||||
|
||||
!$ACC EXIT DATA COPYOUT (cpo_f_p(4:10)) FINALIZE
|
||||
! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(from:\\*\\(c_char \\*\\) parm\\.1\\.data \\\[len: \[^\\\]\]+\\\]\\) map\\(to:cpo_f_p \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(alloc:\\(integer\\(kind=1\\)\\\[0:\\\] \\* restrict\\) cpo_f_p\\.data \\\[pointer assign, bias: \\(sizetype\\) parm\\.1\\.data - \\(sizetype\\) cpo_f_p\\.data\\\]\\) finalize;$" 1 "original" } }
|
||||
! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(force_from:MEM\\\[\\(c_char \\*\\)\[^\\\]\]+\\\] \\\[len: \[^\\\]\]+\\\]\\) finalize$" 1 "gimple" } }
|
||||
! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(force_from:MEM\\\[\\(c_char \\*\\)\[^\\\]\]+\\\] \\\[len: \[^\\\]\]+\\\]\\) map\\(to:cpo_f_p \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(alloc:cpo_f_p\\.data \\\[pointer assign, bias: \[^\\\]\]+\\\]\\) finalize$" 1 "gimple" } }
|
||||
END SUBROUTINE f
|
||||
|
||||
97
libgomp/testsuite/libgomp.oacc-fortran/dynamic-pointer-1.f90
Normal file
97
libgomp/testsuite/libgomp.oacc-fortran/dynamic-pointer-1.f90
Normal file
@ -0,0 +1,97 @@
|
||||
! Verify that a 'enter data'ed 'pointer' object creates a persistent, visible device copy
|
||||
|
||||
! { dg-do run }
|
||||
! { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } }
|
||||
|
||||
module m
|
||||
implicit none
|
||||
contains
|
||||
|
||||
subroutine verify_a (a_ref, a)
|
||||
implicit none
|
||||
integer, dimension (:, :, :), allocatable :: a_ref
|
||||
integer, dimension (:, :, :), pointer :: a
|
||||
|
||||
!$acc routine seq
|
||||
|
||||
if (any (lbound (a) /= lbound (a_ref))) stop 101
|
||||
if (any (ubound (a) /= ubound (a_ref))) stop 102
|
||||
if (size (a) /= size (a_ref)) stop 103
|
||||
end subroutine verify_a
|
||||
|
||||
end module m
|
||||
|
||||
program main
|
||||
use m
|
||||
use openacc
|
||||
implicit none
|
||||
integer, parameter :: n = 30
|
||||
integer, dimension (:, :, :), allocatable, target :: a1, a2
|
||||
integer, dimension (:, :, :), pointer :: p
|
||||
|
||||
allocate (a1(1:n, 0:n-1, 10:n/2))
|
||||
!$acc enter data create(a1)
|
||||
allocate (a2(3:n/3, 10:n, n-10:n+10))
|
||||
!$acc enter data create(a2)
|
||||
|
||||
p => a1
|
||||
call verify_a(a1, p)
|
||||
|
||||
! 'p' object isn't present on the device.
|
||||
!$acc parallel ! Implicit 'copy(p)'; creates 'p' object...
|
||||
call verify_a(a1, p)
|
||||
!$acc end parallel ! ..., and deletes it again.
|
||||
|
||||
p => a2
|
||||
call verify_a(a2, p)
|
||||
|
||||
! 'p' object isn't present on the device.
|
||||
!$acc parallel ! Implicit 'copy(p)'; creates 'p' object...
|
||||
call verify_a(a2, p)
|
||||
!$acc end parallel ! ..., and deletes it again.
|
||||
|
||||
p => a1
|
||||
|
||||
!$acc enter data create(p)
|
||||
! 'p' object is now present on the device (visible device copy).
|
||||
!TODO PR96080 if (.not. acc_is_present (p)) stop 1
|
||||
|
||||
!$acc parallel
|
||||
! On the device, got created as 'p => a1'.
|
||||
call verify_a(a1, p)
|
||||
!$acc end parallel
|
||||
call verify_a(a1, p)
|
||||
|
||||
!$acc parallel
|
||||
p => a2
|
||||
! On the device, 'p => a2' is now set.
|
||||
call verify_a(a2, p)
|
||||
!$acc end parallel
|
||||
! On the host, 'p => a1' persists.
|
||||
call verify_a(a1, p)
|
||||
|
||||
!$acc parallel
|
||||
! On the device, 'p => a2' persists.
|
||||
call verify_a(a2, p)
|
||||
!$acc end parallel
|
||||
! On the host, 'p => a1' still persists.
|
||||
call verify_a(a1, p)
|
||||
|
||||
p => a2
|
||||
|
||||
!$acc parallel
|
||||
p => a1
|
||||
! On the device, 'p => a1' is now set.
|
||||
call verify_a(a1, p)
|
||||
!$acc end parallel
|
||||
! On the host, 'p => a2' persists.
|
||||
call verify_a(a2, p)
|
||||
|
||||
!$acc parallel
|
||||
! On the device, 'p => a1' persists.
|
||||
call verify_a(a1, p)
|
||||
!$acc end parallel
|
||||
! On the host, 'p => a2' still persists.
|
||||
call verify_a(a2, p)
|
||||
|
||||
end program main
|
||||
Loading…
x
Reference in New Issue
Block a user