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