abhinavgaba added subscribers: dreachem, kkwli0.
abhinavgaba added a comment.
> 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.
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)
> 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)`.
One example use-case from @dreachem is this:
#pragma omp requires unified_shared_memory
int x ;
printg("%p\n", &x); // p1h
printf("%d\n", omp_target_is_present(&x, omp_get_default_device())); // 0 (x
is not "present", as per the OpenMP runtime)
#pragma omp target has_device_addr(x)
print("%p\n", &x); // p1h (same address as on the host side)
}
In this case, because of unified shared memory, `x` is accessible on device as
well, even though it is not mapped, or made declare target etc. So we need to
pass the address of x into the region, even though `omp_target_is_present`
would return false for it.
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.
%x = alloca i32 ; Original allocation for x
Map:
<ptr %x, ptr %x, sizeof(ptr), PARAM|LITERAL>
Outlined function:
define void @outlined...(ptr %x) {
...
call i32 @printf(..., ptr %x)
...
}
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
