@@ -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;
}
@@ -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" } } */