OpenMP: Pointers and member mappings

This patch changes the mapping node arrangement used for array components
of derived types in order to accommodate for changes made in the previous
patch, particularly the use of "GOMP_MAP_ATTACH_DETACH" for pointer-typed
derived-type members instead of "GOMP_MAP_ALWAYS_POINTER".

We change the mapping nodes used for a derived-type mapping like this:

  type T
  integer, pointer, dimension(:) :: arrptr
  end type T

  type(T) :: tvar
  [...]
  !$omp target map(tofrom: tvar%arrptr)

So that the nodes used look like this:

  1) map(to: tvar%arrptr)   -->
  GOMP_MAP_TO [implicit]  *tvar%arrptr%data  (the array data)
  GOMP_MAP_TO_PSET        tvar%arrptr        (the descriptor)
  GOMP_MAP_ATTACH_DETACH  tvar%arrptr%data

  2) map(tofrom: tvar%arrptr(3:8)   -->
  GOMP_MAP_TOFROM         *tvar%arrptr%data(3)  (size 8-3+1, etc.)
  GOMP_MAP_TO_PSET        tvar%arrptr
  GOMP_MAP_ATTACH_DETACH  tvar%arrptr%data      (bias 3, etc.)

In this case, we can determine in the front-end that the
whole-array/pointer mapping (1) is only needed to map the pointer
-- so we drop it entirely.  (Note also that we set -- early -- the
OMP_CLAUSE_MAP_RUNTIME_IMPLICIT_P flag for whole-array-via-pointer
mappings. See below.)

In the middle end, we process mappings using the struct sibling-list
handling machinery by moving the "GOMP_MAP_TO_PSET" node from the middle
of the group of three mapping nodes to the proper sorted position after
the GOMP_MAP_STRUCT mapping:

  GOMP_MAP_STRUCT   tvar     (len: 1)
  GOMP_MAP_TO_PSET  tvar%arr (size: 64, etc.)  <--. moved here
  [...]                                           |
  GOMP_MAP_TOFROM         *tvar%arrptr%data(3) ___|
  GOMP_MAP_ATTACH_DETACH  tvar%arrptr%data

In another case, if we have an array of derived-type values "dtarr",
and mappings like:

  i = 1
  j = 1
  map(to: dtarr(i)%arrptr) map(tofrom: dtarr(j)%arrptr(3:8))

We still map the same way, but this time we cannot prove that the base
expressions "dtarr(i) and "dtarr(j)" are the same in the front-end.
So we keep both mappings, but we move the "[implicit]" mapping of the
full-array reference to the end of the clause list in gimplify.cc (by
adjusting the topological sorting algorithm):

  GOMP_MAP_STRUCT         dtvar  (len: 2)
  GOMP_MAP_TO_PSET        dtvar(i)%arrptr
  GOMP_MAP_TO_PSET        dtvar(j)%arrptr
  [...]
  GOMP_MAP_TOFROM         *dtvar(j)%arrptr%data(3)  (size: 8-3+1)
  GOMP_MAP_ATTACH_DETACH  dtvar(j)%arrptr%data
  GOMP_MAP_TO [implicit]  *dtvar(i)%arrptr%data(1)  (size: whole array)
  GOMP_MAP_ATTACH_DETACH  dtvar(i)%arrptr%data

Always moving "[implicit]" full-array mappings after array-section
mappings (without that bit set) means that we'll avoid copying the whole
array unnecessarily -- even in cases where we can't prove that the arrays
are the same.

The patch also fixes some bugs with "enter data" and "exit data"
directives with this new mapping arrangement.  Also now if you have
mappings like this:

  #pragma omp target enter data map(to: dv, dv%arr(1:20))

The whole of the derived-type variable "dv" is mapped, so the
GOMP_MAP_TO_PSET for the array-section mapping can be dropped:

  GOMP_MAP_TO            dv

  GOMP_MAP_TO            *dv%arr%data
  GOMP_MAP_TO_PSET       dv%arr <-- deleted (array section mapping)
  GOMP_MAP_ATTACH_DETACH dv%arr%data

To accommodate for recent changes to mapping nodes made by
Tobias, this version of the patch avoids using GOMP_MAP_TO_PSET
for "exit data" directives, in favour of using the "correct"
GOMP_MAP_RELEASE/GOMP_MAP_DELETE kinds during early expansion.  A new
flag is introduced so the middle-end knows when the latter two kinds
are being used specifically for an array descriptor.

This version of the patch fixes "omp target exit data" handling
for GOMP_MAP_DELETE, and adds pretty-printing dump output
for the OMP_CLAUSE_RELEASE_DESCRIPTOR flag (for a little extra
clarity).

Also I noticed the handling of descriptors on *OpenACC*
exit-data directives was inconsistent, so I've made those use
GOMP_MAP_RELEASE/GOMP_MAP_DELETE with the new flag in the same way as
OpenMP too.  In the end it doesn't actually matter to the runtime,
which handles GOMP_MAP_RELEASE/GOMP_MAP_DELETE/GOMP_MAP_TO_PSET for
array descriptors on OpenACC "exit data" directives the same, anyway,
and doing it this way in the FE avoids needless divergence.

I've added a couple of new tests (gomp/target-enter-exit-data.f90 and
goacc/enter-exit-data-2.f90).

