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