[PATCH 8/8] OpenMP 5.0: [WIP, RFC] Clause ordering for OpenMP 5.0 (topological sorting by base pointer)

[Julian Brown <julian@codesourcery.com>]

This patch reimplements the omp_target_reorder_clauses function in
anticipation of supporting "deeper" struct mappings (that is, with
several structure dereference operators, or similar).

The idea is that in place of the (possibly quadratic) algorithm in
omp_target_reorder_clauses that greedily moves clauses containing
addresses that are subexpressions of other addresses before those other
addresses, we employ a topological sort algorithm to calculate a proper
order for map clauses. This should run in linear time, and hopefully
handles degenerate cases where multiple "levels" of indirect accesses
are present on a given directive.

The new method also takes care to keep clause groups together, addressing
the concerns raised in:

  https://gcc.gnu.org/pipermail/gcc-patches/2021-May/570501.html

To figure out if some given clause depends on a base pointer in another
clause, we strip off the outer layers of the address expression, and check
(via a tree_operand_hash hash table we have built) if the result is a
"base pointer" as defined in OpenMP 5.0 (1.2.6 Data Terminology). There
are some subtleties involved, however:

 - We must treat MEM_REF with zero offset the same as INDIRECT_REF.
   This should probably be fixed in the front ends instead so we always
   use a canonical form (probably INDIRECT_REF). The following patch
   shows one instance of the problem, but there may be others:

   https://gcc.gnu.org/pipermail/gcc-patches/2021-May/571382.html

 - Mapping a whole struct implies mapping each of that struct's
   elements, which may be base pointers. Because those base pointers
   aren't necessarily explicitly referenced in the directive in question,
   we treat the whole-struct mapping as a dependency instead.

 - We also need to special-case handling for "*struct_var" (including
   "*this"), so the un-dereferenced variable is treated as the dependency
   -- which feels a bit wrong. (A subsequent GOMP_MAP_POINTER handles the
   pointer itself for those types of mapping, but the current approach
   only processes the first node in each group.)

Jakub, Chung-Lin -- does this approach seem reasonable? Any comments at
this stage?

2021-08-10  Julian Brown  <julian@codesourcery.com>

gcc/
	* gimplify.c (is_or_contains_p, omp_target_reorder_clauses): Delete
	function.
	(omp_tsort_mark): Add enum.
	(omp_mapping_group): Add struct.
	(omp_get_base_pointer, omp_gather_mapping_groups,
	omp_index_mapping_groups, omp_tsort_mapping_groups_1,
	omp_tsort_mapping_groups, omp_segregate_mapping_groups,
	omp_reorder_mapping_groups): New functions.
	(gimplify_scan_omp_clauses): Call above functions instead of
	omp_target_reorder_clauses.

gcc/testsuite/
	* g++.dg/gomp/target-this-3.C: Adjust expected output.
	* g++.dg/gomp/target-this-4.C: Likewise.
---
 gcc/gimplify.c                            | 598 +++++++++++++++-------
 gcc/testsuite/g++.dg/gomp/target-this-3.C |   2 +-
 gcc/testsuite/g++.dg/gomp/target-this-4.C |   2 +-
 3 files changed, 411 insertions(+), 191 deletions(-)

diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index 0ef2dbde710..ca106ef7acf 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -8534,29 +8534,6 @@ extract_base_bit_offset (tree base, tree *base_ind, tree *base_ref,
   return base;
 }
 
-/* Returns true if EXPR is or contains (as a sub-component) BASE_PTR.  */
-
-static bool
-is_or_contains_p (tree expr, tree base_ptr)
-{
-  if ((TREE_CODE (expr) == INDIRECT_REF && TREE_CODE (base_ptr) == MEM_REF)
-      || (TREE_CODE (expr) == MEM_REF && TREE_CODE (base_ptr) == INDIRECT_REF))
-    return operand_equal_p (TREE_OPERAND (expr, 0),
-			    TREE_OPERAND (base_ptr, 0));
-  while (!operand_equal_p (expr, base_ptr))
-    {
-      if (TREE_CODE (base_ptr) == COMPOUND_EXPR)
-	base_ptr = TREE_OPERAND (base_ptr, 1);
-      if (TREE_CODE (base_ptr) == COMPONENT_REF
-	  || TREE_CODE (base_ptr) == POINTER_PLUS_EXPR
-	  || TREE_CODE (base_ptr) == SAVE_EXPR)
-	base_ptr = TREE_OPERAND (base_ptr, 0);
-      else
-	break;
-    }
-  return operand_equal_p (expr, base_ptr);
-}
-
 /* Remove COMPONENT_REFS and indirections from EXPR.  */
 
 static tree