2023-12-07  Julian Brown  <julian@codesourcery.com>

gcc/fortran/
	* dependency.cc (gfc_omp_expr_prefix_same): New function.
	* dependency.h (gfc_omp_expr_prefix_same): Add prototype.
	* gfortran.h (gfc_omp_namelist): Add "duplicate_of" field to "u2"
	union.
	* trans-openmp.cc (dependency.h): Include.
	(gfc_trans_omp_array_section): Adjust mapping node arrangement for
	array descriptors.  Use GOMP_MAP_TO_PSET or
	GOMP_MAP_RELEASE/GOMP_MAP_DELETE with the OMP_CLAUSE_RELEASE_DESCRIPTOR
	flag set.
	(gfc_symbol_rooted_namelist): New function.
	(gfc_trans_omp_clauses): Check subcomponent and subarray/element
	accesses elsewhere in the clause list for pointers to derived types or
	array descriptors, and adjust or drop mapping nodes appropriately.
	Adjust for changes to mapping node arrangement.
	(gfc_trans_oacc_executable_directive): Pass code op through.

gcc/
	* gimplify.cc (omp_map_clause_descriptor_p): New function.
	(build_omp_struct_comp_nodes, omp_get_attachment, omp_group_base): Use
	above function.
	(omp_tsort_mapping_groups): Process nodes that have
	OMP_CLAUSE_MAP_RUNTIME_IMPLICIT_P set after those that don't.  Add
	enter_exit_data parameter.
	(omp_resolve_clause_dependencies): Remove GOMP_MAP_TO_PSET mappings if
	we're mapping the whole containing derived-type variable.
	(omp_accumulate_sibling_list): Adjust GOMP_MAP_TO_PSET handling.
	Remove GOMP_MAP_ALWAYS_POINTER handling.
	(gimplify_scan_omp_clauses): Pass enter_exit argument to
	omp_tsort_mapping_groups.  Don't adjust/remove GOMP_MAP_TO_PSET
	mappings for derived-type components here.
	* tree.h (OMP_CLAUSE_RELEASE_DESCRIPTOR): New macro.
	* tree-pretty-print.cc (dump_omp_clause): Show
	OMP_CLAUSE_RELEASE_DESCRIPTOR in dump output (with
	GOMP_MAP_TO_PSET-like syntax).

gcc/testsuite/
	* gfortran.dg/goacc/enter-exit-data-2.f90: New test.
	* gfortran.dg/goacc/finalize-1.f: Adjust scan output.
	* gfortran.dg/gomp/map-9.f90: Adjust scan output.
	* gfortran.dg/gomp/map-subarray-2.f90: New test.
	* gfortran.dg/gomp/map-subarray.f90: New test.
	* gfortran.dg/gomp/target-enter-exit-data.f90: New test.

