[OpenMP] Mapping of refs in structs does not work

[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 #141042 and #141046 also.

[llvmbot]

@llvm/issue-subscribers-openmp

Author: Julian Brown (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.
</details>