abhinavgaba added a comment. In D134268#4048357 <https://reviews.llvm.org/D134268#4048357>, @tianshilei1992 wrote:
> Hi there, I'm trying to fix > https://github.com/llvm/llvm-project/issues/59160. The faulty case is > basically like the following: > > 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 > } > > It looks like at runtime, we are trying to copy a (device) pointer to a > device pointer by using host to device data transfer. I noticed that's > because we have `TO` flag marked for the argument. However, since `a` and `b` > are in `has_device_addr`, we are not supposed to map the two variables right? The firstt thing is that in the test itself, `has_device_addr(b[0])` is incorrect. `b` inside the "target data" region refers to the device version of `b`, not the host version. So, it is illegal to do `b[0]` (without unified shared memory) because we cannot load from device `b` on the host (target data is executed on host). After that is changed to `has_device_addr(b)`, the test will likely pass. However, it is still true that `has_device_addr(a, b)` are being treated the same as `map(to:a, b)` for arrays, instead of passing the addresses &a, &b directly into the kernel as LITERALs. I think @jyu2 was working on changing that, so she might be able to say what changes are needed in clang for that. 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
