1 // clang-format off
2 // RUN: %libomptarget-compilexx-generic
3 // RUN: env LIBOMPTARGET_INFO=32 %libomptarget-run-generic 2>&1 | %fcheck-generic --check-prefix=NO-USM
4 //
5 // RUN: %libomptarget-compilexxx-generic-force-usm
6 // RUN: env HSA_XNACK=1 LIBOMPTARGET_INFO=32 \
7 // RUN: %libomptarget-run-generic 2>&1 | %fcheck-generic --check-prefix=FORCE-USM
8 //
9 // REQUIRES: unified_shared_memory
10 //
11 // UNSUPPORTED: nvptx64-nvidia-cuda
12 // UNSUPPORTED: nvptx64-nvidia-cuda-LTO
13 // clang-format on
14
15 #include <cassert>
16 #include <cstdio>
17 #include <cstdlib>
18
19 int GI;
20 #pragma omp declare target
21 int *pGI;
22 #pragma omp end declare target
23
main(void)24 int main(void) {
25
26 GI = 0;
27 // Implicit mappings
28 int alpha = 1;
29 int beta[3] = {2, 5, 8};
30
31 // Require map clauses for non-USM execution
32 pGI = (int *)malloc(sizeof(int));
33 *pGI = 42;
34
35 #pragma omp target map(pGI[ : 1], GI)
36 {
37 GI = 1 * alpha;
38 *pGI = 2 * beta[1];
39 }
40
41 assert(GI == 1);
42 assert(*pGI == 10);
43
44 printf("SUCCESS\n");
45
46 return 0;
47 }
48
49 // clang-format off
50 // NO-USM: omptarget device 0 info: Copying data from host to device, HstPtr={{.*}}, TgtPtr={{.*}}, Size=4
51 // NO-USM-NEXT: omptarget device 0 info: Copying data from host to device, HstPtr={{.*}}, TgtPtr={{.*}}, Size=12
52 // NO-USM-NEXT: omptarget device 0 info: Copying data from host to device, HstPtr={{.*}}, TgtPtr={{.*}}, Size=4
53 // NO-USM-NEXT: omptarget device 0 info: Copying data from host to device, HstPtr={{.*}}, TgtPtr={{.*}}, Size=8, Name=pGI
54 // NO-USM-NEXT: omptarget device 0 info: Copying data from device to host, TgtPtr={{.*}}, HstPtr={{.*}}, Size=4
55 // NO-USM-NEXT: omptarget device 0 info: Copying data from device to host, TgtPtr={{.*}}, HstPtr={{.*}}, Size=12
56 // NO-USM-NEXT: omptarget device 0 info: Copying data from device to host, TgtPtr={{.*}}, HstPtr={{.*}}, Size=4
57 // NO-USM-NEXT: SUCCESS
58
59 // FORCE-USM: SUCCESS
60 //
61 // clang-format on
62