[v3,05/11] OpenMP: push attaches to end of clause list in "target" regions
Commit Message
This patch 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 operation, 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.
2022-09-13 Julian Brown <julian@codesourcery.com>
gcc/
* gimplify.cc (omp_segregate_mapping_groups): Update comment.
(omp_push_attaches_to_end): New function.
(gimplify_scan_omp_clauses): Use omp_push_attaches_to_end for offloaded
regions.
---
gcc/gimplify.cc | 66 ++++++++++++++++++++-
gcc/testsuite/g++.dg/gomp/target-lambda-1.C | 2 +-
2 files changed, 65 insertions(+), 3 deletions(-)
Comments
On Tue, Sep 13, 2022 at 02:03:15PM -0700, Julian Brown wrote:
> This patch 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 operation, 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.
>
> 2022-09-13 Julian Brown <julian@codesourcery.com>
>
> gcc/
> * gimplify.cc (omp_segregate_mapping_groups): Update comment.
> (omp_push_attaches_to_end): New function.
> (gimplify_scan_omp_clauses): Use omp_push_attaches_to_end for offloaded
> regions.
Shouldn't this be done at the end of gimplify_adjust_omp_clauses?
I mean, can't further attach clauses appear because of declare mapper
for implicitly mapped variables?
Other than that, it is yet another walk of the whole clause list, so would
be nicer if it could be done in an existing walk over the clauses or
at least have a flag whether there are any such clauses present and do it
only in that case.
If it could be done in the main gimplify_adjust_omp_clauses loop, nice,
if it can appear also during
splay_tree_foreach (ctx->variables, gimplify_adjust_omp_clauses_1, &data);
that isn't the case.
Jakub
On Wed, 14 Sep 2022 14:44:54 +0200
Jakub Jelinek <jakub@redhat.com> wrote:
> On Tue, Sep 13, 2022 at 02:03:15PM -0700, Julian Brown wrote:
> > This patch 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 operation, 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.
> >
> > 2022-09-13 Julian Brown <julian@codesourcery.com>
> >
> > gcc/
> > * gimplify.cc (omp_segregate_mapping_groups): Update
> > comment. (omp_push_attaches_to_end): New function.
> > (gimplify_scan_omp_clauses): Use omp_push_attaches_to_end
> > for offloaded regions.
>
> Shouldn't this be done at the end of gimplify_adjust_omp_clauses?
> I mean, can't further attach clauses appear because of declare mapper
> for implicitly mapped variables?
> Other than that, it is yet another walk of the whole clause list, so
> would be nicer if it could be done in an existing walk over the
> clauses or at least have a flag whether there are any such clauses
> present and do it only in that case.
> If it could be done in the main gimplify_adjust_omp_clauses loop,
> nice, if it can appear also during
> splay_tree_foreach (ctx->variables, gimplify_adjust_omp_clauses_1,
> &data); that isn't the case.
I don't think any ATTACH clauses can appear during
the gimplify_adjust_omp_clause_1 walk. So, how about this?
Thanks,
Julian
On Sun, Sep 18, 2022 at 08:10:00PM +0100, Julian Brown wrote:
> I don't think any ATTACH clauses can appear during
> the gimplify_adjust_omp_clause_1 walk. So, how about this?
LGTM, thanks.
Jakub
@@ -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" } } */