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