diff mbox series

[v2,01/11] OpenMP 5.0: Clause ordering for OpenMP 5.0 (topological sorting by base pointer)

Message ID d2d49c37b13b0e953a26c0da0a19235336a4da9d.1647619144.git.julian@codesourcery.com
State New
Headers show
Series OpenMP 5.0: C & C++ "declare mapper" support (plus struct rework, etc.) | expand

Commit Message

Julian Brown March 18, 2022, 4:24 p.m. UTC
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.

This version of the patch has been moved to the front of the patch queue,
thus isn't dependent on any of the following struct-rework patches.

2021-11-23  Julian Brown  <julian@codesourcery.com>

gcc/
	* gimplify.c (is_or_contains_p, omp_target_reorder_clauses): Delete
	functions.
	(omp_tsort_mark): Add enum.
	(omp_mapping_group): Add struct.
	(debug_mapping_group, omp_get_base_pointer, omp_get_attachment,
	omp_group_last, omp_gather_mapping_groups, omp_group_base,
	omp_index_mapping_groups, omp_containing_struct,
	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, unless we've seen an error.
	* omp-low.c (scan_sharing_clauses): Avoid strict test if we haven't
	sorted mapping groups.

gcc/testsuite/
	* g++.dg/gomp/target-lambda-1.C: Adjust expected output.
	* g++.dg/gomp/target-this-3.C: Likewise.
	* g++.dg/gomp/target-this-4.C: Likewise.
---
 gcc/gimplify.cc                             | 785 +++++++++++++++++++-
 gcc/omp-low.cc                              |   7 +-
 gcc/testsuite/g++.dg/gomp/target-lambda-1.C |   7 +-
 gcc/testsuite/g++.dg/gomp/target-this-3.C   |   4 +-
 gcc/testsuite/g++.dg/gomp/target-this-4.C   |   4 +-
 5 files changed, 793 insertions(+), 14 deletions(-)

Comments

Jakub Jelinek May 24, 2022, 1:03 p.m. UTC | #1
On Fri, Mar 18, 2022 at 09:24:51AM -0700, Julian Brown wrote:
> 2021-11-23  Julian Brown  <julian@codesourcery.com>
> 
> gcc/
> 	* gimplify.c (is_or_contains_p, omp_target_reorder_clauses): Delete
> 	functions.
> 	(omp_tsort_mark): Add enum.
> 	(omp_mapping_group): Add struct.
> 	(debug_mapping_group, omp_get_base_pointer, omp_get_attachment,
> 	omp_group_last, omp_gather_mapping_groups, omp_group_base,
> 	omp_index_mapping_groups, omp_containing_struct,
> 	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, unless we've seen an error.
> 	* omp-low.c (scan_sharing_clauses): Avoid strict test if we haven't
> 	sorted mapping groups.
> 
> gcc/testsuite/
> 	* g++.dg/gomp/target-lambda-1.C: Adjust expected output.
> 	* g++.dg/gomp/target-this-3.C: Likewise.
> 	* g++.dg/gomp/target-this-4.C: Likewise.
> +

Wouldn't hurt to add a comment on the meanings of the enumerators.

> +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;
> +};
> +
> +__attribute__((used)) static void

I'd use what is used elsewhere,
DEBUG_FUNCTION void
without static.

> +debug_mapping_group (omp_mapping_group *grp)
> +{
> +  tree tmp = OMP_CLAUSE_CHAIN (grp->grp_end);
> +  OMP_CLAUSE_CHAIN (grp->grp_end) = NULL;
> +  debug_generic_expr (*grp->grp_start);
> +  OMP_CLAUSE_CHAIN (grp->grp_end) = tmp;
> +}
> +
> +/* 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) == ARRAY_REF)
> +    expr = TREE_OPERAND (expr, 0);
> +
> +  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)))
> +	     || TREE_CODE (TREE_OPERAND (expr, 0)) == ARRAY_REF))
> +    {
> +      expr = TREE_OPERAND (expr, 0);
> +
> +      while (TREE_CODE (expr) == ARRAY_REF)
> +	expr = TREE_OPERAND (expr, 0);
> +
> +      if (TREE_CODE (expr) == INDIRECT_REF || TREE_CODE (expr) == MEM_REF)
> +	break;
> +    }

I must say I don't see advantages of just a single loop that
looks through all ARRAY_REFs and all COMPONENT_REFs and then just
checks if the expr it got in the end is a decl or INDIRECT_REF
or MEM_REF with offset 0.

> +  if (DECL_P (expr))
> +    return NULL_TREE;
> +
> +  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;
> +}
> +

> +static tree
> +omp_containing_struct (tree expr)
> +{
> +  tree expr0 = expr;
> +
> +  STRIP_NOPS (expr);
> +
> +  tree expr1 = expr;
> +
> +  /* FIXME: other types of accessors?  */
> +  while (TREE_CODE (expr) == ARRAY_REF)
> +    expr = TREE_OPERAND (expr, 0);
> +
> +  if (TREE_CODE (expr) == COMPONENT_REF)
> +    {
> +      if (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)))
> +	  || TREE_CODE (TREE_OPERAND (expr, 0)) == ARRAY_REF)
> +	expr = TREE_OPERAND (expr, 0);
> +      else
> +	internal_error ("unhandled component");
> +    }

