1 | // Test target codegen - host bc file has to be created first. |
2 | // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc |
3 | // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple nvptx64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - -fopenmp-cuda-force-full-runtime | FileCheck %s |
4 | // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm-bc %s -o %t-x86-host.bc |
5 | // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple nvptx-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - -fopenmp-cuda-force-full-runtime | FileCheck %s |
6 | // RUN: %clang_cc1 -verify -fopenmp -fexceptions -fcxx-exceptions -x c++ -triple nvptx-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - -fopenmp-cuda-force-full-runtime | FileCheck %s |
7 | // expected-no-diagnostics |
8 | #ifndef HEADER |
9 | #define HEADER |
10 | |
11 | // CHECK-NOT: @__omp_offloading_{{.+}}_exec_mode = weak constant i8 1 |
12 | |
13 | void foo() { |
14 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) |
15 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) |
16 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) |
17 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) |
18 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) |
19 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) |
20 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) |
21 | #pragma omp target teams distribute parallel for simd |
22 | for (int i = 0; i < 10; ++i) |
23 | ; |
24 | #pragma omp target teams distribute parallel for simd schedule(static) |
25 | for (int i = 0; i < 10; ++i) |
26 | ; |
27 | #pragma omp target teams distribute parallel for simd schedule(static, 1) |
28 | for (int i = 0; i < 10; ++i) |
29 | ; |
30 | #pragma omp target teams distribute parallel for simd schedule(auto) |
31 | for (int i = 0; i < 10; ++i) |
32 | ; |
33 | #pragma omp target teams distribute parallel for simd schedule(runtime) |
34 | for (int i = 0; i < 10; ++i) |
35 | ; |
36 | #pragma omp target teams distribute parallel for simd schedule(dynamic) |
37 | for (int i = 0; i < 10; ++i) |
38 | ; |
39 | #pragma omp target teams distribute parallel for simd schedule(guided) |
40 | for (int i = 0; i < 10; ++i) |
41 | ; |
42 | int a; |
43 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) |
44 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) |
45 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) |
46 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) |
47 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) |
48 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) |
49 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) |
50 | #pragma omp target teams distribute parallel for lastprivate(a) |
51 | for (int i = 0; i < 10; ++i) |
52 | a = i; |
53 | #pragma omp target teams distribute parallel for schedule(static) |
54 | for (int i = 0; i < 10; ++i) |
55 | ; |
56 | #pragma omp target teams distribute parallel for schedule(static, 1) |
57 | for (int i = 0; i < 10; ++i) |
58 | ; |
59 | #pragma omp target teams distribute parallel for schedule(auto) |
60 | for (int i = 0; i < 10; ++i) |
61 | ; |
62 | #pragma omp target teams distribute parallel for schedule(runtime) |
63 | for (int i = 0; i < 10; ++i) |
64 | ; |
65 | #pragma omp target teams distribute parallel for schedule(dynamic) |
66 | for (int i = 0; i < 10; ++i) |
67 | ; |
68 | #pragma omp target teams distribute parallel for schedule(guided) |
69 | for (int i = 0; i < 10; ++i) |
70 | ; |
71 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) |
72 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) |
73 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) |
74 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) |
75 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) |
76 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) |
77 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) |
78 | #pragma omp target teams |
79 | #pragma omp distribute parallel for simd |
80 | for (int i = 0; i < 10; ++i) |
81 | ; |
82 | #pragma omp target teams |
83 | #pragma omp distribute parallel for simd schedule(static) |
84 | for (int i = 0; i < 10; ++i) |
85 | ; |
86 | #pragma omp target teams |
87 | #pragma omp distribute parallel for simd schedule(static, 1) |
88 | for (int i = 0; i < 10; ++i) |
89 | ; |
90 | #pragma omp target teams |
91 | #pragma omp distribute parallel for simd schedule(auto) |
92 | for (int i = 0; i < 10; ++i) |
93 | ; |
94 | #pragma omp target teams |
95 | #pragma omp distribute parallel for simd schedule(runtime) |
96 | for (int i = 0; i < 10; ++i) |
97 | ; |
98 | #pragma omp target teams |
99 | #pragma omp distribute parallel for simd schedule(dynamic) |
100 | for (int i = 0; i < 10; ++i) |
101 | ; |
102 | #pragma omp target teams |
103 | #pragma omp distribute parallel for simd schedule(guided) |
104 | for (int i = 0; i < 10; ++i) |
105 | ; |
106 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) |
107 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) |
108 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) |
109 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) |
110 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) |
111 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) |
112 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) |
113 | #pragma omp target teams |
114 | #pragma omp distribute parallel for |
115 | for (int i = 0; i < 10; ++i) |
116 | ; |
117 | #pragma omp target teams |
118 | #pragma omp distribute parallel for schedule(static) |
119 | for (int i = 0; i < 10; ++i) |
120 | ; |
121 | #pragma omp target teams |
122 | #pragma omp distribute parallel for schedule(static, 1) |
123 | for (int i = 0; i < 10; ++i) |
124 | ; |
125 | #pragma omp target teams |
126 | #pragma omp distribute parallel for schedule(auto) |
127 | for (int i = 0; i < 10; ++i) |
128 | ; |
129 | #pragma omp target teams |
130 | #pragma omp distribute parallel for schedule(runtime) |
131 | for (int i = 0; i < 10; ++i) |
132 | ; |
133 | #pragma omp target teams |
134 | #pragma omp distribute parallel for schedule(dynamic) |
135 | for (int i = 0; i < 10; ++i) |
136 | ; |
137 | #pragma omp target teams |
138 | #pragma omp distribute parallel for schedule(guided) |
139 | for (int i = 0; i < 10; ++i) |
140 | ; |
141 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) |
142 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) |
143 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) |
144 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) |
145 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) |
146 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) |
147 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) |
148 | #pragma omp target |
149 | #pragma omp teams |
150 | #pragma omp distribute parallel for |
151 | for (int i = 0; i < 10; ++i) |
152 | ; |
153 | #pragma omp target |
154 | #pragma omp teams |
155 | #pragma omp distribute parallel for schedule(static) |
156 | for (int i = 0; i < 10; ++i) |
157 | ; |
158 | #pragma omp target |
159 | #pragma omp teams |
160 | #pragma omp distribute parallel for schedule(static, 1) |
161 | for (int i = 0; i < 10; ++i) |
162 | ; |
163 | #pragma omp target |
164 | #pragma omp teams |
165 | #pragma omp distribute parallel for schedule(auto) |
166 | for (int i = 0; i < 10; ++i) |
167 | ; |
168 | #pragma omp target |
169 | #pragma omp teams |
170 | #pragma omp distribute parallel for schedule(runtime) |
171 | for (int i = 0; i < 10; ++i) |
172 | ; |
173 | #pragma omp target |
174 | #pragma omp teams |
175 | #pragma omp distribute parallel for schedule(dynamic) |
176 | for (int i = 0; i < 10; ++i) |
177 | ; |
178 | #pragma omp target |
179 | #pragma omp teams |
180 | #pragma omp distribute parallel for schedule(guided) |
181 | for (int i = 0; i < 10; ++i) |
182 | ; |
183 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) |
184 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) |
185 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) |
186 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) |
187 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) |
188 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) |
189 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) |
190 | #pragma omp target parallel for |
191 | for (int i = 0; i < 10; ++i) |
192 | ; |
193 | #pragma omp target parallel for schedule(static) |
194 | for (int i = 0; i < 10; ++i) |
195 | ; |
196 | #pragma omp target parallel for schedule(static, 1) |
197 | for (int i = 0; i < 10; ++i) |
198 | ; |
199 | #pragma omp target parallel for schedule(auto) |
200 | for (int i = 0; i < 10; ++i) |
201 | ; |
202 | #pragma omp target parallel for schedule(runtime) |
203 | for (int i = 0; i < 10; ++i) |
204 | ; |
205 | #pragma omp target parallel for schedule(dynamic) |
206 | for (int i = 0; i < 10; ++i) |
207 | ; |
208 | #pragma omp target parallel for schedule(guided) |
209 | for (int i = 0; i < 10; ++i) |
210 | ; |
211 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) |
212 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) |
213 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) |
214 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) |
215 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) |
216 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) |
217 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) |
218 | #pragma omp target parallel |
219 | #pragma omp for simd |
220 | for (int i = 0; i < 10; ++i) |
221 | ; |
222 | #pragma omp target parallel |
223 | #pragma omp for simd schedule(static) |
224 | for (int i = 0; i < 10; ++i) |
225 | ; |
226 | #pragma omp target parallel |
227 | #pragma omp for simd schedule(static, 1) |
228 | for (int i = 0; i < 10; ++i) |
229 | ; |
230 | #pragma omp target parallel |
231 | #pragma omp for simd schedule(auto) |
232 | for (int i = 0; i < 10; ++i) |
233 | ; |
234 | #pragma omp target parallel |
235 | #pragma omp for simd schedule(runtime) |
236 | for (int i = 0; i < 10; ++i) |
237 | ; |
238 | #pragma omp target parallel |
239 | #pragma omp for simd schedule(dynamic) |
240 | for (int i = 0; i < 10; ++i) |
241 | ; |
242 | #pragma omp target parallel |
243 | #pragma omp for simd schedule(guided) |
244 | for (int i = 0; i < 10; ++i) |
245 | ; |
246 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) |
247 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) |
248 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) |
249 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) |
250 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) |
251 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) |
252 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) |
253 | #pragma omp target |
254 | #pragma omp parallel |
255 | #pragma omp for simd ordered |
256 | for (int i = 0; i < 10; ++i) |
257 | ; |
258 | #pragma omp target |
259 | #pragma omp parallel |
260 | #pragma omp for simd schedule(static) |
261 | for (int i = 0; i < 10; ++i) |
262 | ; |
263 | #pragma omp target |
264 | #pragma omp parallel |
265 | #pragma omp for simd schedule(static, 1) |
266 | for (int i = 0; i < 10; ++i) |
267 | ; |
268 | #pragma omp target |
269 | #pragma omp parallel |
270 | #pragma omp for simd schedule(auto) |
271 | for (int i = 0; i < 10; ++i) |
272 | ; |
273 | #pragma omp target |
274 | #pragma omp parallel |
275 | #pragma omp for simd schedule(runtime) |
276 | for (int i = 0; i < 10; ++i) |
277 | ; |
278 | #pragma omp target |
279 | #pragma omp parallel |
280 | #pragma omp for simd schedule(dynamic) |
281 | for (int i = 0; i < 10; ++i) |
282 | ; |
283 | #pragma omp target |
284 | #pragma omp parallel |
285 | #pragma omp for simd schedule(guided) |
286 | for (int i = 0; i < 10; ++i) |
287 | ; |
288 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) |
289 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) |
290 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) |
291 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) |
292 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) |
293 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) |
294 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) |
295 | #pragma omp target |
296 | #pragma omp parallel for |
297 | for (int i = 0; i < 10; ++i) |
298 | ; |
299 | #pragma omp target |
300 | #pragma omp parallel for schedule(static) |
301 | for (int i = 0; i < 10; ++i) |
302 | ; |
303 | #pragma omp target |
304 | #pragma omp parallel for schedule(static, 1) |
305 | for (int i = 0; i < 10; ++i) |
306 | ; |
307 | #pragma omp target |
308 | #pragma omp parallel for schedule(auto) |
309 | for (int i = 0; i < 10; ++i) |
310 | ; |
311 | #pragma omp target |
312 | #pragma omp parallel for schedule(runtime) |
313 | for (int i = 0; i < 10; ++i) |
314 | ; |
315 | #pragma omp target |
316 | #pragma omp parallel for schedule(dynamic) |
317 | for (int i = 0; i < 10; ++i) |
318 | ; |
319 | #pragma omp target |
320 | #pragma omp parallel for schedule(guided) |
321 | for (int i = 0; i < 10; ++i) |
322 | ; |
323 | } |
324 | |
325 | #endif |
326 | |
327 | |