@@ -8599,184 +8576,413 @@ aggregate_base_p (tree expr)
   return false;
 }
 
-/* Implement OpenMP 5.x map ordering rules for target directives. There are
-   several rules, and with some level of ambiguity, hopefully we can at least
-   collect the complexity here in one place.  */
+enum omp_tsort_mark {
+  UNVISITED,
+  TEMPORARY,
+  PERMANENT
+};
+
+struct omp_mapping_group {
+  tree *grp_start;
+  tree grp_end;
+  omp_tsort_mark mark;
+  struct omp_mapping_group *sibling;
+  struct omp_mapping_group *next;
+};
+
+/* Return the OpenMP "base pointer" of an expression EXPR, or NULL if there
+   isn't one.  This needs improvement.  */
+
+static tree
+omp_get_base_pointer (tree expr)
+{
+  while (TREE_CODE (expr) == COMPONENT_REF
+	 && (DECL_P (TREE_OPERAND (expr, 0))
+	     || (TREE_CODE (TREE_OPERAND (expr, 0)) == COMPONENT_REF)
+	     || TREE_CODE (TREE_OPERAND (expr, 0)) == INDIRECT_REF
+	     || (TREE_CODE (TREE_OPERAND (expr, 0)) == MEM_REF
+		 && integer_zerop (TREE_OPERAND (TREE_OPERAND (expr, 0), 1)))))
+    expr = TREE_OPERAND (expr, 0);
+
+  if (DECL_P (expr))
+    return NULL_TREE;
+
+  while (TREE_CODE (expr) == ARRAY_REF)
+    expr = TREE_OPERAND (expr, 0);
+
+  if (TREE_CODE (expr) == INDIRECT_REF
+      || TREE_CODE (expr) == MEM_REF)
+    {
+      expr = TREE_OPERAND (expr, 0);
+      while (TREE_CODE (expr) == COMPOUND_EXPR)
+	expr = TREE_OPERAND (expr, 1);
+      if (TREE_CODE (expr) == POINTER_PLUS_EXPR)
+	expr = TREE_OPERAND (expr, 0);
+      if (TREE_CODE (expr) == SAVE_EXPR)
+	expr = TREE_OPERAND (expr, 0);
+      STRIP_NOPS (expr);
+      return expr;
+    }
+
+  return NULL_TREE;
+}
+
+/* Walk through LIST_P, and return a list of groups of mappings found (e.g.
+   OMP_CLAUSE_MAP with GOMP_MAP_{TO/FROM/TOFROM} followed by one or two
+   associated GOMP_MAP_POINTER mappings).  Return a vector of omp_mapping_group
+   if we have more than one such group, else return NULL.  */
+
+static vec<omp_mapping_group> *
+omp_gather_mapping_groups (tree *list_p)
+{
+  vec<omp_mapping_group> *groups = new vec<omp_mapping_group> ();
+
+  for (tree *cp = list_p; *cp; cp = &OMP_CLAUSE_CHAIN (*cp))
+    {
+      tree c = *cp, nc, *grp_last_p = cp;
+
+      switch (OMP_CLAUSE_CODE (c))
+	{
+	default:
+	  continue;
+
+	case OMP_CLAUSE_MAP:
+	  nc = OMP_CLAUSE_CHAIN (c);
+	  while (nc
+		 && OMP_CLAUSE_CODE (nc) == OMP_CLAUSE_MAP
+		 && ((OMP_CLAUSE_MAP_KIND (nc)
+		      == GOMP_MAP_FIRSTPRIVATE_REFERENCE)
+		     || (OMP_CLAUSE_MAP_KIND (nc)
+			 == GOMP_MAP_FIRSTPRIVATE_POINTER)
+		     || OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_ATTACH_DETACH
+		     || OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_POINTER
+		     || (OMP_CLAUSE_MAP_KIND (nc)
+			 == GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION)
+		     || OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_ALWAYS_POINTER
+		     || OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_TO_PSET))
+	    {
+	      grp_last_p = &OMP_CLAUSE_CHAIN (c);
+	      c = nc;
+	      tree nc2 = OMP_CLAUSE_CHAIN (nc);
+	      if (nc2
+		  && OMP_CLAUSE_CODE (nc2) == OMP_CLAUSE_MAP
+		  && (OMP_CLAUSE_MAP_KIND (nc)
+		      == GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION)
+		  && OMP_CLAUSE_MAP_KIND (nc2) == GOMP_MAP_ATTACH)
+		{
+		  grp_last_p = &OMP_CLAUSE_CHAIN (nc);
+		  c = nc2;
+		  nc2 = OMP_CLAUSE_CHAIN (nc2);
+		}
+	       nc = nc2;
+	    }
+	  break;
+	}
+
+      omp_mapping_group grp;
+
+      grp.grp_start = cp;
+      grp.grp_end = *grp_last_p;
+      grp.mark = UNVISITED;
+      grp.sibling = NULL;
+      grp.next = NULL;
+      groups->safe_push (grp);
+
+      cp = grp_last_p;
+    }
+
+  if (groups->length () > 1)
+    return groups;
+  else
+    {
+      delete groups;
+      return NULL;
+    }
+}
+
+/* Given a vector of omp_mapping_groups, build a hash table so we can look up
+   nodes by tree_operand_hash.  */
+
+static hash_map<tree_operand_hash, omp_mapping_group *> *
+omp_index_mapping_groups (vec<omp_mapping_group> *groups)
+{
+  hash_map<tree_operand_hash, omp_mapping_group *> *grpmap
+    = new hash_map<tree_operand_hash, omp_mapping_group *>;
+
+  omp_mapping_group *grp;
+  unsigned int i;
+
+  FOR_EACH_VEC_ELT (*groups, i, grp)
+    {
+      tree decl = OMP_CLAUSE_DECL (*grp->grp_start);
+
+      /* Sometimes we see zero-offset MEM_REF instead of INDIRECT_REF, meaning
+	 node-hash lookups don't work.  This is a workaround for that, but
+	 ideally we should just create the INDIRECT_REF at source instead.
+	 FIXME.  */
+      if (TREE_CODE (decl) == MEM_REF
+	  && integer_zerop (TREE_OPERAND (decl, 1)))
+	decl = build1 (INDIRECT_REF, TREE_TYPE (decl), TREE_OPERAND (decl, 0));
+
+      omp_mapping_group **prev = grpmap->get (decl);
+
+      if (prev)
+	{
+	  /* Mapping the same thing twice is normally diagnosed as an error,
+	     but can happen under some circumstances, e.g. in pr99928-16.c,
+	     the directive:
+
+	     #pragma omp target simd reduction(+:a[:3]) \
+				     map(always, tofrom: a[:6])
+	     ...
+
+	     will result in two "a[0]" mappings (of different sizes).  */
+
+	  grp->sibling = (*prev)->sibling;
+	  (*prev)->sibling = grp;
+	}
+      else
+	grpmap->put (decl, grp);
+    }
+  return grpmap;
+}
+
+/* Helper function for omp_tsort_mapping_groups.  */
 
 static void
