offload/test/sanitizer/use_after_free_1.c

   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[N / 2];
  27     *P = 3;
  28   }
  29   llvm_omp_target_free_host(A, omp_get_default_device());
  30   // clang-format off
  31 // CHECK: OFFLOAD ERROR: Memory access fault by GPU {{.*}} (agent 0x{{.*}}) at virtual address [[PTR:0x[0-9a-z]*]]. Reasons: {{.*}}
  32 // NTRCE: Use 'OFFLOAD_TRACK_ALLOCATION_TRACES=true' to track device allocations
  33 // TRACE: Device pointer [[PTR]] points into prior host-issued allocation:
  34 // TRACE: Last deallocation:
  35 // TRACE: Last allocation of size 1073741824
  36 // clang-format on
  37 #pragma omp target
  38   { *P = 5; }
  39 }