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 | |
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(size: 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(format: "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 | |