1 | // clang-format off |
2 | // RUN: %libomptarget-compileopt-generic |
3 | // RUN: %not --crash env -u LLVM_DISABLE_SYMBOLIZATION %libomptarget-run-generic 2>&1 | %fcheck-generic --check-prefixes=CHECK,NTRCE |
4 | // RUN: %libomptarget-compileopt-generic |
5 | // RUN: %not --crash env -u LLVM_DISABLE_SYMBOLIZATION OFFLOAD_TRACK_ALLOCATION_TRACES=1 %libomptarget-run-generic 2>&1 | %fcheck-generic --check-prefixes=CHECK,TRACE |
6 | // clang-format on |
7 | |
8 | // UNSUPPORTED: aarch64-unknown-linux-gnu |
9 | // UNSUPPORTED: aarch64-unknown-linux-gnu-LTO |
10 | // UNSUPPORTED: x86_64-unknown-linux-gnu |
11 | // UNSUPPORTED: x86_64-unknown-linux-gnu-LTO |
12 | // UNSUPPORTED: s390x-ibm-linux-gnu |
13 | // UNSUPPORTED: s390x-ibm-linux-gnu-LTO |
14 | |
15 | #include <omp.h> |
16 | |
17 | void *llvm_omp_target_alloc_host(size_t Size, int DeviceNum); |
18 | void llvm_omp_target_free_host(void *Ptr, int DeviceNum); |
19 | |
20 | int main() { |
21 | int N = (1 << 30); |
22 | char *A = (char *)llvm_omp_target_alloc_host(N, omp_get_default_device()); |
23 | char *P; |
24 | #pragma omp target map(from : P) |
25 | { |
26 | P = &A[0]; |
27 | *P = 3; |
28 | } |
29 | // clang-format off |
30 | // CHECK: OFFLOAD ERROR: memory access fault by GPU {{.*}} (agent 0x{{.*}}) at virtual address [[PTR:0x[0-9a-z]*]]. Reasons: {{.*}} |
31 | // NTRCE: Use 'OFFLOAD_TRACK_ALLOCATION_TRACES=true' to track device allocations |
32 | // TRACE: Device pointer [[PTR]] does not point into any (current or prior) host-issued allocation. |
33 | // TRACE: Closest host-issued allocation (distance 4096 bytes; might be by page): |
34 | // TRACE: Last allocation of size 1073741824 |
35 | // clang-format on |
36 | #pragma omp target |
37 | { P[-4] = 5; } |
38 | |
39 | llvm_omp_target_free_host(Ptr: A, DeviceNum: omp_get_default_device()); |
40 | } |
41 | |