libgomp/
	* testsuite/libgomp.fortran/map-subarray.f90: New test.
	* testsuite/libgomp.fortran/map-subarray-2.f90: New test.
	* testsuite/libgomp.fortran/map-subarray-3.f90: New test.
	* testsuite/libgomp.fortran/map-subarray-4.f90: New test.
	* testsuite/libgomp.fortran/map-subarray-6.f90: New test.
	* testsuite/libgomp.fortran/map-subarray-7.f90: New test.
	* testsuite/libgomp.fortran/map-subarray-8.f90: New test.
	* testsuite/libgomp.fortran/map-subcomponents.f90: New test.
	* testsuite/libgomp.fortran/struct-elem-map-1.f90: Adjust for
	descriptor-mapping changes.  Remove XFAIL.

1 files changed, 126 insertions, 37 deletions

diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc
index 26df5b0..9665746 100644
--- a/gcc/gimplify.cc
+++ b/gcc/gimplify.cc
@@ -9151,6 +9151,25 @@ gimplify_omp_depend (tree *list_p, gimple_seq *pre_p)
   return 1;
 }
 
+/* True if mapping node C maps, or unmaps, a (Fortran) array descriptor.  */
+
+static bool
+omp_map_clause_descriptor_p (tree c)
+{
+  if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP)
+    return false;
+
+  if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_TO_PSET)
+    return true;
+
+  if ((OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_RELEASE
+       || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DELETE)
+      && OMP_CLAUSE_RELEASE_DESCRIPTOR (c))
+    return true;
+
+  return false;
+}
+
 /* For a set of mappings describing an array section pointed to by a struct
    (or derived type, etc.) component, create an "alloc" or "release" node to
    insert into a list following a GOMP_MAP_STRUCT node.  For some types of
@@ -9186,9 +9205,7 @@ build_omp_struct_comp_nodes (enum tree_code code, tree grp_start, tree grp_end,
   if (OMP_CLAUSE_CHAIN (grp_start) != grp_end)
     grp_mid = OMP_CLAUSE_CHAIN (grp_start);
 
-  if (grp_mid
-      && OMP_CLAUSE_CODE (grp_mid) == OMP_CLAUSE_MAP
-      && OMP_CLAUSE_MAP_KIND (grp_mid) == GOMP_MAP_TO_PSET)
+  if (grp_mid && omp_map_clause_descriptor_p (grp_mid))
     OMP_CLAUSE_SIZE (c2) = OMP_CLAUSE_SIZE (grp_mid);
   else
     OMP_CLAUSE_SIZE (c2) = TYPE_SIZE_UNIT (ptr_type_node);
@@ -9374,7 +9391,7 @@ omp_get_attachment (omp_mapping_group *grp)
 	return NULL_TREE;
 
       node = OMP_CLAUSE_CHAIN (node);
-      if (node && OMP_CLAUSE_MAP_KIND (node) == GOMP_MAP_TO_PSET)
+      if (node && omp_map_clause_descriptor_p (node))
 	{
 	  gcc_assert (node != grp->grp_end);
 	  node = OMP_CLAUSE_CHAIN (node);
@@ -9469,7 +9486,7 @@ omp_group_last (tree *start_p)
 		     == GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION)
 		 || OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_DETACH
 		 || OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_ALWAYS_POINTER
-		 || OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_TO_PSET))
+		 || omp_map_clause_descriptor_p (nc)))
 	{
 	  tree nc2 = OMP_CLAUSE_CHAIN (nc);
 	  if (OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_DETACH)
@@ -9636,33 +9653,32 @@ omp_group_base (omp_mapping_group *grp, unsigned int *chained,
 	return node;
 
       node = OMP_CLAUSE_CHAIN (node);
-      if (node && OMP_CLAUSE_MAP_KIND (node) == GOMP_MAP_TO_PSET)
+      if (!node)
+	internal_error ("unexpected mapping node");
+      if (omp_map_clause_descriptor_p (node))
 	{
 	  if (node == grp->grp_end)
 	    return *grp->grp_start;
 	  node = OMP_CLAUSE_CHAIN (node);
 	}
-      if (node)
-	switch (OMP_CLAUSE_MAP_KIND (node))
-	  {
-	  case GOMP_MAP_POINTER:
-	  case GOMP_MAP_FIRSTPRIVATE_POINTER:
-	  case GOMP_MAP_FIRSTPRIVATE_REFERENCE:
-	  case GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION:
-	    *firstprivate = OMP_CLAUSE_DECL (node);
-	    return *grp->grp_start;
+      switch (OMP_CLAUSE_MAP_KIND (node))
+	{
+	case GOMP_MAP_POINTER:
+	case GOMP_MAP_FIRSTPRIVATE_POINTER:
+	case GOMP_MAP_FIRSTPRIVATE_REFERENCE:
+	case GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION:
+	  *firstprivate = OMP_CLAUSE_DECL (node);
+	  return *grp->grp_start;
 
-	  case GOMP_MAP_ALWAYS_POINTER:
-	  case GOMP_MAP_ATTACH_DETACH:
-	  case GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION:
-	  case GOMP_MAP_DETACH:
-	    return *grp->grp_start;
+	case GOMP_MAP_ALWAYS_POINTER:
+	case GOMP_MAP_ATTACH_DETACH:
+	case GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION:
+	case GOMP_MAP_DETACH:
+	  return *grp->grp_start;
 
-	  default:
-	    internal_error ("unexpected mapping node");
-	  }
-      else
-	internal_error ("unexpected mapping node");
+	default:
+	  internal_error ("unexpected mapping node");
+	}
       return error_mark_node;
 
     case GOMP_MAP_TO_PSET:
@@ -10010,18 +10026,45 @@ omp_tsort_mapping_groups_1 (omp_mapping_group ***outlist,
 static omp_mapping_group *
 omp_tsort_mapping_groups (vec<omp_mapping_group> *groups,
 			  hash_map<tree_operand_hash_no_se, omp_mapping_group *>
-			    *grpmap)
+			    *grpmap,
+			  bool enter_exit_data)
 {
   omp_mapping_group *grp, *outlist = NULL, **cursor;
   unsigned int i;
+  bool saw_runtime_implicit = false;
 
   cursor = &outlist;
 
   FOR_EACH_VEC_ELT (*groups, i, grp)
     {
       if (grp->mark != PERMANENT)
-	if (!omp_tsort_mapping_groups_1 (&cursor, groups, grpmap, grp))
-	  return NULL;
+	{
+	  if (OMP_CLAUSE_MAP_RUNTIME_IMPLICIT_P (*grp->grp_start))
+	    {
+	      saw_runtime_implicit = true;
+	      continue;
+	    }
+	  if (!omp_tsort_mapping_groups_1 (&cursor, groups, grpmap, grp))
+	    return NULL;
+	}
+    }
+
+  if (!saw_runtime_implicit)
+    return outlist;
+
+  FOR_EACH_VEC_ELT (*groups, i, grp)
+    {
+      if (grp->mark != PERMANENT
+	  && OMP_CLAUSE_MAP_RUNTIME_IMPLICIT_P (*grp->grp_start))
+	{
+	  /* Clear the flag for enter/exit data because it is currently
+	     meaningless for those operations in libgomp.  */
+	  if (enter_exit_data)
+	    OMP_CLAUSE_MAP_RUNTIME_IMPLICIT_P (*grp->grp_start) = 0;
+
+	  if (!omp_tsort_mapping_groups_1 (&cursor, groups, grpmap, grp))
+	    return NULL;
+	}
     }
 
   return outlist;
