https://gcc.gnu.org/bugzilla/show_bug.cgi?id=113724

--- Comment #4 from Tobias Burnus <burnus at gcc dot gnu.org> ---
Created attachment 57377
  --> https://gcc.gnu.org/bugzilla/attachment.cgi?id=57377&action=edit
Fixes the ICE – might paper over a real issue; doesn't fix the run-time issue →
TODO + 'data'-issue in PR comment 4

The patch fixes the issue

    #pragma omp target data map(S1.p[:N],S1.p,S1.a,S1.b)

This gets split into the groups (reverse order!) 'S1.b' (i = 0), 'S1.a' (i =
1), 'S1.p' (i = 3) and 'S1.p[:N]' (i = 4; map + attach).

In omp_build_struct_sibling_lists the collecting and reordering happens but for
'S1.p' there should be an 'alloc' not a 'tofrom' - such 'S1.p' has grp->deleted
and the attach code will add the 'alloc' code.

Until i = 2, everything is fine:
  *(*group[2])->deleted == true
  *(*group[2])->grp_start == (*group[2])->grp_end

and this encloses 'map(tofrom:S1.p [len: 8])' – which should be removed in
favor of a later (i = 3) added 'map(alloc:S1.p [len: 8])'.

In principle, everything looks fine, until i =3 calls
omp_accumulate_sibling_list, which in turn calls:

          continue_at
            = cl ? omp_siblist_move_concat_nodes_after (cl, tail_chain,
                                                        grp_start_p, grp_end,
                                                        sc)
                 : omp_siblist_move_nodes_after (grp_start_p, grp_end, sc);

where 'cl' != NULL_TREE.

After the call, 'tail_chain' alias 'list_p' looks fine - except for the tailing
'map(tofrom:S1.p [len: 8])'.

In principle, 'groups' is no longer touched - except for the the 'grp->deleted'
handling, which fails (deletes the wrong stuff) because grp_begin points to the
wrong tree.

Solution: Do the OMP_CLAUSE_DECL nullifying earlier such that messing around
with groups won't cause issues.


TODO: We should really find out WHY i=2's grp_begin gets updated. If it happens
just for previously processed grp items, that's fine - but what will happen if
it also affects a still to be processed item? - If that indeed happens,
everything will be messed up again!

* * *

The testcase shows another issue:

  target data map(to: S2.p[:N])

gets mapped as:

  map(struct:S2 [len: 1]) map(alloc:S2.p [len: 0])
       map(tofrom:*_14 [len: 400])    map(attach:S2.p [bias: 0])

before:

  map(tofrom:*_14 [len: 400]) map(attach:S2.p [bias: 0]

The problem of the former is of course that 'S' is already partially mapped and
that an alloc of length 0 will then fail already in 'target data' for the
attach:S2.p as 0 bytes aren't sufficient for a pointer attachment.

This applies both to target data and target, except that for 'target', 'S'
might appear implicitly - while for 'data' it can only appear explicitly or not
at all.

Reply via email to