Again?

> @@ -9063,11 +9820,29 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
>  	break;
>        }
>  
> -  if (code == OMP_TARGET
> -      || code == OMP_TARGET_DATA
> -      || code == OMP_TARGET_ENTER_DATA
> -      || code == OMP_TARGET_EXIT_DATA)
> -    omp_target_reorder_clauses (list_p);
> +  /* Topological sorting may fail if we have duplicate nodes, which
> +     we should have detected and shown an error for already.  Skip
> +     sorting in that case.  */
> +  if (!seen_error ()
> +      && (code == OMP_TARGET
> +	  || code == OMP_TARGET_DATA
> +	  || code == OMP_TARGET_ENTER_DATA
> +	  || code == OMP_TARGET_EXIT_DATA))
> +    {
> +      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;
> +	}
> +    }

I think big question is if we do want to do this map clause reordering
before processing the  omp target etc. clauses, or after (during
gimplify_adjust_omp_clauses, when clauses from the implicit mappings
are added too and especially with the declare mapper expansions),
or both before and after.

>    while ((c = *list_p) != NULL)
>      {
> diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc
> index c33b3daa439..ffeb1f34fd7 100644
> --- a/gcc/omp-low.cc
> +++ b/gcc/omp-low.cc
> @@ -1537,8 +1537,11 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
>  	    {
>  	      /* If this is an offloaded region, an attach operation should
>  		 only exist when the pointer variable is mapped in a prior
> -		 clause.  */
> -	      if (is_gimple_omp_offloaded (ctx->stmt))
> +		 clause.
> +		 If we had an error, we may not have attempted to sort clauses
> +		 properly, so avoid the test.  */
> +	      if (is_gimple_omp_offloaded (ctx->stmt)
> +		  && !seen_error ())

If we encounter a major error during processing map clauses, we should consider
just leaving out the offloading construct from the IL.

	Jakub
Julian Brown June 8, 2022, 3 p.m. UTC | #2
Hi Jakub,

Thanks for review!

On Tue, 24 May 2022 15:03:07 +0200
Jakub Jelinek via Fortran <fortran@gcc.gnu.org> wrote:

> On Fri, Mar 18, 2022 at 09:24:51AM -0700, Julian Brown wrote:
> > 2021-11-23  Julian Brown  <julian@codesourcery.com>
> > 
> > gcc/
> > 	* gimplify.c (is_or_contains_p,
> > omp_target_reorder_clauses): Delete functions.
> > 	(omp_tsort_mark): Add enum.
> > 	(omp_mapping_group): Add struct.
> > 	(debug_mapping_group, omp_get_base_pointer,
> > omp_get_attachment, omp_group_last, omp_gather_mapping_groups,
> > omp_group_base, omp_index_mapping_groups, omp_containing_struct,
> > 	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, unless we've seen an error.
> > 	* omp-low.c (scan_sharing_clauses): Avoid strict test if we
> > haven't sorted mapping groups.
> > 
> > gcc/testsuite/
> > 	* g++.dg/gomp/target-lambda-1.C: Adjust expected output.
> > 	* g++.dg/gomp/target-this-3.C: Likewise.
> > 	* g++.dg/gomp/target-this-4.C: Likewise.
> > +  
> 
> Wouldn't hurt to add a comment on the meanings of the enumerators.

Added.

> > +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;
> > +};
> > +
> > +__attribute__((used)) static void  
> 
> I'd use what is used elsewhere,
> DEBUG_FUNCTION void
> without static.

Fixed.

> > +static tree
> > +omp_get_base_pointer (tree expr)

> I must say I don't see advantages of just a single loop that
> looks through all ARRAY_REFs and all COMPONENT_REFs and then just
> checks if the expr it got in the end is a decl or INDIRECT_REF
> or MEM_REF with offset 0.
> 
> > +static tree
> > +omp_containing_struct (tree expr)
> Again?

I've simplified these loops, and removed the "needs improvement"
comment.

> > @@ -9063,11 +9820,29 @@ gimplify_scan_omp_clauses (tree *list_p,
> > gimple_seq *pre_p, break;
> >        }
> >  
> > -  if (code == OMP_TARGET
> > -      || code == OMP_TARGET_DATA
> > -      || code == OMP_TARGET_ENTER_DATA
> > -      || code == OMP_TARGET_EXIT_DATA)
> > -    omp_target_reorder_clauses (list_p);
> > +  /* Topological sorting may fail if we have duplicate nodes, which
> > +     we should have detected and shown an error for already.  Skip
> > +     sorting in that case.  */
> > +  if (!seen_error ()
> > +      && (code == OMP_TARGET
> > +	  || code == OMP_TARGET_DATA
> > +	  || code == OMP_TARGET_ENTER_DATA
> > +	  || code == OMP_TARGET_EXIT_DATA))
> > +    {
> > +      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;
> > +	}
> > +    }  
> 
> I think big question is if we do want to do this map clause reordering
> before processing the  omp target etc. clauses, or after (during
> gimplify_adjust_omp_clauses, when clauses from the implicit mappings
> are added too and especially with the declare mapper expansions),
> or both before and after.

The existing code constrains us a bit here, unless we want to
completely rewrite it!

We can only do sorting on clauses before gimplification, otherwise the
"structural" matching of the parsed syntax of base pointers inside other
clauses on the directive, etc. will certainly fail.

(Semi-relatedly, I asked this on the omp-lang mailing list:

  "When we have mappings that represent base pointers, and other
  mappings that use those base pointers, the former must be ordered to
  take place before the latter -- but should we determine that relation
  purely syntactically? How about if we write e.g. "p->" on one vs.
  "(*p)." on the other?"

but no reply...)

So, this is fine for sorting explicit mapping clauses. When planning
the approach I've used for "declare mapper" support, I wrote this (in
an internal email):

