[v3,05/11] OpenMP: push attaches to end of clause list in "target" regions

diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc
index c7998c2ccbd..bc7848843b3 100644
--- a/gcc/gimplify.cc
+++ b/gcc/gimplify.cc
@@ -9641,7 +9641,9 @@  omp_tsort_mapping_groups (vec<omp_mapping_group> *groups,
 /* 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.  */
+   order of the original list.
+   See also omp_push_attaches_to_end below -- we call that later after scanning
+   omp clauses.  */
 
 static omp_mapping_group *
 omp_segregate_mapping_groups (omp_mapping_group *inlist)
@@ -9681,6 +9683,55 @@  omp_segregate_mapping_groups (omp_mapping_group *inlist)
   return tf_groups;
 }
 
+/* This function moves GOMP_MAP_ATTACH{_ZERO_LENGTH_ARRAY_SECTION} nodes to the
+   end of the clause list, for offload regions.  This ensures that when we do
+   the attach, both the "attachment point" and the target region have both
+   already been mapped on the target.  This avoids a pathological case that can
+   otherwise happen with struct sibling-list handling.
+
+   Do not call this for non-offload regions, e.g. for "enter data" or
+   "exit data" directives.
+
+   The order of attach nodes and of non-attach nodes is otherwise retained.  */
+
+static tree
+omp_push_attaches_to_end (tree list)
+{
+  tree nonattach_list = NULL_TREE, attach_list = NULL_TREE;
+  tree *nonattach_tail = &nonattach_list, *attach_tail = &attach_list;
+
+  for (tree w = list; w;)
+    {
+      tree next = OMP_CLAUSE_CHAIN (w);
+
+      if (OMP_CLAUSE_CODE (w) != OMP_CLAUSE_MAP)
+	goto nonattach;
+
+      switch (OMP_CLAUSE_MAP_KIND (w))
+	{
+	case GOMP_MAP_ATTACH:
+	case GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION:
+	  *attach_tail = w;
+	  OMP_CLAUSE_CHAIN (w) = NULL_TREE;
+	  attach_tail = &OMP_CLAUSE_CHAIN (w);
+	  break;
+
+	default:
+	nonattach:
+	  *nonattach_tail = w;
+	  OMP_CLAUSE_CHAIN (w) = NULL_TREE;
+	  nonattach_tail = &OMP_CLAUSE_CHAIN (w);
+	}
+
+      w = next;
+    }
+
+  /* Splice lists together.  */
+  *nonattach_tail = attach_list;
+
+  return nonattach_list;
+}
+
 /* 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
@@ -11950,7 +12001,18 @@  gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 	list_p = &OMP_CLAUSE_CHAIN (c);
     }
 
-  ctx->clauses = *orig_list_p;
+  if ((region_type & ORT_TARGET) != 0)
+    /* If we have a target region, we can push all the attaches to the end of
+       the list (we may have standalone "attach" operations synthesized for
+       GOMP_MAP_STRUCT nodes that must be processed after the attachment point
+       AND the pointed-to block have been mapped).  */
+    ctx->clauses = omp_push_attaches_to_end (*orig_list_p);
+  else
+    /* ...but if we have something else, e.g. "enter data", we need to keep
+       "attach" nodes together with the previous node they attach to so that
+       separate "exit data" operations work properly (see libgomp/target.c).  */
+    ctx->clauses = *orig_list_p;
+
   gimplify_omp_ctxp = ctx;
 }
 
diff --git a/gcc/testsuite/g++.dg/gomp/target-lambda-1.C b/gcc/testsuite/g++.dg/gomp/target-lambda-1.C
index bff7fa7c669..5ce8ceadb19 100644
--- a/gcc/testsuite/g++.dg/gomp/target-lambda-1.C
+++ b/gcc/testsuite/g++.dg/gomp/target-lambda-1.C
@@ -87,7 +87,7 @@  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\(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\(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\(alloc:\*_[0-9]+ \[len: 0\]\) map\(attach_zero_length_array_section:__closure->__iptr \[bias: 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\(alloc:\*_[0-9]+ \[len: 0\]\) map\(attach_zero_length_array_section:loop\.__data1 \[bias: 0\]\)} "gimple" } } */