Skip to content

[OpenMP] Mapping of refs in structs does not work #141054

Open
@jtb20

Description

@jtb20
Contributor

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.

Activity

self-assigned this
on May 22, 2025
llvmbot

llvmbot commented on May 22, 2025

@llvmbot
Member

@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, `&amp;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>
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Metadata

Metadata

Assignees

Labels

Type

No type

Projects

No projects

Milestone

No milestone

Relationships

None yet

    Development

    No branches or pull requests

      Participants

      @jtb20@EugeneZelenko@llvmbot

      Issue actions

        [OpenMP] Mapping of refs in structs does not work · Issue #141054 · llvm/llvm-project