MatCat.OpenSource/8sa1-gcc
1
0
Fork 0
Code Issues Pull Requests Projects Releases Wiki Activity

openacc: Fix standalone attach for Fortran assumed-shape array pointers

Browse Source 
This patch makes it so that an "attach" operation for a Fortran pointer
with an array descriptor copies that array descriptor to the target,
and similarly that detach operations release the array descriptor.

2020-07-16  Julian Brown  <julian@codesourcery.com>
	    Thomas Schwinge  <thomas@codesourcery.com>

gcc/fortran/
	* trans-openmp.c (gfc_trans_omp_clauses): Rework OpenACC
	attach/detach handling for arrays with descriptors.

gcc/testsuite/
	* gfortran.dg/goacc/attach-descriptor.f90: New test.

libgomp/
	* testsuite/libgomp.oacc-fortran/attach-descriptor-1.f90: New test.
	* testsuite/libgomp.oacc-fortran/attach-descriptor-2.f90: New test.

Co-Authored-By: Thomas Schwinge <thomas@codesourcery.com>
This commit is contained in:
Julian Brown 2020-06-09 15:53:39 -07:00
parent 0f66b8486c
commit 39dda00208
4 changed files with 107 additions and 4 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

22
gcc/fortran/trans-openmp.c
View File

@ -2685,9 +2685,7 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
decl = build_fold_indirect_ref (decl);
}
}
if (GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (decl))
&& n->u.map_op != OMP_MAP_ATTACH
&& n->u.map_op != OMP_MAP_DETACH)
if (GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (decl)))
{
tree type = TREE_TYPE (decl);
tree ptr = gfc_conv_descriptor_data_get (decl);
@ -2705,7 +2703,23 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
OMP_CLAUSE_SIZE (node2) = TYPE_SIZE_UNIT (type);
node3 = build_omp_clause (input_location,
OMP_CLAUSE_MAP);
OMP_CLAUSE_SET_MAP_KIND (node3, GOMP_MAP_POINTER);
if (n->u.map_op == OMP_MAP_ATTACH)
{
/* Standalone attach clauses used with arrays with
descriptors must copy the descriptor to the target,
else they won't have anything to perform the
attachment onto (see OpenACC 2.6, "2.6.3. Data
Structures with Pointers"). */
OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_ALLOC);
OMP_CLAUSE_SET_MAP_KIND (node3, GOMP_MAP_ATTACH);
}
else if (n->u.map_op == OMP_MAP_DETACH)
{
OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_RELEASE);
OMP_CLAUSE_SET_MAP_KIND (node3, GOMP_MAP_DETACH);
}
else
OMP_CLAUSE_SET_MAP_KIND (node3, GOMP_MAP_POINTER);
if (present)
{
ptr = gfc_conv_descriptor_data_get (decl);

18
gcc/testsuite/gfortran.dg/goacc/attach-descriptor.f90 Normal file
View File

@ -0,0 +1,18 @@
! { dg-additional-options "-fdump-tree-original" }
program att
implicit none
type t
integer :: arr1(10)
integer, allocatable :: arr2(:)
end type t
type(t) :: myvar
integer, target :: tarr(10)
integer, pointer :: myptr(:)
!$acc enter data attach(myvar%arr2, myptr)
! { dg-final { scan-tree-dump-times "(?n)#pragma acc enter data map\\(attach:myvar\\.arr2 \\\[bias: 0\\\]\\) map\\(alloc:\\*\\(c_char \\*\\) myptr\\.data \\\[len: \[^\\\]\]+\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(attach:\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) myptr\\.data \\\[bias: 0\\\]\\);$" 1 "original" } }
!$acc exit data detach(myvar%arr2, myptr)
! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(detach:myvar\\.arr2 \\\[bias: 0\\\]\\) map\\(release:\\*\\(c_char \\*\\) myptr\\.data \\\[len: \[^\\\]\]+\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(detach:\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) myptr\\.data \\\[bias: 0\\\]\\);$" 1 "original" } }
end program att

53
libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-1.f90 Normal file
View File

@ -0,0 +1,53 @@
! { dg-do run }
program att
use openacc
implicit none
type t
integer :: arr1(10)
integer, allocatable :: arr2(:)
end type t
integer :: i
type(t) :: myvar
integer, target :: tarr(10)
integer, pointer :: myptr(:)
allocate(myvar%arr2(10))
do i=1,10
myvar%arr1(i) = 0
myvar%arr2(i) = 0
tarr(i) = 0
end do
call acc_copyin(myvar)
call acc_copyin(myvar%arr2)
call acc_copyin(tarr)
myptr => tarr
!$acc enter data attach(myvar%arr2, myptr)
! FIXME: This warning is emitted on the wrong line number.
! { dg-warning "using vector_length \\(32\\), ignoring 1" "" { target openacc_nvidia_accel_selected } 38 }
!$acc serial present(myvar%arr2)
do i=1,10
myvar%arr1(i) = i
myvar%arr2(i) = i
end do
myptr(3) = 99
!$acc end serial
!$acc exit data detach(myvar%arr2, myptr)
call acc_copyout(myvar%arr2)
call acc_copyout(myvar)
call acc_copyout(tarr)
do i=1,10
if (myvar%arr1(i) .ne. i) stop 1
if (myvar%arr2(i) .ne. i) stop 2
end do
if (tarr(3) .ne. 99) stop 3
end program att

18
libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-2.f90 Normal file
View File

@ -0,0 +1,18 @@
! { dg-do run }
program main
use openacc
implicit none
! TODO Per PR96080, data types chosen so that we can create a
! "pointer object 'data_p'" on the device.
integer, dimension(:), target :: data(1)
integer, dimension(:), pointer :: data_p
!TODO Per PR96080, not using OpenACC/Fortran runtime library routines.
!$acc enter data create(data)
data_p => data
!$acc enter data copyin(data_p)
!$acc enter data attach(data_p)
end program main