"At the moment, gimplifying OMP workshare regions proceeds in three
phases:

 1. Clauses are processed (gimplify_scan_omp_clauses), creating
    records of mapped variables in a splay tree, with associated flags.

 2. The body of the workshare region is processed (gimplified),
    augmenting the same splay tree with information about variables
    which are used implicitly (and maybe also modifying the "explicit"
    mappings from the first step).

 3. The clauses are modified based on the results of the second stage
    (gimplify_adjust_omp_clauses). E.g. clauses are removed that refer
    to variables that aren't actually used in the region, or new
    clauses created for implicitly-referenced variables without mapping
    clauses on the construct.

The problem with this with regards to mappers is that the "expanded"
mappers should undergo some of the processing we currently perform
during phase 1 (struct sibling list handling, and so on), but we don't
know which variables are implicitly referenced until phase 2.

[description of a plan that didn't work removed]

So the new plan is to do:

phase 1  (scan original clauses)
phase 2  (scan workshare body)
phase 1  (use variables from "2" to instantiate mappers, and process
          new clauses only. Prepend new list to original clauses)
phase 3  (as before)

I was concerned that this would upset the sorting code -- but I think
actually, as long as implicitly-created clauses are inserted at the
front of the clause list, there can't be a case where a pointer base is
mapped after a use of that base. If that assumption turns out to be
wrong, then things might get a little more complicated."

...and so far, the plan seems to be working out. The assumption, to
state it in other words, is that an implicitly-added map clause *cannot*
have a dependency on an explicit map clause, in terms of relying on a
base pointer in that explicit clause, by construction.

> >    while ((c = *list_p) != NULL)
> >      {
> > diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc
> > index c33b3daa439..ffeb1f34fd7 100644
> > --- a/gcc/omp-low.cc
> > +++ b/gcc/omp-low.cc
> > @@ -1537,8 +1537,11 @@ scan_sharing_clauses (tree clauses,
> > omp_context *ctx) {
> >  	      /* If this is an offloaded region, an attach
> > operation should only exist when the pointer variable is mapped in
> > a prior
> > -		 clause.  */
> > -	      if (is_gimple_omp_offloaded (ctx->stmt))
> > +		 clause.
> > +		 If we had an error, we may not have attempted to
> > sort clauses
> > +		 properly, so avoid the test.  */
> > +	      if (is_gimple_omp_offloaded (ctx->stmt)
> > +		  && !seen_error ())  
> 
> If we encounter a major error during processing map clauses, we
> should consider just leaving out the offloading construct from the IL.

I experimented with that idea, but I think it may be a much more
invasive change (I saw lots of testsuite fall-out relating to things
that no longer raise "cascaded" errors, though maybe the approach I
took was too crude). I think if we want to do that, it's probably better
handled with a separate patch.

I've re-tested the attached with offloading to NVPTX. OK?

Thanks,

Julian
Jakub Jelinek June 9, 2022, 2:45 p.m. UTC | #3
On Wed, Jun 08, 2022 at 04:00:39PM +0100, Julian Brown wrote:
> > I think big question is if we do want to do this map clause reordering
> > before processing the  omp target etc. clauses, or after (during
> > gimplify_adjust_omp_clauses, when clauses from the implicit mappings
> > are added too and especially with the declare mapper expansions),
> > or both before and after.
> 
> The existing code constrains us a bit here, unless we want to
> completely rewrite it!
> 
> We can only do sorting on clauses before gimplification, otherwise the
> "structural" matching of the parsed syntax of base pointers inside other
> clauses on the directive, etc. will certainly fail.
> 
> (Semi-relatedly, I asked this on the omp-lang mailing list:
> 
>   "When we have mappings that represent base pointers, and other
>   mappings that use those base pointers, the former must be ordered to
>   take place before the latter -- but should we determine that relation
>   purely syntactically? How about if we write e.g. "p->" on one vs.
>   "(*p)." on the other?"
> 
> but no reply...)
> 
> So, this is fine for sorting explicit mapping clauses. When planning
> the approach I've used for "declare mapper" support, I wrote this (in
> an internal email):
> 
> "At the moment, gimplifying OMP workshare regions proceeds in three
> phases:
> 
>  1. Clauses are processed (gimplify_scan_omp_clauses), creating
>     records of mapped variables in a splay tree, with associated flags.
> 
>  2. The body of the workshare region is processed (gimplified),
>     augmenting the same splay tree with information about variables
>     which are used implicitly (and maybe also modifying the "explicit"
>     mappings from the first step).
> 
>  3. The clauses are modified based on the results of the second stage
>     (gimplify_adjust_omp_clauses). E.g. clauses are removed that refer
>     to variables that aren't actually used in the region, or new
>     clauses created for implicitly-referenced variables without mapping
>     clauses on the construct.
> 
> The problem with this with regards to mappers is that the "expanded"
> mappers should undergo some of the processing we currently perform
> during phase 1 (struct sibling list handling, and so on), but we don't
> know which variables are implicitly referenced until phase 2.
> 
> [description of a plan that didn't work removed]
> 
> So the new plan is to do:
> 
> phase 1  (scan original clauses)
> phase 2  (scan workshare body)
> phase 1  (use variables from "2" to instantiate mappers, and process
>           new clauses only. Prepend new list to original clauses)
> phase 3  (as before)
> 
> I was concerned that this would upset the sorting code -- but I think
> actually, as long as implicitly-created clauses are inserted at the
> front of the clause list, there can't be a case where a pointer base is
> mapped after a use of that base. If that assumption turns out to be
> wrong, then things might get a little more complicated."
> 
> ...and so far, the plan seems to be working out. The assumption, to
> state it in other words, is that an implicitly-added map clause *cannot*
> have a dependency on an explicit map clause, in terms of relying on a
> base pointer in that explicit clause, by construction.

I don't think there is any need to add extra phases, but we can move
some code from gimplify_scan_omp_clauses to gimplify_adjust_omp_clauses.
What must be done in gimplify_scan_omp_clauses is stuff that will or
could affect the gimplification of the region's body, in that phase 2
we want to know say that some variable was privatized explicitly or
explicitly mapped or none of that, so we can based on that decide if we
should note implicit data sharing or implicit mapping etc.
But e.g. the sorting of the OMP_CLAUSE_MAP clauses is something that can
IMHO be deferred until we have all those clauses, probably it is done
in gimplify_scan_omp_clauses right now was just that the sorting at least
initially was only needed for struct mapping (map (tofrom: a.b, a.c, a.d.e, a.d.f))
and that could appear only explicitly, not implicitly, implicit mapping
would only map the whole var.
But declare mapper changes this substantially, declare mapper can add
similar mappings even from the implicit maps.
So, I think we should keep in phase 1 for OMP_CLAUSE_MAP only the stuff that
perhaps gimplifies some expressions used in those and puts records about
them into splay trees and sorting and ideally some kind of merging of
adjacent mappings can be done only when we have even the implicit
mappings all collected (so that would be after
  splay_tree_foreach (ctx->variables, gimplify_adjust_omp_clauses_1, &data);
finishes).

	Jakub
diff mbox series

Patch

diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc
index 875b115d02d..968cbd263f5 100644
--- a/gcc/gimplify.cc
+++ b/gcc/gimplify.cc
@@ -8738,6 +8738,7 @@  extract_base_bit_offset (tree base, tree *base_ref, poly_int64 *bitposp,
   return base;
 }
 
+#if 0
 /* Returns true if EXPR is or contains (as a sub-component) BASE_PTR.  */
 
 static bool
@@ -8761,6 +8762,7 @@  is_or_contains_p (tree expr, tree base_ptr)
   return operand_equal_p (expr, base_ptr);
 }
 
+
 /* 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.  */
@@ -8940,6 +8942,761 @@  omp_target_reorder_clauses (tree *list_p)
 	    }
       }
 }
