Re: [PATCH v5 2/4] OpenMP/OpenACC: Rework clause expansion and nested struct handling

[Tobias Burnus <tobias@codesourcery.com>]

Hi Julian,

On 07.12.22 16:16, Julian Brown wrote:
> On Wed, 7 Dec 2022 15:54:42 +0100 Tobias Burnus <tobias@codesourcery.com> wrote:
>> If I understand Deepak's comment (on OpenMP.org's omp-lang list, sorry
>> it is a nonpublic list) correctly, the following wording implies that
>> a 'from: s.w[z:4]' for a pointer 's.w' also implies a mapping of
>> 's.w' - if 's' is used inside the target region and, thus, gets
>> implicitly mapped.
>>
>> [TR11 157:21-26] (approx. [5.2 154:22-27], [5.1 352:17-22], [5.0
>> 320:22-27])
>>
>> "If a list item with an implicit data-mapping attribute does not have
>> any corresponding storage in the device data environment prior to a
>> task encountering the construct associated with the map clause, and
>> one or more contiguous parts of the original storage are either list
>> items or base pointers to list items that are explicitly mapped on
>> the construct, only those parts of the original storage will have
>> corresponding storage in the device data environment as a result of
>> the map clauses on the construct."
> Hmmm... IIRC that is a different conclusion than the one we have
> understood previously, leading to e.g. the patch here (Chung-Lin CC'ed):
>
> https://gcc.gnu.org/pipermail/gcc-patches/2021-May/570075.html

This seems to be the "Target directive struct mapping question" omp-lang thread,
started on 2021-03-22.

I think we need to distinguish:

   #pragma omp target enter data map(to: s.w[:10])

from

   #pragma omp target map(tofrom: s.arr[:20])
     s.arr[0] = 5;

As in the latter case 's' gets implicitly mapped and then applies to
the base pointer 's.arr' of 's.arr[:20]'. While in the former case,
only the pointee gets mapped without the pointer 's.arr' (and, hence,
there is also no pointer attachment).

At least that's what I get from the wording above and reading Deepak's last
email - and it does not seem to clash with the discussion in the lengthy
omp-lang thread. (Maybe there are other threads – or I completely misread them.)

I think it makes sense to have a clarifying example in OpenMP; hence,
I filed the OpenMP.org example issue #342, starting with essentially
what I wrote above: 'target enter data' needs more work to get the pointer
handling done, 'target' + accessing 's' works as is.

I hope it makes sense.

> Follow-on discussion then questioned whether the change was really the
> intention of the spec, but we thought it was.  Has that changed now?

No idea – I find it difficult to track all the language changes and find
mapping complex and unclear.

However, it does seem to make sense in the way written above without
contradicting to all previous discussions, minus the common confusion.
(As least as I gathered from browsing both omp-lang and gcc-patches.)

> (I think actually changing the behaviour is a matter of flipping a
> switch, but let's make sure we choose the right setting!)

That sounds great!

Tobias

-----------------
Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634 München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht München, HRB 106955