Re: [PATCH, OG10, OpenMP 5.0, committed] Implement relaxation of implicit map vs. existing device mappings

[Chung-Lin Tang <cltang@codesourcery.com>]

On 2021/5/7 8:35 PM, Thomas Schwinge wrote:
> On 2021-05-05T23:17:25+0800, Chung-Lin Tang via Gcc-patches<gcc-patches@gcc.gnu.org>  wrote:
>> This patch implements relaxing the requirements when a map with the implicit attribute encounters
>> an overlapping existing map.  [...]
> Oh, oh, these data mapping interfaces/semantics ares getting more and
> more "convoluted"...  %-\ (Not your fault, of course.)
> 
> Haven't looked in too much detail in the patch/implementation (I'm not
> very well-versend in the exact OpenMP semantics anyway), but I suppose we
> should do similar things for OpenACC, too.  I think we even currently do
> have a gimplification-level "hack" to replicate data clauses' array
> bounds for implicit data clauses on compute constructs, if the default
> "complete" mapping is going to clash with a "limited" mapping that's
> specified in an outer OpenACC 'data' directive.  (That, of course,
> doesn't work for the general case of non-lexical scoping, or dynamic
> OpenACC 'enter data', etc., I suppose) I suppose your method could easily
> replace and improve that; we shall look into that later.
> 
> That said, in your patch, is this current implementation (explicitly)
> meant or not meant to be active for OpenACC, too, or just OpenMP (I
> couldn't quickly tell), and/or is it (implicitly?) a no-op for OpenACC?

It appears that I have inadvertently enabled it for OpenACC as well!
But everything was tested together, so I assume it works okay for that mode as well.

The entire set of implicit-specific actions are enabled by the setting of
'OMP_CLAUSE_MAP_IMPLICIT_P (clause) = 1' inside gimplify.c:gimplify_adjust_omp_clauses_1,
so in case you want to disable it for OpenACC again, that's where you need to add the guard condition.

>> Also, another adjustment in this patch is how implicitly created clauses are added to the current
>> clause list in gimplify_adjust_omp_clauses(). Instead of simply appending the new clauses to the end,
>> this patch adds them at the position "after initial non-map clauses, but right before any existing
>> map clauses".
> Probably you haven't been testing such a configuration; I've just pushed
> "Fix up 'c-c++-common/goacc/firstprivate-mappings-1.c' for C, non-LP64"
> to devel/omp/gcc-10 branch in commit
> c51cc3b96f0b562deaffcfbcc51043aed216801a, see attached.

Thanks, I was relying on eyeballing to know where to fix testcases like this;
I did fix another similar case, but missed this one.

> 
>> The reason for this is: when combined with other map clauses, for example:
>>
>>     #pragma omp target map(rec.ptr[:N])
>>     for (int i = 0; i < N; i++)
>>       rec.ptr[i] += 1;
>>
>> There will be an implicit map created for map(rec), because of the access inside the target region.
>> The expectation is that 'rec' is implicitly mapped, and then the pointed array-section part by 'rec.ptr'
>> will be mapped, and then attachment to the 'rec.ptr' field of the mapped 'rec' (in that order).
>>
>> If the implicit 'map(rec)' is appended to the end, instead of placed before other maps, the attachment
>> operation will not find anything to attach to, and the entire region will fail.
> But that doesn't (negatively) affect user-visible semantics (OpenMP, and
> also OpenACC, if applicable), in that more/bigger objects then get mapped
> than were before?  (I suppose not?)

It probably won't affect user level semantics, although we should look out if this change in convention
exposes some other bugs.

Chung-Lin