[llvm-bugs] [Bug 141054] [OpenMP] Mapping of refs in structs does not work

[LLVM Bugs via llvm-bugs]

Issue	141054
Summary	[OpenMP] Mapping of refs in structs does not work
Labels	new issue
Assignees	jtb20
Reporter	jtb20

    Consider this test case:
```
#include <cstdio>
#include <cstddef>

struct C {
  int x;
  int &xr;
  int *p;
 int *&pr;
  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.

_______________________________________________
llvm-bugs mailing list
llvm-bugs@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-bugs