@@ -10424,6 +10467,11 @@ omp_check_mapping_compatibility (location_t loc,
    mapping.  However, if we have a reference to pointer, make other appropriate
    adjustments to the mapping nodes instead.
 
+   If we have an ATTACH_DETACH node with a Fortran pointer-set (array
+   descriptor) mapping for a derived-type component, and we're also mapping the
+   whole of the derived-type variable on another clause, the pointer-set
+   mapping is removed.
+
    If we have a component access but we're also mapping the whole of the
    containing struct, drop the former access.
 
@@ -10603,6 +10651,17 @@ omp_resolve_clause_dependencies (enum tree_code code,
 		   GOMP_MAP_ATTACH_ZLAS for it.  */
 		if (!base_mapped_to && referenced_ptr_node)
 		  OMP_CLAUSE_SET_MAP_KIND (referenced_ptr_node, zlas_kind);
+
+		omp_mapping_group *struct_group;
+		tree desc;
+		if ((desc = OMP_CLAUSE_CHAIN (*grp->grp_start))
+		    && omp_map_clause_descriptor_p (desc)
+		    && omp_mapped_by_containing_struct (grpmap, decl,
+							&struct_group))
+		  /* If we have a pointer set but we're mapping (or unmapping)
+		     the whole of the containing struct, we can remove the
+		     pointer set mapping.  */
+		  OMP_CLAUSE_CHAIN (*grp->grp_start) = OMP_CLAUSE_CHAIN (desc);
 	      }
 	    else if (TREE_CODE (TREE_TYPE (base_ptr)) == REFERENCE_TYPE
 		     && (TREE_CODE (TREE_TYPE (TREE_TYPE (base_ptr)))
@@ -11001,11 +11060,17 @@ omp_accumulate_sibling_list (enum omp_region_type region_type,
      for the purposes of gathering sibling lists, etc.  */
   /* gcc_assert (base == addr_tokens[base_token]->expr);  */
 
-  bool ptr = (OMP_CLAUSE_MAP_KIND (grp_end) == GOMP_MAP_ALWAYS_POINTER);
   bool attach_detach = ((OMP_CLAUSE_MAP_KIND (grp_end)
 			 == GOMP_MAP_ATTACH_DETACH)
 			|| (OMP_CLAUSE_MAP_KIND (grp_end)
 			    == GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION));
+  bool has_descriptor = false;
+  if (OMP_CLAUSE_CHAIN (*grp_start_p) != grp_end)
+    {
+      tree grp_mid = OMP_CLAUSE_CHAIN (*grp_start_p);
+      if (grp_mid && omp_map_clause_descriptor_p (grp_mid))
+	has_descriptor = true;
+    }
 
   if (!struct_map_to_clause || struct_map_to_clause->get (base) == NULL)
     {
@@ -11028,7 +11093,18 @@ omp_accumulate_sibling_list (enum omp_region_type region_type,
 	 GOMP_MAP_STRUCT into the middle of the old one.  */
       tree *insert_node_pos = reprocessing_struct ? *added_tail : grp_start_p;
 
-      if (ptr || attach_detach)
+      if (has_descriptor)
+	{
+	  tree desc = OMP_CLAUSE_CHAIN (*grp_start_p);
+	  if (code == OMP_TARGET_EXIT_DATA || code == OACC_EXIT_DATA)
+	    OMP_CLAUSE_SET_MAP_KIND (desc, GOMP_MAP_RELEASE);
+	  tree sc = *insert_node_pos;
+	  OMP_CLAUSE_CHAIN (l) = desc;
+	  OMP_CLAUSE_CHAIN (*grp_start_p) = OMP_CLAUSE_CHAIN (desc);
+	  OMP_CLAUSE_CHAIN (desc) = sc;
+	  *insert_node_pos = l;
+	}
+      else if (attach_detach)
 	{
 	  tree extra_node;
 	  tree alloc_node
@@ -11259,7 +11335,7 @@ omp_accumulate_sibling_list (enum omp_region_type region_type,
 	  || OMP_CLAUSE_MAP_KIND (*sc) == GOMP_MAP_ATTACH_DETACH)
 	sc = &OMP_CLAUSE_CHAIN (*sc);
       for (i = 0; i < elems; i++, sc = &OMP_CLAUSE_CHAIN (*sc))
-	if ((ptr || attach_detach) && sc == grp_start_p)
+	if (attach_detach && sc == grp_start_p)
 	  break;
 	else if (TREE_CODE (OMP_CLAUSE_DECL (*sc)) != COMPONENT_REF
 		 && TREE_CODE (OMP_CLAUSE_DECL (*sc)) != INDIRECT_REF
@@ -11315,7 +11391,7 @@ omp_accumulate_sibling_list (enum omp_region_type region_type,
 		|| (known_eq (coffset, offset)
 		    && maybe_lt (cbitpos, bitpos)))
 	      {
-		if (ptr || attach_detach)
+		if (attach_detach)
 		  scp = sc;
 		else
 		  break;
@@ -11331,7 +11407,9 @@ omp_accumulate_sibling_list (enum omp_region_type region_type,
 	     the list manipulation below.  We only need to handle the (pointer
 	     or reference) attach/detach case.  */
 	  tree extra_node, alloc_node;
-	  if (attach_detach)
+	  if (has_descriptor)
+	    gcc_unreachable ();
+	  else if (attach_detach)
 	    alloc_node = build_omp_struct_comp_nodes (code, *grp_start_p,
 						      grp_end, &extra_node);
 	  else
@@ -11364,7 +11442,17 @@ omp_accumulate_sibling_list (enum omp_region_type region_type,
 	  return NULL;
 	}
 
-      if (ptr || attach_detach)
+      if (has_descriptor)
+	{
+	  tree desc = OMP_CLAUSE_CHAIN (*grp_start_p);
+	  if (code == OMP_TARGET_EXIT_DATA
+	      || code == OACC_EXIT_DATA)
+	    OMP_CLAUSE_SET_MAP_KIND (desc, GOMP_MAP_RELEASE);
+	  omp_siblist_move_node_after (desc,
+				       &OMP_CLAUSE_CHAIN (*grp_start_p),
+				       scp ? scp : sc);
+	}
+      else if (attach_detach)
 	{
 	  tree cl = NULL_TREE, extra_node;
 	  tree alloc_node = build_omp_struct_comp_nodes (code, *grp_start_p,
@@ -11509,8 +11597,7 @@ omp_build_struct_sibling_lists (enum tree_code code,
 	     as a struct (the GOMP_MAP_POINTER following will have the form
 	     "var.data", but such mappings are handled specially).  */
 	  tree grpmid = OMP_CLAUSE_CHAIN (*grp_start_p);
-	  if (OMP_CLAUSE_CODE (grpmid) == OMP_CLAUSE_MAP
-	      && OMP_CLAUSE_MAP_KIND (grpmid) == GOMP_MAP_TO_PSET
+	  if (omp_map_clause_descriptor_p (grpmid)
 	      && DECL_P (OMP_CLAUSE_DECL (grpmid)))
 	    continue;
 	}
@@ -11786,6 +11873,8 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 					  list_p);
 
 	  omp_mapping_group *outlist = NULL;
+	  bool enter_exit = (code == OMP_TARGET_ENTER_DATA
+			     || code == OMP_TARGET_EXIT_DATA);
 
 	  /* Topological sorting may fail if we have duplicate nodes, which
 	     we should have detected and shown an error for already.  Skip
@@ -11800,7 +11889,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 	  groups = omp_gather_mapping_groups (list_p);
 	  grpmap = omp_index_mapping_groups (groups);
 
-	  outlist = omp_tsort_mapping_groups (groups, grpmap);
+	  outlist = omp_tsort_mapping_groups (groups, grpmap, enter_exit);
 	  outlist = omp_segregate_mapping_groups (outlist);
 	  list_p = omp_reorder_mapping_groups (groups, outlist, list_p);