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 _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
