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