+#endif
+
+
+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;
+};
+
+__attribute__((used)) static void
+debug_mapping_group (omp_mapping_group *grp)
+{
+  tree tmp = OMP_CLAUSE_CHAIN (grp->grp_end);
+  OMP_CLAUSE_CHAIN (grp->grp_end) = NULL;
+  debug_generic_expr (*grp->grp_start);
+  OMP_CLAUSE_CHAIN (grp->grp_end) = tmp;
+}
+
+/* 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) == ARRAY_REF)
+    expr = TREE_OPERAND (expr, 0);
+
+  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)))
+	     || TREE_CODE (TREE_OPERAND (expr, 0)) == ARRAY_REF))
+    {
+      expr = TREE_OPERAND (expr, 0);
+
+      while (TREE_CODE (expr) == ARRAY_REF)
+	expr = TREE_OPERAND (expr, 0);
+
+      if (TREE_CODE (expr) == INDIRECT_REF || TREE_CODE (expr) == MEM_REF)
+	break;
+    }
+
+  if (DECL_P (expr))
+    return NULL_TREE;
+
+  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;
+}
+
+/* An attach or detach operation depends directly on the address being
+   attached/detached.  Return that address, or none if there are no
+   attachments/detachments.  */
+
+static tree
+omp_get_attachment (omp_mapping_group *grp)
+{
+  tree node = *grp->grp_start;
+
+  switch (OMP_CLAUSE_MAP_KIND (node))
+    {
+    case GOMP_MAP_TO:
+    case GOMP_MAP_FROM:
+    case GOMP_MAP_TOFROM:
+    case GOMP_MAP_ALWAYS_FROM:
+    case GOMP_MAP_ALWAYS_TO:
+    case GOMP_MAP_ALWAYS_TOFROM:
+    case GOMP_MAP_FORCE_FROM:
+    case GOMP_MAP_FORCE_TO:
+    case GOMP_MAP_FORCE_TOFROM:
+    case GOMP_MAP_FORCE_PRESENT:
+    case GOMP_MAP_ALLOC:
+    case GOMP_MAP_RELEASE:
+    case GOMP_MAP_DELETE:
+    case GOMP_MAP_FORCE_ALLOC:
+      if (node == grp->grp_end)
+	return NULL_TREE;
+
+      node = OMP_CLAUSE_CHAIN (node);
+      if (node && OMP_CLAUSE_MAP_KIND (node) == GOMP_MAP_TO_PSET)
+	{
+	  gcc_assert (node != grp->grp_end);
+	  node = OMP_CLAUSE_CHAIN (node);
+	}
+      if (node)
+	switch (OMP_CLAUSE_MAP_KIND (node))
+	  {
+	  case GOMP_MAP_POINTER:
+	  case GOMP_MAP_ALWAYS_POINTER:
+	  case GOMP_MAP_FIRSTPRIVATE_POINTER:
+	  case GOMP_MAP_FIRSTPRIVATE_REFERENCE:
+	  case GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION:
+	    return NULL_TREE;
+
+	  case GOMP_MAP_ATTACH_DETACH:
+	  case GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION:
+	    return OMP_CLAUSE_DECL (node);
+
+	  default:
+	    internal_error ("unexpected mapping node");
+	  }
+      return error_mark_node;
+
+    case GOMP_MAP_TO_PSET:
+      gcc_assert (node != grp->grp_end);
+      node = OMP_CLAUSE_CHAIN (node);
+      if (OMP_CLAUSE_MAP_KIND (node) == GOMP_MAP_ATTACH
+	  || OMP_CLAUSE_MAP_KIND (node) == GOMP_MAP_DETACH)
+	return OMP_CLAUSE_DECL (node);
+      else
+	internal_error ("unexpected mapping node");
+      return error_mark_node;
+
+    case GOMP_MAP_ATTACH:
+    case GOMP_MAP_DETACH:
+      node = OMP_CLAUSE_CHAIN (node);
+      if (!node || *grp->grp_start == grp->grp_end)
+	return OMP_CLAUSE_DECL (*grp->grp_start);
+      if (OMP_CLAUSE_MAP_KIND (node) == GOMP_MAP_FIRSTPRIVATE_POINTER
+	  || OMP_CLAUSE_MAP_KIND (node) == GOMP_MAP_FIRSTPRIVATE_REFERENCE)
+	return OMP_CLAUSE_DECL (*grp->grp_start);
+      else
+	internal_error ("unexpected mapping node");
+      return error_mark_node;
+
+    case GOMP_MAP_STRUCT:
+    case GOMP_MAP_FORCE_DEVICEPTR:
+    case GOMP_MAP_DEVICE_RESIDENT:
+    case GOMP_MAP_LINK:
+    case GOMP_MAP_IF_PRESENT:
+    case GOMP_MAP_FIRSTPRIVATE:
+    case GOMP_MAP_FIRSTPRIVATE_INT:
+    case GOMP_MAP_USE_DEVICE_PTR:
+    case GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION:
+      return NULL_TREE;
+
+    default:
+      internal_error ("unexpected mapping node");
+    }
+
+  return error_mark_node;
+}
+
+/* Given a pointer START_P to the start of a group of related (e.g. pointer)
+   mappings, return the chain pointer to the end of that group in the list.  */
+
+static tree *
+omp_group_last (tree *start_p)
+{
+  tree c = *start_p, nc, *grp_last_p = start_p;
+
+  gcc_assert (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP);
+
+  nc = OMP_CLAUSE_CHAIN (c);
+
+  if (!nc || OMP_CLAUSE_CODE (nc) != OMP_CLAUSE_MAP)
+    return grp_last_p;
+
+  switch (OMP_CLAUSE_MAP_KIND (c))
+    {
+    default:
+      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_ATTACH_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;
+
+    case GOMP_MAP_ATTACH:
+    case GOMP_MAP_DETACH:
+      /* This is a weird artifact of how directives are parsed: bare attach or
+	 detach clauses get a subsequent (meaningless) FIRSTPRIVATE_POINTER or
+	 FIRSTPRIVATE_REFERENCE node.  FIXME.  */
+      if (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))
+	grp_last_p = &OMP_CLAUSE_CHAIN (c);
+      break;
+
+    case GOMP_MAP_TO_PSET:
+      if (OMP_CLAUSE_CODE (nc) == OMP_CLAUSE_MAP
+	  && (OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_ATTACH
+	      || OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_DETACH))
+	grp_last_p = &OMP_CLAUSE_CHAIN (c);
+      break;
+    }
+
+  return grp_last_p;
+}
+
+/* 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))
+    {
+      if (OMP_CLAUSE_CODE (*cp) != OMP_CLAUSE_MAP)
+	continue;
+
+      tree *grp_last_p = omp_group_last (cp);
+      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 () > 0)
+    return groups;
+  else
+    {
+      delete groups;
+      return NULL;
+    }
+}
+
+/* A pointer mapping group GRP may define a block of memory starting at some
+   base address, and maybe also define a firstprivate pointer or firstprivate
+   reference that points to that block.  The return value is a node containing
+   the former, and the *FIRSTPRIVATE pointer is set if we have the latter.
+   If we define several base pointers, i.e. for a GOMP_MAP_STRUCT mapping,
+   return the number of consecutive chained nodes in CHAINED.  */
+
+static tree
+omp_group_base (omp_mapping_group *grp, unsigned int *chained,
+		tree *firstprivate)
+{
+  tree node = *grp->grp_start;
+
+  *firstprivate = NULL_TREE;
+  *chained = 1;
+
+  switch (OMP_CLAUSE_MAP_KIND (node))
+    {
+    case GOMP_MAP_TO:
+    case GOMP_MAP_FROM:
+    case GOMP_MAP_TOFROM:
+    case GOMP_MAP_ALWAYS_FROM:
+    case GOMP_MAP_ALWAYS_TO:
+    case GOMP_MAP_ALWAYS_TOFROM:
+    case GOMP_MAP_FORCE_FROM:
+    case GOMP_MAP_FORCE_TO:
+    case GOMP_MAP_FORCE_TOFROM:
+    case GOMP_MAP_FORCE_PRESENT:
+    case GOMP_MAP_ALLOC:
+    case GOMP_MAP_RELEASE:
+    case GOMP_MAP_DELETE:
+    case GOMP_MAP_FORCE_ALLOC:
+      if (node == grp->grp_end)
+	return node;
+
+      node = OMP_CLAUSE_CHAIN (node);
+      if (node && OMP_CLAUSE_MAP_KIND (node) == GOMP_MAP_TO_PSET)
+	{
+	  gcc_assert (node != grp->grp_end);
+	  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;
+
+	  case GOMP_MAP_ALWAYS_POINTER:
+	  case GOMP_MAP_ATTACH_DETACH:
+	  case GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION:
+	    return *grp->grp_start;
+
+	  default:
+	    internal_error ("unexpected mapping node");
+	  }
+      else
+	internal_error ("unexpected mapping node");
+      return error_mark_node;
+
+    case GOMP_MAP_TO_PSET:
+      gcc_assert (node != grp->grp_end);
+      node = OMP_CLAUSE_CHAIN (node);
+      if (OMP_CLAUSE_MAP_KIND (node) == GOMP_MAP_ATTACH
+	  || OMP_CLAUSE_MAP_KIND (node) == GOMP_MAP_DETACH)
+	return NULL_TREE;
+      else
+	internal_error ("unexpected mapping node");
+      return error_mark_node;
+
+    case GOMP_MAP_ATTACH:
+    case GOMP_MAP_DETACH:
+      node = OMP_CLAUSE_CHAIN (node);
+      if (!node || *grp->grp_start == grp->grp_end)
+	return NULL_TREE;
+      if (OMP_CLAUSE_MAP_KIND (node) == GOMP_MAP_FIRSTPRIVATE_POINTER
+	  || OMP_CLAUSE_MAP_KIND (node) == GOMP_MAP_FIRSTPRIVATE_REFERENCE)
+	{
+	  /* We're mapping the base pointer itself in a bare attach or detach
+	     node.  This is a side effect of how parsing works, and the mapping
+	     will be removed anyway (at least for enter/exit data directives).
+	     We should ignore the mapping here.  FIXME.  */
+	  return NULL_TREE;
+	}
+      else
+	internal_error ("unexpected mapping node");
+      return error_mark_node;
+
+    case GOMP_MAP_FORCE_DEVICEPTR:
+    case GOMP_MAP_DEVICE_RESIDENT:
+    case GOMP_MAP_LINK:
+    case GOMP_MAP_IF_PRESENT:
+    case GOMP_MAP_FIRSTPRIVATE:
+    case GOMP_MAP_FIRSTPRIVATE_INT:
+    case GOMP_MAP_USE_DEVICE_PTR:
+    case GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION:
+      return NULL_TREE;
+
+    case GOMP_MAP_FIRSTPRIVATE_POINTER:
+    case GOMP_MAP_FIRSTPRIVATE_REFERENCE:
+    case GOMP_MAP_POINTER:
+    case GOMP_MAP_ALWAYS_POINTER:
+    case GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION:
+      /* These shouldn't appear by themselves.  */
+      if (!seen_error ())
+	internal_error ("unexpected pointer mapping node");
+      return error_mark_node;
+
+    default:
+      gcc_unreachable ();
+    }
+
+  return error_mark_node;
+}
+
+/* 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 fpp;
+      unsigned int chained;
+      tree node = omp_group_base (grp, &chained, &fpp);
+
+      if (node == error_mark_node || (!node && !fpp))
+	continue;
+
+      for (unsigned j = 0;
+	   node && j < chained;
+	   node = OMP_CLAUSE_CHAIN (node), j++)
+	{
+	  tree decl = OMP_CLAUSE_DECL (node);
+
+	  /* 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 && *prev == grp)
+	    /* Empty.  */;
+	  else 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);
+	}
+
+      if (!fpp)
+	continue;
+
+      omp_mapping_group **prev = grpmap->get (fpp);
+      if (prev)
+	{
+	  grp->sibling = (*prev)->sibling;
+	  (*prev)->sibling = grp;
+	}
+      else
+	grpmap->put (fpp, grp);
+    }
+  return grpmap;
+}
+
+/* Find the immediately-containing struct for a component ref (etc.)
+   expression EXPR.  */
+
+static tree
+omp_containing_struct (tree expr)
+{
+  tree expr0 = expr;
+
+  STRIP_NOPS (expr);
+
+  tree expr1 = expr;
+
+  /* FIXME: other types of accessors?  */
+  while (TREE_CODE (expr) == ARRAY_REF)
+    expr = TREE_OPERAND (expr, 0);
+
+  if (TREE_CODE (expr) == COMPONENT_REF)
+    {
+      if (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)))
+	  || TREE_CODE (TREE_OPERAND (expr, 0)) == ARRAY_REF)
+	expr = TREE_OPERAND (expr, 0);
+      else
+	internal_error ("unhandled component");
+    }
+
+  return (expr == expr1) ? expr0 : expr;
+}
+
+/* Helper function for omp_tsort_mapping_groups.  Returns TRUE on success, or
+   FALSE on error.  */
+
+static bool
+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)
+{
+  if (grp->mark == PERMANENT)
+    return true;
+  if (grp->mark == TEMPORARY)
+    {
+      fprintf (stderr, "when processing group:\n");
+      debug_mapping_group (grp);
+      internal_error ("base pointer cycle detected");
+      return false;
+    }
+  grp->mark = TEMPORARY;
+
+  tree attaches_to = omp_get_attachment (grp);
+
+  if (attaches_to)
+    {
+      omp_mapping_group **basep = grpmap->get (attaches_to);
+
+      if (basep)
+	{
+	  gcc_assert (*basep != grp);
+	  for (omp_mapping_group *w = *basep; w; w = w->sibling)
+	    if (!omp_tsort_mapping_groups_1 (outlist, groups, grpmap, w))
+	      return false;
+	}
+    }
+
+  tree decl = OMP_CLAUSE_DECL (*grp->grp_start);
+
+  while (decl)
+    {
+      tree base = omp_get_base_pointer (decl);
+
+      if (!base)
+	break;
+
+      omp_mapping_group **innerp = grpmap->get (base);
+
+      /* 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)
+	{
+	  base = omp_containing_struct (base);
+	  innerp = grpmap->get (base);
+
+	  if (!innerp
+	      && TREE_CODE (base) == MEM_REF
+	      && integer_zerop (TREE_OPERAND (base, 1)))
+	    {
+	      tree ind = TREE_OPERAND (base, 0);
+	      ind = build1 (INDIRECT_REF, TREE_TYPE (base), ind);
+	      innerp = grpmap->get (ind);
+	    }
+	}
+
+      if (innerp && *innerp != grp)
+	{
+	  for (omp_mapping_group *w = *innerp; w; w = w->sibling)
+	    if (!omp_tsort_mapping_groups_1 (outlist, groups, grpmap, w))
+	      return false;
+	  break;
+	}
+
+      decl = base;
+    }
+
+  grp->mark = PERMANENT;
+
+  /* Emit grp to output list.  */
+
+  **outlist = grp;
+  *outlist = &grp->next;
+
+  return true;
+}
+
+/* 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)
+	if (!omp_tsort_mapping_groups_1 (&cursor, groups, grpmap, grp))
+	  return NULL;
+    }
+
+  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 *> old_headps (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_headps.quick_push (grp->grp_start);
+      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];
+	  *old_headps[i] = new_heads[i];
+	}
+      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
    of combined/composite constructs, starting with OCTX.
@@ -9063,11 +9820,29 @@  gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 	break;
       }
 
-  if (code == OMP_TARGET
-      || code == OMP_TARGET_DATA
-      || code == OMP_TARGET_ENTER_DATA
-      || code == OMP_TARGET_EXIT_DATA)
-    omp_target_reorder_clauses (list_p);
+  /* Topological sorting may fail if we have duplicate nodes, which
+     we should have detected and shown an error for already.  Skip
+     sorting in that case.  */
+  if (!seen_error ()
+      && (code == OMP_TARGET
+	  || code == OMP_TARGET_DATA
+	  || code == OMP_TARGET_ENTER_DATA
+	  || code == OMP_TARGET_EXIT_DATA))
+    {
+      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/omp-low.cc b/gcc/omp-low.cc
index c33b3daa439..ffeb1f34fd7 100644
--- a/gcc/omp-low.cc
+++ b/gcc/omp-low.cc
@@ -1537,8 +1537,11 @@  scan_sharing_clauses (tree clauses, omp_context *ctx)
 	    {
 	      /* If this is an offloaded region, an attach operation should
 		 only exist when the pointer variable is mapped in a prior
-		 clause.  */
-	      if (is_gimple_omp_offloaded (ctx->stmt))
+		 clause.
+		 If we had an error, we may not have attempted to sort clauses
+		 properly, so avoid the test.  */
+	      if (is_gimple_omp_offloaded (ctx->stmt)
+		  && !seen_error ())
 		gcc_assert
 		  (maybe_lookup_decl (decl, ctx)
 		   || (is_global_var (maybe_lookup_decl_in_outer_ctx (decl, ctx))
diff --git a/gcc/testsuite/g++.dg/gomp/target-lambda-1.C b/gcc/testsuite/g++.dg/gomp/target-lambda-1.C
index 7f83f92ec93..279dab1d8e8 100644
--- a/gcc/testsuite/g++.dg/gomp/target-lambda-1.C
+++ b/gcc/testsuite/g++.dg/gomp/target-lambda-1.C
@@ -87,8 +87,9 @@  int main (void)
   return 0;
 }
 
-/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(b\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:iptr \[pointer assign, bias: 0\]\) 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\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\) map\(always_pointer:__closure->__this \[pointer assign, bias: 0\]\) map\(attach_zero_length_array_section:__closure->__iptr \[bias: 0\]\) map\(attach_zero_length_array_section:_[0-9]+->ptr \[bias: 0\]\) map\(from:mapped \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\)} "gimple" } } */
+/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(b\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:iptr \[pointer assign, bias: 0\]\) 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\(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\(attach_zero_length_array_section:__closure->__iptr \[bias: 0\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) map\(attach_zero_length_array_section:_[0-9]+->ptr \[bias: 0\]\) 
+} "gimple" } } */
 
