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
17void *llvm_omp_target_alloc_host(size_t Size, int DeviceNum);
18void llvm_omp_target_free_host(void *Ptr, int DeviceNum);
19
20int 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

source code of offload/test/sanitizer/ptr_outside_alloc_1.c