[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#4052323 <https://reviews.llvm.org/D141627#4052323>, @abhinavgaba wrote:

>> 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).
>
> I think you are talking about `is_device_ptr` clause. There is no `is_device_address` clause in OpenMP.

Oh that's correct. I directly copied from Jennifer's comment. ;-)

> The is_device_ptr clause is meant only for "ptrs" (pointers). For example:
>
>   int *p = omp_target_alloc(...);
>   #pragma omp target is_device_ptr(p)

That's true, but what about the case I mentioned? It is also supposed to use `is_device_ptr`.

>> 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.
>
> Based on a brief discussion with some members of the OpenMP spec committee, the idea for "has_device_addr" is to have the address passed-in directly (as a literal, similar to `is_device_ptr`) into the target region, without any map lookup. So, there is no requirement that the variable has to be mapped, or tracked by libomptarget. That requirement is for `map(present:x)`.
> In terms of the code emitted, the original idea of passing the address in as a LITERAL, similar to `is_device_ptr` is the right way to think about it.

No. I think you are mixing things up. The spec says:

> The has_device_addr clause indicates that its list items already have device addresses and therefore they may be directly accessed from a target device.

It only indicates the list items already have device addresses. I don't think it has another level of meaning that, the list of variables listed are device addresses. The second part above is, "they may be directly accessed from a target device". My reading is, they may be directly accessed from a target device "without a mapping", which exactly the `map(present:x)` you suggested indicates. And yes, we don't need extra flag for that. `present` is exactly we need here.


Repository:
  rG LLVM Github Monorepo

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

https://reviews.llvm.org/D141627