tianshilei1992 added a comment.
Hi there, I'm trying to fix https://github.com/llvm/llvm-project/issues/59160.
The faulty case is basically like following:
cpp
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?
================
Comment at: clang/lib/CodeGen/CGOpenMPRuntime.cpp:9098
+ CombinedInfo.Types.push_back(
+ (Cap->capturesVariable() ? OMP_MAP_TO : OMP_MAP_LITERAL) |
+ OMP_MAP_TARGET_PARAM);
----------------
The variable used in `has_device_addr` indicates it already has device address
and can be accessed directly. Why do we need `OMP_MAP_TO` here? `OMP_MAP_TO`
indicates we are gonna transfer data to the device right?
Repository:
rG LLVM Github Monorepo
CHANGES SINCE LAST ACTION
https://reviews.llvm.org/D134268/new/
https://reviews.llvm.org/D134268
_______________________________________________
cfe-commits mailing list
[email protected]
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
