[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 D134268#4048357 <https://reviews.llvm.org/D134268#4048357>, @tianshilei1992 wrote:

> Hi there, I'm trying to fix https://github.com/llvm/llvm-project/issues/59160. The faulty case is basically like the following:
>
>   void xoo() {
>     short a[10], b[10];
>     a[1] = 111;
>     b[1] = 111;
>   #pragma omp target data map(to : a[0 : 2], b[0 : 2]) use_device_addr(a, b)
>   #pragma omp target has_device_addr(a) has_device_addr(b[0])
>     {
>       a[1] = 222;
>       b[1] = 222;
>       // CHECK: 222 222
>       printf("%hd %hd %p %p %p\n", a[1], b[1], &a, b, &b);
>     }
>     // CHECK:111
>     printf("%hd %hd %p %p %p\n", a[1], b[1], &a, b, &b); // 111 111 p1d p2d p3d
>   }
>
> It looks like at runtime, we are trying to copy a (device) pointer to a device pointer by using host to device data transfer. I noticed that's because we have `TO` flag marked for the argument. However, since `a` and `b` are in `has_device_addr`, we are not supposed to map the two variables right?

The firstt thing is that in the test itself, `has_device_addr(b[0])` is incorrect. `b` inside the "target data" region refers to the device version of `b`, not the host version. So, it is illegal to do `b[0]` (without unified shared memory) because we cannot load from device `b` on the host (target data is executed on host).

After that is changed to `has_device_addr(b)`, the test will likely pass.

However, it is still true that `has_device_addr(a, b)` are being treated the same as `map(to:a, b)` for arrays, instead of passing the addresses &a, &b directly into the kernel as LITERALs. I think @jyu2 was working on changing that, so she might be able to say what changes are needed in clang for that.


Repository:
  rG LLVM Github Monorepo

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

https://reviews.llvm.org/D141627