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