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 // REQUIRES: amdgpu
11 // REQUIRES: unified_shared_memory
12
13 // clang-format on
14
15 #include <cstdio>
16
main()17 int main() {
18 int n = 1024;
19
20 // test various mapping types
21 int *a = new int[n];
22 int k = 3;
23 int b[n];
24
25 for (int i = 0; i < n; i++)
26 b[i] = i;
27
28 // clang-format off
29 // INFO_ZERO: Return HstPtrBegin 0x{{.*}} Size=4096 for unified shared memory
30 // INFO_ZERO: Return HstPtrBegin 0x{{.*}} Size=4096 for unified shared memory
31
32 // INFO_COPY: Creating new map entry with HstPtrBase=0x{{.*}}, HstPtrBegin=0x{{.*}}, TgtAllocBegin=0x{{.*}}, TgtPtrBegin=0x{{.*}}, Size=4096,
33 // INFO_COPY: Creating new map entry with HstPtrBase=0x{{.*}}, HstPtrBegin=0x{{.*}}, TgtAllocBegin=0x{{.*}}, TgtPtrBegin=0x{{.*}}, Size=4096,
34 // INFO_COPY: Mapping exists with HstPtrBegin=0x{{.*}}, TgtPtrBegin=0x{{.*}}, Size=4096, DynRefCount=1 (update suppressed)
35 // INFO_COPY: Mapping exists with HstPtrBegin=0x{{.*}}, TgtPtrBegin=0x{{.*}}, Size=4096, DynRefCount=1 (update suppressed)
36 // clang-format on
37 #pragma omp target teams distribute parallel for map(tofrom : a[ : n]) \
38 map(to : b[ : n])
39 for (int i = 0; i < n; i++)
40 a[i] = i + b[i] + k;
41
42 int err = 0;
43 for (int i = 0; i < n; i++)
44 if (a[i] != i + b[i] + k)
45 err++;
46
47 // CHECK: PASS
48 if (err == 0)
49 printf("PASS\n");
50 return err;
51 }
52