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

[Shilei Tian via Phabricator via cfe-commits]

tianshilei1992 added a comment.

In D141627#4051851 <https://reviews.llvm.org/D141627#4051851>, @jyu2 wrote:

> That part of code is original add for is_device_address, so I just wonder, if the change could break is_device_address?

Now I kinda think it is not right to mix `is_device_address` and `has_device_addr`.

Basically, `is_device_address` means the list items are device address, so the address should be taken as literal, directly passed to the kernel.

On the other hand, `has_device_addr` indicates that the list items *should* have device address, which means there has to be an entry for that. Note that it is different from the OpenMP `map` clause. OpenMP's `map` clause (w/o `always` of course) means if the list items are not mapped, do it, and transfer the data accordingly; otherwise, use the one in the map table. I think `has_device_addr` only means map table lookup. Use it if found, otherwise undefined behavior (per spec). We are not supposed to update mapping table.

So back this patch, or clang front end, I think the correct way to handle this is to create a new flag, indicating the mapping is supposed to exist. The runtime needs to be changed accordingly in a way that if the flag is set, it should error out if it doesn't find any mapping.

Meanwhile, I think the test case for `has_device_addr` is not correct.

  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
  }

In `target data` we already put `a` and `b` in `use_device_addr`. That indicates all use of `a` and `b` will be the corresponding device addresses. Therefore, in `target` directive, we should use `is_device_address` instead of `has_device_addr`. The correct way to use `has_device_addr` is, we already map the list items by using `target data` w/o `use_device_addr`. Then when we launch a kernel using `target` directive with `has_device_addr`, we tell the `target` region, the list items *should* be there, and use them, otherwise it is an error (we choose to error out for the undefined behavior).

@jyu2 WDYT?


Repository:
  rG LLVM Github Monorepo

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

https://reviews.llvm.org/D141627