-omp_target_reorder_clauses (tree *list_p)
+omp_tsort_mapping_groups_1 (omp_mapping_group ***outlist,
+			    vec<omp_mapping_group> *groups,
+			    hash_map<tree_operand_hash, omp_mapping_group *>
+			      *grpmap,
+			    omp_mapping_group *grp)
 {
-  /* Collect refs to alloc/release/delete maps.  */
-  auto_vec<tree, 32> ard;
-  tree *cp = list_p;
-  while (*cp != NULL_TREE)
-    if (OMP_CLAUSE_CODE (*cp) == OMP_CLAUSE_MAP
-	&& (OMP_CLAUSE_MAP_KIND (*cp) == GOMP_MAP_ALLOC
-	    || OMP_CLAUSE_MAP_KIND (*cp) == GOMP_MAP_RELEASE
-	    || OMP_CLAUSE_MAP_KIND (*cp) == GOMP_MAP_DELETE))
-      {
-	/* Unlink cp and push to ard.  */
-	tree c = *cp;
-	tree nc = OMP_CLAUSE_CHAIN (c);
-	*cp = nc;
-	ard.safe_push (c);
-
-	/* Any associated pointer type maps should also move along.  */
-	while (*cp != NULL_TREE
-	       && OMP_CLAUSE_CODE (*cp) == OMP_CLAUSE_MAP
-	       && (OMP_CLAUSE_MAP_KIND (*cp) == GOMP_MAP_FIRSTPRIVATE_REFERENCE
-		   || OMP_CLAUSE_MAP_KIND (*cp) == GOMP_MAP_FIRSTPRIVATE_POINTER
-		   || OMP_CLAUSE_MAP_KIND (*cp) == GOMP_MAP_ATTACH_DETACH
-		   || OMP_CLAUSE_MAP_KIND (*cp) == GOMP_MAP_POINTER
-		   || OMP_CLAUSE_MAP_KIND (*cp) == GOMP_MAP_ALWAYS_POINTER
-		   || OMP_CLAUSE_MAP_KIND (*cp) == GOMP_MAP_TO_PSET))
-	  {
-	    c = *cp;
-	    nc = OMP_CLAUSE_CHAIN (c);
-	    *cp = nc;
-	    ard.safe_push (c);
-	  }
-      }
-    else
-      cp = &OMP_CLAUSE_CHAIN (*cp);
-
-  /* Link alloc/release/delete maps to the end of list.  */
-  for (unsigned int i = 0; i < ard.length (); i++)
+  if (grp->mark == PERMANENT)
+    return;
+  if (grp->mark == TEMPORARY)
     {
-      *cp = ard[i];
-      cp = &OMP_CLAUSE_CHAIN (ard[i]);
+      error ("not a DAG?");
+      return;
     }
-  *cp = NULL_TREE;
+  grp->mark = TEMPORARY;
 
-  /* OpenMP 5.0 requires that pointer variables are mapped before
-     its use as a base-pointer.  */
-  auto_vec<tree *, 32> atf;
-  for (tree *cp = list_p; *cp; cp = &OMP_CLAUSE_CHAIN (*cp))
-    if (OMP_CLAUSE_CODE (*cp) == OMP_CLAUSE_MAP)
-      {
-	/* Collect alloc, to, from, to/from clause tree pointers.  */
-	gomp_map_kind k = OMP_CLAUSE_MAP_KIND (*cp);
-	if (k == GOMP_MAP_ALLOC
-	    || k == GOMP_MAP_TO
-	    || k == GOMP_MAP_FROM
-	    || k == GOMP_MAP_TOFROM
-	    || k == GOMP_MAP_ALWAYS_TO
-	    || k == GOMP_MAP_ALWAYS_FROM
-	    || k == GOMP_MAP_ALWAYS_TOFROM)
-	  atf.safe_push (cp);
-      }
+  tree decl = OMP_CLAUSE_DECL (*grp->grp_start);
 
-  for (unsigned int i = 0; i < atf.length (); i++)
-    if (atf[i])
-      {
-	tree *cp = atf[i];
-	tree decl = OMP_CLAUSE_DECL (*cp);
-	if (TREE_CODE (decl) == INDIRECT_REF || TREE_CODE (decl) == MEM_REF)
-	  {
-	    tree base_ptr = TREE_OPERAND (decl, 0);
-	    STRIP_TYPE_NOPS (base_ptr);
-	    for (unsigned int j = i + 1; j < atf.length (); j++)
-	      if (atf[j])
-		{
-		  tree *cp2 = atf[j];
-		  tree decl2 = OMP_CLAUSE_DECL (*cp2);
+  while (decl)
+    {
+      tree base = omp_get_base_pointer (decl);
 
-		  decl2 = OMP_CLAUSE_DECL (*cp2);
-		  if (is_or_contains_p (decl2, base_ptr))
-		    {
-		      /* Move *cp2 to before *cp.  */
-		      tree c = *cp2;
-		      *cp2 = OMP_CLAUSE_CHAIN (c);
-		      OMP_CLAUSE_CHAIN (c) = *cp;
-		      *cp = c;
+      if (!base)
+	break;
 
-		      if (*cp2 != NULL_TREE
-			  && OMP_CLAUSE_CODE (*cp2) == OMP_CLAUSE_MAP
-			  && OMP_CLAUSE_MAP_KIND (*cp2) == GOMP_MAP_ALWAYS_POINTER)
-			{
-			  tree c2 = *cp2;
-			  *cp2 = OMP_CLAUSE_CHAIN (c2);
-			  OMP_CLAUSE_CHAIN (c2) = OMP_CLAUSE_CHAIN (c);
-			  OMP_CLAUSE_CHAIN (c) = c2;
-			}
+      omp_mapping_group **innerp = grpmap->get (base);
 
-		      atf[j] = NULL;
-		  }
-		}
-	  }
-      }
+      /* We should treat whole-structure mappings as if all (pointer, in this
+	 case) members are mapped as individual list items.  Check if we have
+	 such a whole-structure mapping, if we don't have an explicit reference
+	 to the pointer member itself.  */
+      if (!innerp && TREE_CODE (base) == COMPONENT_REF)
+	{
+	  while (TREE_CODE (base) == COMPONENT_REF
+		 && (DECL_P (TREE_OPERAND (base, 0))
+		     || (TREE_CODE (TREE_OPERAND (base, 0)) == COMPONENT_REF)))
+	    base = TREE_OPERAND (base, 0);
 
-  /* For attach_detach map clauses, if there is another map that maps the
-     attached/detached pointer, make sure that map is ordered before the
-     attach_detach.  */
-  atf.truncate (0);
-  for (tree *cp = list_p; *cp; cp = &OMP_CLAUSE_CHAIN (*cp))
-    if (OMP_CLAUSE_CODE (*cp) == OMP_CLAUSE_MAP)
-      {
-	/* Collect alloc, to, from, to/from clauses, and
-	   always_pointer/attach_detach clauses.  */
-	gomp_map_kind k = OMP_CLAUSE_MAP_KIND (*cp);
-	if (k == GOMP_MAP_ALLOC
-	    || k == GOMP_MAP_TO
-	    || k == GOMP_MAP_FROM
-	    || k == GOMP_MAP_TOFROM
-	    || k == GOMP_MAP_ALWAYS_TO
-	    || k == GOMP_MAP_ALWAYS_FROM
-	    || k == GOMP_MAP_ALWAYS_TOFROM
-	    || k == GOMP_MAP_ATTACH_DETACH
-	    || k == GOMP_MAP_ALWAYS_POINTER)
-	  atf.safe_push (cp);
-      }
+	  innerp = grpmap->get (base);
+	}
+      else if (!innerp
+	       && TREE_CODE (decl) == COMPONENT_REF
+	       && TREE_CODE (TREE_OPERAND (decl, 0)) == INDIRECT_REF)
+	{
+	  /* As a special case, handle mappings like "((struct T *) this)->ptr",
+	     where "*(struct T *) this", but not "this" by itself, is mapped by
+	     another clause.  */
+	  innerp = grpmap->get (TREE_OPERAND (decl, 0));
+	}
 
-  for (unsigned int i = 0; i < atf.length (); i++)
-    if (atf[i])
-      {
-	tree *cp = atf[i];
-	tree ptr = OMP_CLAUSE_DECL (*cp);
-	STRIP_TYPE_NOPS (ptr);
-	if (OMP_CLAUSE_MAP_KIND (*cp) == GOMP_MAP_ATTACH_DETACH)
-	  for (unsigned int j = i + 1; j < atf.length (); j++)
-	    {
-	      tree *cp2 = atf[j];
-	      tree decl2 = OMP_CLAUSE_DECL (*cp2);
-	      if (OMP_CLAUSE_MAP_KIND (*cp2) != GOMP_MAP_ATTACH_DETACH
-		  && OMP_CLAUSE_MAP_KIND (*cp2) != GOMP_MAP_ALWAYS_POINTER
-		  && is_or_contains_p (decl2, ptr))
-		{
-		  /* Move *cp2 to before *cp.  */
-		  tree c = *cp2;
-		  *cp2 = OMP_CLAUSE_CHAIN (c);
-		  OMP_CLAUSE_CHAIN (c) = *cp;
-		  *cp = c;
-		  atf[j] = NULL;
+      if (innerp)
+	{
+	  for (omp_mapping_group *w = *innerp; w; w = w->sibling)
+	    omp_tsort_mapping_groups_1 (outlist, groups, grpmap, w);
+	  break;
+	}
 
-		  /* If decl2 is of the form '*decl2_opnd0', and followed by an
-		     ALWAYS_POINTER or ATTACH_DETACH of 'decl2_opnd0', move the
-		     pointer operation along with *cp2. This can happen for C++
-		     reference sequences.  */
-		  if (j + 1 < atf.length ()
-		      && (TREE_CODE (decl2) == INDIRECT_REF
-			  || TREE_CODE (decl2) == MEM_REF))
-		    {
-		      tree *cp3 = atf[j + 1];
-		      tree decl3 = OMP_CLAUSE_DECL (*cp3);
-		      tree decl2_opnd0 = TREE_OPERAND (decl2, 0);
-		      if ((OMP_CLAUSE_MAP_KIND (*cp3) == GOMP_MAP_ALWAYS_POINTER
-			   || OMP_CLAUSE_MAP_KIND (*cp3) == GOMP_MAP_ATTACH_DETACH)
-			  && operand_equal_p (decl3, decl2_opnd0))
-			{
-			  /* Also move *cp3 to before *cp.  */
-			  c = *cp3;
-			  *cp2 = OMP_CLAUSE_CHAIN (c);
-			  OMP_CLAUSE_CHAIN (c) = *cp;
-			  *cp = c;
-			  atf[j + 1] = NULL;
-			  j += 1;
-			}
-		    }
-		}
-	    }
-      }
+      decl = base;
+    }
+
+  grp->mark = PERMANENT;
+
+  /* Emit grp to output list.  */
+
+  **outlist = grp;
+  *outlist = &grp->next;
+}
+
+/* Topologically sort GROUPS, so that OMP 5.0-defined base pointers come
+   before mappings that use those pointers.  This is an implementation of the
+   depth-first search algorithm, described e.g. at:
+
+     https://en.wikipedia.org/wiki/Topological_sorting
+*/
+
+static omp_mapping_group *
+omp_tsort_mapping_groups (vec<omp_mapping_group> *groups,
+			  hash_map<tree_operand_hash, omp_mapping_group *>
+			    *grpmap)
+{
+  omp_mapping_group *grp, *outlist = NULL, **cursor;
+  unsigned int i;
+
+  cursor = &outlist;
+
+  FOR_EACH_VEC_ELT (*groups, i, grp)
+    {
+      if (grp->mark != PERMANENT)
+	omp_tsort_mapping_groups_1 (&cursor, groups, grpmap, grp);
+    }
+
+  return outlist;
+}
+
+/* Split INLIST into two parts, moving groups corresponding to
+   ALLOC/RELEASE/DELETE mappings to one list, and other mappings to another.
+   The former list is then appended to the latter.  Each sub-list retains the
+   order of the original list.  */
+
+static omp_mapping_group *
+omp_segregate_mapping_groups (omp_mapping_group *inlist)
+{
+  omp_mapping_group *ard_groups = NULL, *tf_groups = NULL;
+  omp_mapping_group **ard_tail = &ard_groups, **tf_tail = &tf_groups;
+
+  for (omp_mapping_group *w = inlist; w;)
+    {
+      tree c = *w->grp_start;
+      omp_mapping_group *next = w->next;
+
+      gcc_assert (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP);
+
+      switch (OMP_CLAUSE_MAP_KIND (c))
+	{
+	case GOMP_MAP_ALLOC:
+	case GOMP_MAP_RELEASE:
+	case GOMP_MAP_DELETE:
+	  *ard_tail = w;
+	  w->next = NULL;
+	  ard_tail = &w->next;
+	  break;
+
+	default:
+	  *tf_tail = w;
+	  w->next = NULL;
+	  tf_tail = &w->next;
+	}
+
+      w = next;
+    }
+
+  /* Now splice the lists together...  */
+  *tf_tail = ard_groups;
+
+  return tf_groups;
+}
+
+/* Given a list LIST_P containing groups of mappings given by GROUPS, reorder
+   those groups based on the output list of omp_tsort_mapping_groups --
+   singly-linked, threaded through each element's NEXT pointer starting at
+   HEAD.  Each list element appears exactly once in that linked list.
+
+   Each element of GROUPS may correspond to one or several mapping nodes.
+   Node groups are kept together, and in the reordered list, the positions of
+   the original groups are reused for the positions of the reordered list.
+   Hence if we have e.g.
+
+     {to ptr ptr} firstprivate {tofrom ptr} ...
+      ^             ^           ^
+      first group  non-"map"    second group
+
+   and say the second group contains a base pointer for the first so must be
+   moved before it, the resulting list will contain:
+
+     {tofrom ptr} firstprivate {to ptr ptr} ...
+      ^ prev. second group      ^ prev. first group
+*/
+
+static tree *
+omp_reorder_mapping_groups (vec<omp_mapping_group> *groups,
+			    omp_mapping_group *head,
+			    tree *list_p)
+{
+  omp_mapping_group *grp;
+  unsigned int i;
+  unsigned numgroups = groups->length ();
+  auto_vec<tree> old_heads (numgroups);
+  auto_vec<tree> new_heads (numgroups);
+  auto_vec<tree> old_succs (numgroups);
+  bool map_at_start = (list_p == (*groups)[0].grp_start);
+
+  tree *new_grp_tail = NULL;
+
+  /* Stash the start & end nodes of each mapping group before we start
+     modifying the list.  */
+  FOR_EACH_VEC_ELT (*groups, i, grp)
+    {
+      old_heads.quick_push (*grp->grp_start);
+      old_succs.quick_push (OMP_CLAUSE_CHAIN (grp->grp_end));
+    }
+
+  /* And similarly, the heads of the groups in the order we want to rearrange
+     the list to.  */
+  for (omp_mapping_group *w = head; w; w = w->next)
+    new_heads.quick_push (*w->grp_start);
+
+  FOR_EACH_VEC_ELT (*groups, i, grp)
+    {
+      gcc_assert (head);
+
+      if (new_grp_tail && old_succs[i - 1] == old_heads[i])
+	{
+	  /* a {b c d} {e f g} h i j   (original)
+	     -->
+	     a {k l m} {e f g} h i j   (inserted new group on last iter)
+	     -->
+	     a {k l m} {n o p} h i j   (this time, chain last group to new one)
+		      ^new_grp_tail
+	  */
+	  *new_grp_tail = new_heads[i];
+	}
+      else if (new_grp_tail)
+	{
+	  /* a {b c d} e {f g h} i j k   (original)
+	     -->
+	     a {l m n} e {f g h} i j k   (gap after last iter's group)
+	     -->
+	     a {l m n} e {o p q} h i j   (chain last group to old successor)
+		      ^new_grp_tail
+	   */
+	  *new_grp_tail = old_succs[i - 1];
+	}
+      else
+	{
+	  /* The first inserted group -- point to new group, and leave end
+	     open.
+	     a {b c d} e f
+	     -->
+	     a {g h i...
+	  */
+	  *grp->grp_start = new_heads[i];
+	}
+
+      new_grp_tail = &OMP_CLAUSE_CHAIN (head->grp_end);
+
+      head = head->next;
+    }
+
+  if (new_grp_tail)
+    *new_grp_tail = old_succs[numgroups - 1];
+
+  gcc_assert (!head);
+
+  return map_at_start ? (*groups)[0].grp_start : list_p;
 }
 
 /* DECL is supposed to have lastprivate semantics in the outer contexts
@@ -9503,7 +9709,21 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
       || code == OMP_TARGET_DATA
       || code == OMP_TARGET_ENTER_DATA
       || code == OMP_TARGET_EXIT_DATA)
-    omp_target_reorder_clauses (list_p);
+    {
+      vec<omp_mapping_group> *groups;
+      groups = omp_gather_mapping_groups (list_p);
+      if (groups)
+	{
+	  hash_map<tree_operand_hash, omp_mapping_group *> *grpmap;
+	  grpmap = omp_index_mapping_groups (groups);
+	  omp_mapping_group *outlist
+	    = omp_tsort_mapping_groups (groups, grpmap);
+	  outlist = omp_segregate_mapping_groups (outlist);
+	  list_p = omp_reorder_mapping_groups (groups, outlist, list_p);
+	  delete grpmap;
+	  delete groups;
+	}
+    }
 
   while ((c = *list_p) != NULL)
     {
diff --git a/gcc/testsuite/g++.dg/gomp/target-this-3.C b/gcc/testsuite/g++.dg/gomp/target-this-3.C
index 2755b4b58bd..c2af5338a1a 100644
--- a/gcc/testsuite/g++.dg/gomp/target-this-3.C
+++ b/gcc/testsuite/g++.dg/gomp/target-this-3.C
@@ -100,6 +100,6 @@ int main (void)
   return 0;
 }
 
-/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(n\) map\(tofrom:\*this \[len: [0-9]+\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\) map\(alloc:\*_[0-9]+ \[pointer assign, zero-length array section, bias: 0\]\) map\(attach:this->refptr \[bias: 0\]\) map\(from:mapped \[len: [0-9]+\]\) map\(alloc:\*_[0-9+] \[len: 0\]\)} "gimple" } } */
+/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(n\) map\(tofrom:\*this \[len: [0-9]+\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\) map\(from:mapped \[len: [0-9]+\]\) map\(alloc:\*_[0-9+] \[len: 0\]\) map\(alloc:\*_[0-9]+ \[pointer assign, zero-length array section, bias: 0\]\) map\(attach:this->refptr \[bias: 0\]\)} "gimple" } } */
 
 /* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(n\) map\(tofrom:\*this \[len: [0-9]+\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\) map\(attach_zero_length_array_section:this->ptr \[bias: 0\]\) map\(from:mapped \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\)} "gimple" } } */
