openacc: No attach/detach present/release mappings for array descriptors

Standalone attach and detach clauses should not create present/release
mappings for Fortran array descriptors (e.g. used when we have a pointer
to an array), both because it is unnecessary and because those mappings
will be incorrectly subject to reference counting. Simply omitting the
mappings means we just use GOMP_MAP_TO_PSET and GOMP_MAP_{ATTACH,DETACH}
mappings for array descriptors.

That requires a tweak in gimplify.c, since we may now see GOMP_MAP_TO_PSET
without a preceding data-movement mapping.

2020-08-03  Julian Brown  <julian@codesourcery.com>
	    Thomas Schwinge  <thomas@codesourcery.com>

gcc/fortran/
	* trans-openmp.c (gfc_trans_omp_clauses): Don't create present/release
	mappings for array descriptors.

gcc/
	* gimplify.c (gimplify_omp_target_update): Allow GOMP_MAP_TO_PSET
	without a preceding data-movement mapping.

gcc/testsuite/
	* gfortran.dg/goacc/attach-descriptor.f90: Update pattern output. Add
	scanning of gimplify dump.

libgomp/
	* testsuite/libgomp.oacc-fortran/attach-descriptor-1.f90: Don't run for
	shared-memory devices.  Extend with further checking.

Co-Authored-By: Thomas Schwinge <thomas@codesourcery.com>

3 files changed, 47 insertions, 21 deletions

diff --git a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c
index 076efb0..98702b1 100644
--- a/gcc/fortran/trans-openmp.c
+++ b/gcc/fortran/trans-openmp.c
@@ -2718,23 +2718,6 @@ 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);
-		      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);
@@ -2748,6 +2731,33 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
 			OMP_CLAUSE_DECL (node3)
 			  = gfc_conv_descriptor_data_get (decl);
 		      OMP_CLAUSE_SIZE (node3) = size_int (0);
+		      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 (node3, GOMP_MAP_ATTACH);
+			  /* We don't want to map PTR at all in this case, so
+			     delete its node and shuffle the others down.  */
+			  node = node2;
+			  node2 = node3;
+			  node3 = NULL;
+			  goto finalize_map_clause;
+			}
+		      else if (n->u.map_op == OMP_MAP_DETACH)
+			{
+			  OMP_CLAUSE_SET_MAP_KIND (node3, GOMP_MAP_DETACH);
+			  /* Similarly to above, we don't want to unmap PTR
+			     here.  */
+			  node = node2;
+			  node2 = node3;
+			  node3 = NULL;
+			  goto finalize_map_clause;
+			}
+		      else
+			OMP_CLAUSE_SET_MAP_KIND (node3, GOMP_MAP_POINTER);
 
 		      /* We have to check for n->sym->attr.dimension because
 			 of scalar coarrays.  */
diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index 15dfee9..6a5349c 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -13013,8 +13013,13 @@ gimplify_omp_target_update (tree *expr_p, gimple_seq *pre_p)
 	      OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_DELETE);
 	      have_clause = true;
 	      break;
-	    case GOMP_MAP_POINTER:
 	    case GOMP_MAP_TO_PSET:
+	      /* Fortran arrays with descriptors must map that descriptor when
+		 doing standalone "attach" operations (in OpenACC).  In that
+		 case GOMP_MAP_TO_PSET appears by itself with no preceding
+		 clause (see trans-openmp.c:gfc_trans_omp_clauses).  */
+	      break;
+	    case GOMP_MAP_POINTER:
 	      /* TODO PR92929: we may see these here, but they'll always follow
 		 one of the clauses above, and will be handled by libgomp as
 		 one group, so no handling required here.  */
diff --git a/gcc/testsuite/gfortran.dg/goacc/attach-descriptor.f90 b/gcc/testsuite/gfortran.dg/goacc/attach-descriptor.f90
index 9ca36f7..373bdcb 100644
--- a/gcc/testsuite/gfortran.dg/goacc/attach-descriptor.f90
+++ b/gcc/testsuite/gfortran.dg/goacc/attach-descriptor.f90
@@ -1,4 +1,4 @@
-! { dg-additional-options "-fdump-tree-original" }
+! { dg-additional-options "-fdump-tree-original -fdump-tree-gimple" }
 
 program att
   implicit none
@@ -11,8 +11,19 @@ program att
   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" } }
+! { dg-final { scan-tree-dump-times "(?n)#pragma acc enter data map\\(attach:myvar\\.arr2 \\\[bias: 0\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(attach:\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) myptr\\.data \\\[bias: 0\\\]\\);$" 1 "original" } }
+! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(attach:myvar\\.arr2 \\\[bias: 0\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(attach:myptr\\.data \\\[bias: 0\\\]\\)$" 1 "gimple" } }
 
   !$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" } }
+! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(detach:myvar\\.arr2 \\\[bias: 0\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(detach:\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) myptr\\.data \\\[bias: 0\\\]\\);$" 1 "original" } }
+! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(detach:myvar\\.arr2 \\\[bias: 0\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(detach:myptr\\.data \\\[bias: 0\\\]\\)$" 1 "gimple" } }
+
+  ! Test valid usage and processing of the finalize clause.
+  !$acc exit data detach(myvar%arr2, myptr) finalize
+! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(detach:myvar\\.arr2 \\\[bias: 0\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(detach:\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) myptr\\.data \\\[bias: 0\\\]\\) finalize;$" 1 "original" } }
+  ! For array-descriptor detaches, we no longer generate a "release" mapping
+  ! for the pointed-to data for gimplify.c to turn into "delete".  Make sure
+  ! the mapping still isn't there.
+! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(force_detach:myvar\\.arr2 \\\[bias: 0\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(force_detach:myptr\\.data \\\[bias: 0\\\]\\) finalize$" 1 "gimple" } }
+
 end program att