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  <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>

diff --git a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c
index f35ff4191c27e60e0bb5558647614696ce2c8492..80929c77cc6543b57c8c6c88889483c4dcdc2f06 100644 (file)
--- 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 (file)
index 0000000..9ca36f7
--- /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 (file)
index 0000000..5d79cbc
--- /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 (file)
index 0000000..58ef44e
--- /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