-/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(end\) firstprivate\(begin\) map\(to:loop \[len: [0-9]+\]\) map\(attach_zero_length_array_section:loop\.__data1 \[bias: 0\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\)} "gimple" } } */
+/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(end\) firstprivate\(begin\) map\(to:loop \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) map\(attach_zero_length_array_section:loop\.__data1 \[bias: 0\]\)} "gimple" } } */
 
-/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(end\) firstprivate\(begin\) map\(to:loop \[len: [0-9]+\]\) map\(attach_zero_length_array_section:loop\.__data2 \[bias: 0\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\)} "gimple" } } */
+/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(end\) firstprivate\(begin\) map\(to:loop \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) map\(attach_zero_length_array_section:loop\.__data2 \[bias: 0\]\)} "gimple" } } */
diff --git a/gcc/testsuite/g++.dg/gomp/target-this-3.C b/gcc/testsuite/g++.dg/gomp/target-this-3.C
index 91cfbd6ef20..bc2cc0b297d 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.* map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\) 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]+\]\)} "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.* map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\) 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" } } */
+/* { 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\(attach_zero_length_array_section:this->ptr \[bias: 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 e4b2a71bbb4..9ade3cc0b2b 100644
--- a/gcc/testsuite/g++.dg/gomp/target-this-4.C
+++ b/gcc/testsuite/g++.dg/gomp/target-this-4.C
@@ -102,6 +102,6 @@  int main (void)
   return 0;
 }
 
-/* { 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\(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: [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\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\) map\(always_pointer:__closure->__this \[pointer assign, bias: 0\]\) map\(from:mapped \[len: 1\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) map\(attach_zero_length_array_section:_[0-9]+->ptr \[bias: 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\(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\(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" } } */