abhinavgaba added a comment. In D141627#4048584 <https://reviews.llvm.org/D141627#4048584>, @tianshilei1992 wrote:
> I agree that b is not right here, but that doesn’t matter because I stepped > into the runtime library and it crashed when processing a. > > And why are they treated as to? Treating `has_device_addr(a)` as `map(a)` is incorrect. I think it is just a vestige of the prior implementation, where `has_device_addr `was fully ignored and instead `(map(tofrom))` kicked in for `a` (which is the implicit map for arrays). The test likely passed on x86_64 plugin because it just re-mapped the output of `use_device_addr(a)`, which is a device address, again, but on architectures without unified memory, this re-mapping won't work, hence the failure you see with Cuda. #include <stdio.h> int main() { short a[10], b[10]; a[1] = 111; b[1] = 111; printf("%hd %hd %p %p\n", a[1], b[1], &a, &b); // 111 111 p1h p2h #pragma omp target data map(to : a[0 : 2], b[0 : 2]) use_device_addr(a, b) { printf("%p %p\n", &a, &b); // p1d p2d #pragma omp target has_device_addr(a) has_device_addr(b) { a[1] = 222; b[1] = 222; printf("%hd %hd %p %p\n", a[1], b[1], &a, &b); // 222 222 p1d p2d } } // CHECK:111 printf("%hd %hd %p %p\n", a[1], b[1], &a, &b); // 111 111 p1h p2h } $ clang -O0 -fopenmp -fopenmp-targets=x86_64 hda_test.c -fopenmp-version=51 && ./a.out 111 111 0x7fff2a47ecb0 0x7fff2a47ec90 // 111 111 p1h p2h 0x55f3cf685b10 0x55f3cf685c10 // p1d p2d: device versions of p1h, p2h 222 222 0x55f3cf685d70 0x55f3cf685e70 // p1dd p2dd: another different device version of the two, because of tthe remapping. These should have been "p1d p2d". 111 111 0x7fff2a47ecb0 0x7fff2a47ec90 // 111 111 p1h p2h 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