[PATCH] D141627: [Clang][OpenMP] Fix the issue that list items in `has_device_addr` are still mapped to the target device

[Abhinav Gaba via Phabricator via cfe-commits]

abhinavgaba added a comment.

In D141627#4048584 <https://reviews.llvm.org/D141627#4048584>, @tianshilei1992 wrote:

> I agree that b is not right here, but that doesn’t matter because I stepped into the runtime library and it crashed when processing a.
>
> And why are they treated as to?

Treating `has_device_addr(a)` as `map(a)` is incorrect. I think it is just a vestige of the prior implementation, where `has_device_addr `was fully ignored and instead  `(map(tofrom))` kicked in for `a` (which is the implicit map for arrays).

The test likely passed on x86_64 plugin because it just re-mapped the output of `use_device_addr(a)`, which is a device address, again, but on architectures without unified memory, this re-mapping won't work, hence the failure you see with Cuda.

  #include <stdio.h>
  int main() {
      short a[10], b[10];
      a[1] = 111;
      b[1] = 111;
      printf("%hd %hd %p %p\n", a[1], b[1], &a, &b); // 111 111 p1h p2h
    #pragma omp target data map(to : a[0 : 2], b[0 : 2]) use_device_addr(a, b)
      {
      printf("%p %p\n", &a, &b); // p1d p2d
    #pragma omp target has_device_addr(a) has_device_addr(b)
      {
        a[1] = 222;
        b[1] = 222;
        printf("%hd %hd %p %p\n", a[1], b[1], &a, &b); // 222 222 p1d p2d
      }
      }
      // CHECK:111
      printf("%hd %hd %p %p\n", a[1], b[1], &a, &b); // 111 111 p1h p2h
    }



  $ clang -O0 -fopenmp -fopenmp-targets=x86_64 hda_test.c -fopenmp-version=51 && ./a.out
  111 111 0x7fff2a47ecb0 0x7fff2a47ec90 // 111 111 p1h p2h
  0x55f3cf685b10 0x55f3cf685c10 // p1d p2d: device versions of p1h, p2h
  222 222 0x55f3cf685d70 0x55f3cf685e70 // p1dd p2dd:  another different device version of the two, because of tthe remapping. These should have been "p1d p2d".
  111 111 0x7fff2a47ecb0 0x7fff2a47ec90 // 111 111 p1h p2h

  


Repository:
  rG LLVM Github Monorepo

CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D141627/new/

https://reviews.llvm.org/D141627