From 39dda0020801045d9a604575b2a2593c05310015 Mon Sep 17 00:00:00 2001 From: Julian Brown Date: Tue, 9 Jun 2020 15:53:39 -0700 Subject: [PATCH] openacc: Fix standalone attach for Fortran assumed-shape array pointers 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 Thomas Schwinge 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 --- gcc/fortran/trans-openmp.c | 22 ++++++-- .../gfortran.dg/goacc/attach-descriptor.f90 | 18 +++++++ .../attach-descriptor-1.f90 | 53 +++++++++++++++++++ .../attach-descriptor-2.f90 | 18 +++++++ 4 files changed, 107 insertions(+), 4 deletions(-) create mode 100644 gcc/testsuite/gfortran.dg/goacc/attach-descriptor.f90 create mode 100644 libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-1.f90 create mode 100644 libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-2.f90 diff --git a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c index f35ff4191c2..80929c77cc6 100644 --- a/gcc/fortran/trans-openmp.c +++ b/gcc/fortran/trans-openmp.c @@ -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); diff --git a/gcc/testsuite/gfortran.dg/goacc/attach-descriptor.f90 b/gcc/testsuite/gfortran.dg/goacc/attach-descriptor.f90 new file mode 100644 index 00000000000..9ca36f770c7 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/goacc/attach-descriptor.f90 @@ -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 diff --git a/libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-1.f90 new file mode 100644 index 00000000000..5d79cbc14fc --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-1.f90 @@ -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 diff --git a/libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-2.f90 b/libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-2.f90 new file mode 100644 index 00000000000..58ef44e7c76 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-2.f90 @@ -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 -- 2.30.2