diff --git a/gcc/testsuite/g++.dg/gomp/target-this-4.C b/gcc/testsuite/g++.dg/gomp/target-this-4.C
index 3703762f45a..432b0b35bad 100644
--- a/gcc/testsuite/g++.dg/gomp/target-this-4.C
+++ b/gcc/testsuite/g++.dg/gomp/target-this-4.C
@@ -104,4 +104,4 @@ int main (void)
 
 /* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(n\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\) map\(to:\*__closure \[len: [0-9]+\]\) map\(firstprivate:__closure \[pointer assign, bias: 0\]\) map\(struct:\*__closure \[len: 1\]\) map\(alloc:__closure->__this \[len: [0-9]+\]\) map\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\) map\(always_pointer:__closure->__this \[pointer assign, bias: 0\]\) map\(attach_zero_length_array_section:_[0-9]+->ptr \[bias: 0\]\) map\(from:mapped \[len: 1\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\)} "gimple" } } */
 
-/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(n\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\) map\(to:\*__closure \[len: [0-9]+\]\) map\(firstprivate:__closure \[pointer assign, bias: 0\]\) map\(struct:\*__closure \[len: 1\]\) map\(alloc:__closure->__this \[len: [0-9]+\]\) map\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\) map\(always_pointer:__closure->__this \[pointer assign, bias: 0\]\) map\(alloc:\*_[0-9]+ \[pointer assign, zero-length array section, bias: 0\]\) map\(attach:_[0-9]+->refptr \[bias: 0\]\) map\(from:mapped \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\)} "gimple" } } */
+/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(n\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\) map\(to:\*__closure \[len: [0-9]+\]\) map\(firstprivate:__closure \[pointer assign, bias: 0\]\) map\(struct:\*__closure \[len: 1\]\) map\(alloc:__closure->__this \[len: [0-9]+\]\) map\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\) map\(always_pointer:__closure->__this \[pointer assign, bias: 0\]\) map\(from:mapped \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) map\(alloc:\*_[0-9]+ \[pointer assign, zero-length array section, bias: 0\]\) map\(attach:_[0-9]+->refptr \[bias: 0\]\)} "gimple" } } */
-- 
2.29.2