<table border="1" cellspacing="0" cellpadding="8">
<tr>
<th>Issue</th>
<td>
<a href=https://github.com/llvm/llvm-project/issues/141054>141054</a>
</td>
</tr>
<tr>
<th>Summary</th>
<td>
[OpenMP] Mapping of refs in structs does not work
</td>
</tr>
<tr>
<th>Labels</th>
<td>
new issue
</td>
</tr>
<tr>
<th>Assignees</th>
<td>
jtb20
</td>
</tr>
<tr>
<th>Reporter</th>
<td>
jtb20
</td>
</tr>
</table>
<pre>
Consider this test case:
```
#include <cstdio>
#include <cstddef>
struct C {
int x;
int &xr;
int *p;
int *≺
int (&ar)[20];
C(int &axr, int *ap, int *&apr, int (&aar)[20]) : xr(axr), p(ap), pr(apr), ar(aar)
{
x = 1;
xr = 1;
*p = 1;
*pr = 1;
ar[0] = 1;
}
};
int ref2()
{
int v;
int tgt1, tgt2;
int *ptr = &tgt2;
int arr[20];
C c(v, &tgt1, ptr, arr);
printf("\n------ ref2 ------\n");
#pragma omp target map(tofrom: c.x)
{
++c.x;
}
printf("ref itself is %p, should refer to %p\n", ((char*)&c + offsetof(C, xr)), &v);
printf("&c is %p\n", (void*) &c);
printf("ref...\n");
fflush(stdout);
#pragma omp target map(tofrom: c.xr)
{
++c.xr;
}
printf("pointer...\n");
fflush(stdout);
#pragma omp target map(tofrom: c.p[:1])
{
++(*c.p);
}
printf("ref to pointer...\n");
printf("ref itself is %p, addr is %p, points to %p\n", ((char*)&c + offsetof(C, pr)), &c.pr, c.pr);
printf("struct base is %p\n", &c);
fflush(stdout);
#pragma omp target map(to: c.pr) map(tofrom: c.pr[:1])
{
++(*c.pr);
}
printf ("ref to array...\n");
printf ("ref itself is %p, &arr[0] is %p\n", ((char*)&c + offsetof(C, ar)), &c.ar[0]);
printf ("struct base is %p\n", &c);
fflush(stdout);
#pragma omp target map(tofrom: c.ar[:1])
{
++c.ar[0];
}
#pragma omp target map(tofrom: c.ar)
{
++c.ar[0];
}
printf ("c.x=%d, c.xr=%d, v=%d, *c.p=%d, tgt1=%d, tgt2=%d, *c.pr=%d, c.ar[0]=%d\n", c.x, c.xr, v, *c.p, tgt1, tgt2, *c.pr, c.ar[0]);
return 0;
}
int main()
{
ref2();
return 0;
}
```
Compiling with offloading to AMDGCN, this crashes on the second target region, mapping `c.xr`.
Running with `LIBOMPTARGET_DEBUG=5`, we can see the problem:
```
omptarget --> Entry 0: Base=0x00007fffbccf0628, Begin=0x00007fffbccf0630, Size=8, Type=0x20, Name=unknown
(base is &c, begin is the actual ref address for c.xr)
omptarget --> Entry 1: Base=0x00007fffbccf0628, Begin=0x00007fffbccf06bc, Size=4, Type=0x1000000000013, Name=unknown
(base is &c, begin is the referent of c.xr)
...
omptarget --> Update pointer (0x000077a901602000) -> [0x000077a901600000]
(mapped ptr &c -> mapped referent of c.xr)
```
(Excerpt, parenthesized lines edited in)
The `Update pointer` should have written back to the actual ref address for `c.xr`, but actually writes back to the base pointer of the struct, `&c`.
I have a rough sketch of a patch to fix this in simple cases, but I think more complicated cases involve fixing https://github.com/llvm/llvm-project/issues/141042 and https://github.com/llvm/llvm-project/issues/141046 also.
</pre>
<img width="1" height="1" alt="" src="http://email.email.llvm.org/o/eJysV19v2zgM_zTKC9FAlmMnechD_rTDgOs27Lrng2LLsVZbEiQ5TffpD5TtJE7TXrtdUKCWSJM_8kdSMndO7pQQC5KsCGM__ZZRwhhJNiPe-FLbRdgabXX-vFhr5WQuLPhSOvDCeci4EyReErokKe3-6JKwWKqsanIBJF5nzudSk_j2qiQXRSeiS-dtk3lYA5muCF0CSOXhQOLTgrD0YIcbS9OtuyVhqblQmRGWckvYnCQrRkmyacVBY03YrLPMD5awdW-Hm7MFSs2ZNFgcmmRzIPES0MYsWJqjusGV6RdBZnoZD8vWCkLpowY4AIk3EPVRABzs5Q4Gfm3vhSK3JFkhwqGATDeYgukpFxiZFQUL0c2D8JjX_XlC_c5HiN_vPLvkwrf-CUsvpdzaYfoB1pARNtujrfaFYNZ422YnJObElLFS-SKgYyRZq5vwC5ChfcZdFA5eIyw2lu9qDro24LndCQ81cjLzurC6Rtay8eElC4StCFuhaJizSzBWFCC9E1UB0gFhSagcV-qmyhEedoxu93uA67aGZlmJ9C9DQaQZegRdFE54jbbXqNiV0rzL0v4Y3UVG8P3e_cDNXsu89YHvZ6-8b0UxHo8vMwhQFFXjSsJmzue68YPcvje1Vyr8mFv7MrnnsIyWygv7YWjvxWZw8MXLqO3gqyADkCWqnly_UQlew1ug_7tweJ7b83Ww5v6khMywhLJxO8ra_1eRdZN4y524WlSDMvp9HjoOEMU1buwHyLFvsQMDeri1_PkNcuAtdsJRchypVxvufaTwS1KOg_o1QB8h5f_oDf56_vv2PWF-mfr3e_m47UFewoTeEJbkbVXjTDku92fPXRufNsKRM1iyF-p2YPsMVLt7IiCcIZ3_4PjksXfV-zgzfmH1RL0VvrEKaMdbHzsepTWX6vKUPj-637Jwdk1b69rISqodPElfYoFWmue49hqW95tP6y8BMd72MstdKRxoBb4U4ESmVd6TasVOaoW6NTcGDZCUhjykdNzC_t4odfREUvrX59XX-28Py--fbh_-2dyufnwi8SZBXGwNTwIyrsAJEZwZq7eVqF_eM3VtOgQ3NyS-hVvl7TNgwEtYhZvphh4opXRaFMU2ywqaYoLWsBI7qa5I4-D-b_kLXw2aD8-mNcOC6AuvcdmoR6WfVLjQzk79iA24hi0axw3EzjPf8ArJCXNdOAeFtmeH4isxRL8dwzY7i2EyiCGix18U_1Y44TYjlAddnMWAs_RlJD9Mzr3oD0Ps1Q7qlM9plFJGKcXBH5Sx_gdS_GFDBExYViKHcLvEeRpe6TavQxp-jsxuD5mwxofDkKN2KZz8JXKopBIORC69yAG7at7W60MpsEyHMZCU9te6ku8FPFnpvVCw5dkj9swbhJ8aIuS08Z1i9RysCDcwEjjoE6eLtuXC9A-jA42k2am5PrdwOFjd7Epwj8Jn2M_AwXB89BoKeWg7WSpwsjaVCB9vrofzGaXqEWptBWS6NpXMOCYlaIFUe13tBZrBNi69Nw4bkt0RdreTvmy240zXhN1V1b7_d2Os_ikQ9J10rkFnd9EkohMGXOV_aiQFXjk9HuWLOJ_Hcz4Si2g6mc6jNE7YqFykeTJjURrH020R5el8Op3Np4JlST5Ncp4UI7lglCU0YSxi0Txi41k2yyfZNiu28yhPJymZUFFzWY0Rx1jb3SggWKD7ZDKq-FZUrvt0VuIJgrT7fLaLAH7b7ByZ0Eo6705mvPRV-Ob-aoS6_4Y3iftucuoCa6elKTDuINfCgdIenrR9HDW2WvxZ5pLJaL9g_wYAAP__D-p6sw">