xref: /llvm-project/offload/test/unified_shared_memory/api.c (revision 904b1a850536d273b0e11bd17a7ea642ba3b5bc4)
1*330d8983SJohannes Doerfert // RUN: %libomptarget-compile-generic
2*330d8983SJohannes Doerfert // RUN: env HSA_XNACK=1 \
3*330d8983SJohannes Doerfert // RUN: %libomptarget-run-generic | %fcheck-generic
4*330d8983SJohannes Doerfert // XFAIL: nvptx64-nvidia-cuda
5*330d8983SJohannes Doerfert // XFAIL: nvptx64-nvidia-cuda-LTO
6*330d8983SJohannes Doerfert 
7*330d8983SJohannes Doerfert // REQUIRES: unified_shared_memory
8*330d8983SJohannes Doerfert 
9*330d8983SJohannes Doerfert #include <omp.h>
10*330d8983SJohannes Doerfert #include <stdio.h>
11*330d8983SJohannes Doerfert 
12*330d8983SJohannes Doerfert // End of definitions copied from OpenMP RTL.
13*330d8983SJohannes Doerfert // ---------------------------------------------------------------------------
14*330d8983SJohannes Doerfert 
15*330d8983SJohannes Doerfert #pragma omp requires unified_shared_memory
16*330d8983SJohannes Doerfert 
17*330d8983SJohannes Doerfert #define N 1024
18*330d8983SJohannes Doerfert 
init(int A[],int B[],int C[])19*330d8983SJohannes Doerfert void init(int A[], int B[], int C[]) {
20*330d8983SJohannes Doerfert   for (int i = 0; i < N; ++i) {
21*330d8983SJohannes Doerfert     A[i] = 0;
22*330d8983SJohannes Doerfert     B[i] = 1;
23*330d8983SJohannes Doerfert     C[i] = i;
24*330d8983SJohannes Doerfert   }
25*330d8983SJohannes Doerfert }
26*330d8983SJohannes Doerfert 
main(int argc,char * argv[])27*330d8983SJohannes Doerfert int main(int argc, char *argv[]) {
28*330d8983SJohannes Doerfert   const int device = omp_get_default_device();
29*330d8983SJohannes Doerfert 
30*330d8983SJohannes Doerfert   // CHECK: Initial device: [[INITIAL_DEVICE:[0-9]+]]
31*330d8983SJohannes Doerfert   printf("Initial device: %d\n", omp_get_initial_device());
32*330d8983SJohannes Doerfert   // CHECK: Num devices: [[INITIAL_DEVICE]]
33*330d8983SJohannes Doerfert   printf("Num devices: %d\n", omp_get_num_devices());
34*330d8983SJohannes Doerfert 
35*330d8983SJohannes Doerfert   //
36*330d8983SJohannes Doerfert   // Target alloc & target memcpy
37*330d8983SJohannes Doerfert   //
38*330d8983SJohannes Doerfert   int A[N], B[N], C[N];
39*330d8983SJohannes Doerfert 
40*330d8983SJohannes Doerfert   // Init
41*330d8983SJohannes Doerfert   init(A, B, C);
42*330d8983SJohannes Doerfert 
43*330d8983SJohannes Doerfert   int *pA, *pB, *pC;
44*330d8983SJohannes Doerfert 
45*330d8983SJohannes Doerfert   // map ptrs
46*330d8983SJohannes Doerfert   pA = &A[0];
47*330d8983SJohannes Doerfert   pB = &B[0];
48*330d8983SJohannes Doerfert   pC = &C[0];
49*330d8983SJohannes Doerfert 
50*330d8983SJohannes Doerfert   int *d_A = (int *)omp_target_alloc(N * sizeof(int), device);
51*330d8983SJohannes Doerfert   int *d_B = (int *)omp_target_alloc(N * sizeof(int), device);
52*330d8983SJohannes Doerfert   int *d_C = (int *)omp_target_alloc(N * sizeof(int), device);
53*330d8983SJohannes Doerfert 
54*330d8983SJohannes Doerfert   // CHECK: omp_target_alloc succeeded
55*330d8983SJohannes Doerfert   printf("omp_target_alloc %s\n", d_A && d_B && d_C ? "succeeded" : "failed");
56*330d8983SJohannes Doerfert 
57*330d8983SJohannes Doerfert   omp_target_memcpy(d_B, pB, N * sizeof(int), 0, 0, device,
58*330d8983SJohannes Doerfert                     omp_get_initial_device());
59*330d8983SJohannes Doerfert   omp_target_memcpy(d_C, pC, N * sizeof(int), 0, 0, device,
60*330d8983SJohannes Doerfert                     omp_get_initial_device());
61*330d8983SJohannes Doerfert 
62*330d8983SJohannes Doerfert #pragma omp target is_device_ptr(d_A, d_B, d_C) device(device)
63*330d8983SJohannes Doerfert   {
64*330d8983SJohannes Doerfert #pragma omp parallel for schedule(static, 1)
65*330d8983SJohannes Doerfert     for (int i = 0; i < N; i++) {
66*330d8983SJohannes Doerfert       d_A[i] = d_B[i] + d_C[i] + 1;
67*330d8983SJohannes Doerfert     }
68*330d8983SJohannes Doerfert   }
69*330d8983SJohannes Doerfert 
70*330d8983SJohannes Doerfert   omp_target_memcpy(pA, d_A, N * sizeof(int), 0, 0, omp_get_initial_device(),
71*330d8983SJohannes Doerfert                     device);
72*330d8983SJohannes Doerfert 
73*330d8983SJohannes Doerfert   // CHECK: Test omp_target_memcpy: Succeeded
74*330d8983SJohannes Doerfert   int fail = 0;
75*330d8983SJohannes Doerfert   for (int i = 0; i < N; ++i) {
76*330d8983SJohannes Doerfert     if (A[i] != i + 2)
77*330d8983SJohannes Doerfert       fail++;
78*330d8983SJohannes Doerfert   }
79*330d8983SJohannes Doerfert   if (fail) {
80*330d8983SJohannes Doerfert     printf("Test omp_target_memcpy: Failed\n");
81*330d8983SJohannes Doerfert   } else {
82*330d8983SJohannes Doerfert     printf("Test omp_target_memcpy: Succeeded\n");
83*330d8983SJohannes Doerfert   }
84*330d8983SJohannes Doerfert 
85*330d8983SJohannes Doerfert   //
86*330d8983SJohannes Doerfert   // target_is_present and target_associate/disassociate_ptr
87*330d8983SJohannes Doerfert   //
88*330d8983SJohannes Doerfert   init(A, B, C);
89*330d8983SJohannes Doerfert 
90*330d8983SJohannes Doerfert   // CHECK: B is not present, associating it...
91*330d8983SJohannes Doerfert   // CHECK: omp_target_associate_ptr B succeeded
92*330d8983SJohannes Doerfert   if (!omp_target_is_present(B, device)) {
93*330d8983SJohannes Doerfert     printf("B is not present, associating it...\n");
94*330d8983SJohannes Doerfert     int rc = omp_target_associate_ptr(B, d_B, N * sizeof(int), 0, device);
95*330d8983SJohannes Doerfert     printf("omp_target_associate_ptr B %s\n", !rc ? "succeeded" : "failed");
96*330d8983SJohannes Doerfert   }
97*330d8983SJohannes Doerfert 
98*330d8983SJohannes Doerfert   // CHECK: C is not present, associating it...
99*330d8983SJohannes Doerfert   // CHECK: omp_target_associate_ptr C succeeded
100*330d8983SJohannes Doerfert   if (!omp_target_is_present(C, device)) {
101*330d8983SJohannes Doerfert     printf("C is not present, associating it...\n");
102*330d8983SJohannes Doerfert     int rc = omp_target_associate_ptr(C, d_C, N * sizeof(int), 0, device);
103*330d8983SJohannes Doerfert     printf("omp_target_associate_ptr C %s\n", !rc ? "succeeded" : "failed");
104*330d8983SJohannes Doerfert   }
105*330d8983SJohannes Doerfert 
106*330d8983SJohannes Doerfert // CHECK: Inside target data: A is not present
107*330d8983SJohannes Doerfert // CHECK: Inside target data: B is present
108*330d8983SJohannes Doerfert // CHECK: Inside target data: C is present
109*330d8983SJohannes Doerfert #pragma omp target data map(from : B, C) device(device)
110*330d8983SJohannes Doerfert   {
111*330d8983SJohannes Doerfert     printf("Inside target data: A is%s present\n",
112*330d8983SJohannes Doerfert            omp_target_is_present(A, device) ? "" : " not");
113*330d8983SJohannes Doerfert     printf("Inside target data: B is%s present\n",
114*330d8983SJohannes Doerfert            omp_target_is_present(B, device) ? "" : " not");
115*330d8983SJohannes Doerfert     printf("Inside target data: C is%s present\n",
116*330d8983SJohannes Doerfert            omp_target_is_present(C, device) ? "" : " not");
117*330d8983SJohannes Doerfert 
118*330d8983SJohannes Doerfert #pragma omp target map(from : A) device(device)
119*330d8983SJohannes Doerfert     {
120*330d8983SJohannes Doerfert #pragma omp parallel for schedule(static, 1)
121*330d8983SJohannes Doerfert       for (int i = 0; i < N; i++)
122*330d8983SJohannes Doerfert         A[i] = B[i] + C[i] + 1;
123*330d8983SJohannes Doerfert     }
124*330d8983SJohannes Doerfert   }
125*330d8983SJohannes Doerfert 
126*330d8983SJohannes Doerfert   // CHECK: B is present, disassociating it...
127*330d8983SJohannes Doerfert   // CHECK: omp_target_disassociate_ptr B succeeded
128*330d8983SJohannes Doerfert   // CHECK: C is present, disassociating it...
129*330d8983SJohannes Doerfert   // CHECK: omp_target_disassociate_ptr C succeeded
130*330d8983SJohannes Doerfert   if (omp_target_is_present(B, device)) {
131*330d8983SJohannes Doerfert     printf("B is present, disassociating it...\n");
132*330d8983SJohannes Doerfert     int rc = omp_target_disassociate_ptr(B, device);
133*330d8983SJohannes Doerfert     printf("omp_target_disassociate_ptr B %s\n", !rc ? "succeeded" : "failed");
134*330d8983SJohannes Doerfert   }
135*330d8983SJohannes Doerfert   if (omp_target_is_present(C, device)) {
136*330d8983SJohannes Doerfert     printf("C is present, disassociating it...\n");
137*330d8983SJohannes Doerfert     int rc = omp_target_disassociate_ptr(C, device);
138*330d8983SJohannes Doerfert     printf("omp_target_disassociate_ptr C %s\n", !rc ? "succeeded" : "failed");
139*330d8983SJohannes Doerfert   }
140*330d8983SJohannes Doerfert 
141*330d8983SJohannes Doerfert   // CHECK: Test omp_target_associate_ptr: Succeeded
142*330d8983SJohannes Doerfert   fail = 0;
143*330d8983SJohannes Doerfert   for (int i = 0; i < N; ++i) {
144*330d8983SJohannes Doerfert     if (A[i] != i + 2)
145*330d8983SJohannes Doerfert       fail++;
146*330d8983SJohannes Doerfert   }
147*330d8983SJohannes Doerfert   if (fail) {
148*330d8983SJohannes Doerfert     printf("Test omp_target_associate_ptr: Failed\n");
149*330d8983SJohannes Doerfert   } else {
150*330d8983SJohannes Doerfert     printf("Test omp_target_associate_ptr: Succeeded\n");
151*330d8983SJohannes Doerfert   }
152*330d8983SJohannes Doerfert 
153*330d8983SJohannes Doerfert   omp_target_free(d_A, device);
154*330d8983SJohannes Doerfert   omp_target_free(d_B, device);
155*330d8983SJohannes Doerfert   omp_target_free(d_C, device);
156*330d8983SJohannes Doerfert 
157*330d8983SJohannes Doerfert   printf("Done!\n");
158*330d8983SJohannes Doerfert 
159*330d8983SJohannes Doerfert   return 0;
160*330d8983SJohannes Doerfert }
161