https://gcc.gnu.org/bugzilla/show_bug.cgi?id=104696
Tobias Burnus <burnus at gcc dot gnu.org> changed:
What |Removed |Added
----------------------------------------------------------------------------
Summary|[12 Regression][OpenMP] |[OpenMP] Implicit mapping
|Implicit mapping breaks |breaks struct mapping
|struct mapping |
Keywords| |missed-optimization
--- Comment #1 from Tobias Burnus <burnus at gcc dot gnu.org> ---
Looking closer at it, I no longer think it is a regression (to be checked if
deemed important).
But it looks as if there are two problems - one wrong-code one and one
missed-optimization one.
Namely, I think reason for both issues is that
map(to:var3.r[1].d [len: 88]) ...
is not turned into
map(struct:var3 [len: 1]) map(to:var3.r[1].d [len: 88])
but into 'to' + pointer assign. This does not even work when using the 'always'
modifier.
('struct:' appears when using 'Q' instead of 'R(2)'.)
That that the struct not detected seems to be because array and component refs
are mixed – hiding that the var is memory wise in the same struct.
I believe with that fixed, it would work correctly.
* * *
For C/C++, it "works".
But: it still does not detect that the member is part of the whole struct - and
allocates pointlessly too much memory. To illustrate this, I added a large
'arr' element.
Otherwise, that the reason that it works in C/C++ is that ATTACH and not
pointer assign is used.
Namely:
------------ test.c--------------
struct s { int *d; };
struct s2 { struct s r[5], q, arr[1024][1024]; };
int main () {
struct s2 x;
x.q.d = __builtin_malloc(sizeof(int));
x.r[1].d = __builtin_malloc(sizeof(int));
#pragma omp target map(tofrom: x.q.d[:1])
*x.q.d = 2;
#pragma omp target map(tofrom: x.r[1].d[:1])
*x.r[1].d = 3;
__builtin_printf("%d, %d\n", *x.q.d, *x.r[1].d);
return 0;
}
------------ test.c--------------
gives:
#pragma omp target num_teams(1) thread_limit(0)
map(tofrom:x [len: 8388656][implicit]) map(tofrom:*_3 [len: 4])
map(attach:x.q.d [bias: 0])
#pragma omp target num_teams(1) thread_limit(0)
map(tofrom:x [len: 8388656][implicit])
map(tofrom:*_4 [len: 4]) map(attach:x.r[1].d [bias: 0])
