1// clang-format off
2// RUN: %libomptarget-compilexx-generic
3// RUN: env OMPX_APU_MAPS=1 HSA_XNACK=1 LIBOMPTARGET_INFO=30 %libomptarget-run-generic 2>&1 \
4// RUN: | %fcheck-generic -check-prefix=INFO_ZERO -check-prefix=CHECK
5
6// RUN: %libomptarget-compilexx-generic
7// RUN: env HSA_XNACK=0 LIBOMPTARGET_INFO=30 %libomptarget-run-generic 2>&1 \
8// RUN: | %fcheck-generic -check-prefix=INFO_COPY -check-prefix=CHECK
9
10// UNSUPPORTED: aarch64-unknown-linux-gnu
11// UNSUPPORTED: aarch64-unknown-linux-gnu-LTO
12// UNSUPPORTED: nvptx64-nvidia-cuda
13// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
14// UNSUPPORTED: x86_64-pc-linux-gnu
15// UNSUPPORTED: x86_64-pc-linux-gnu-LTO
16// UNSUPPORTED: s390x-ibm-linux-gnu
17// UNSUPPORTED: s390x-ibm-linux-gnu-LTO
18
19// REQUIRES: unified_shared_memory
20
21// clang-format on
22
23#include <cstdio>
24
25int main() {
26 int n = 1024;
27
28 // test various mapping types
29 int *a = new int[n];
30 int k = 3;
31 int b[n];
32
33 for (int i = 0; i < n; i++)
34 b[i] = i;
35
36 // clang-format off
37 // INFO_ZERO: Return HstPtrBegin 0x{{.*}} Size=4096 for unified shared memory
38 // INFO_ZERO: Return HstPtrBegin 0x{{.*}} Size=4096 for unified shared memory
39
40 // INFO_COPY: Creating new map entry with HstPtrBase=0x{{.*}}, HstPtrBegin=0x{{.*}}, TgtAllocBegin=0x{{.*}}, TgtPtrBegin=0x{{.*}}, Size=4096,
41 // INFO_COPY: Creating new map entry with HstPtrBase=0x{{.*}}, HstPtrBegin=0x{{.*}}, TgtAllocBegin=0x{{.*}}, TgtPtrBegin=0x{{.*}}, Size=4096,
42 // INFO_COPY: Mapping exists with HstPtrBegin=0x{{.*}}, TgtPtrBegin=0x{{.*}}, Size=4096, DynRefCount=1 (update suppressed)
43 // INFO_COPY: Mapping exists with HstPtrBegin=0x{{.*}}, TgtPtrBegin=0x{{.*}}, Size=4096, DynRefCount=1 (update suppressed)
44// clang-format on
45#pragma omp target teams distribute parallel for map(tofrom : a[ : n]) \
46 map(to : b[ : n])
47 for (int i = 0; i < n; i++)
48 a[i] = i + b[i] + k;
49
50 int err = 0;
51 for (int i = 0; i < n; i++)
52 if (a[i] != i + b[i] + k)
53 err++;
54
55 // CHECK: PASS
56 if (err == 0)
57 printf(format: "PASS\n");
58 return err;
59}
60

source code of offload/test/mapping/auto_zero_copy.cpp