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 - | 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 - | 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 - | 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 | // CHECK-DAG: [[DISTR_LIGHT:@.+]] = private unnamed_addr constant %struct.ident_t { i32 0, i32 2050, i32 3, i32 0, i8* getelementptr inbounds |
13 | // CHECK-DAG: [[FOR_LIGHT:@.+]] = private unnamed_addr constant %struct.ident_t { i32 0, i32 514, i32 3, i32 0, i8* getelementptr inbounds |
14 | // CHECK-DAG: [[LIGHT:@.+]] = private unnamed_addr constant %struct.ident_t { i32 0, i32 2, i32 3, i32 0, i8* getelementptr inbounds |
15 | // CHECK-DAG: [[DISTR_FULL:@.+]] = private unnamed_addr constant %struct.ident_t { i32 0, i32 2050, i32 1, i32 0, i8* getelementptr inbounds |
16 | // CHECK-DAG: [[FULL:@.+]] = private unnamed_addr constant %struct.ident_t { i32 0, i32 2, i32 1, i32 0, i8* getelementptr inbounds |
17 | // CHECK-DAG: [[BAR_LIGHT:@.+]] = private unnamed_addr constant %struct.ident_t { i32 0, i32 66, i32 3, i32 0, i8* getelementptr inbounds |
18 | // CHECK-DAG: [[BAR_FULL:@.+]] = private unnamed_addr constant %struct.ident_t { i32 0, i32 66, i32 1, i32 0, i8* getelementptr inbounds |
19 | // CHECK-NOT: @__omp_offloading_{{.+}}_exec_mode = weak constant i8 1 |
20 | |
21 | void foo() { |
22 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 0, i16 0) |
23 | // CHECK-DAG: [[DISTR_LIGHT]] |
24 | // CHECK-DAG: [[FOR_LIGHT]] |
25 | // CHECK-DAG: [[LIGHT]] |
26 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 0, i16 0) |
27 | // CHECK-DAG: [[DISTR_LIGHT]] |
28 | // CHECK-DAG: [[FOR_LIGHT]] |
29 | // CHECK-DAG: [[LIGHT]] |
30 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 0, i16 0) |
31 | // CHECK-DAG: [[DISTR_LIGHT]] |
32 | // CHECK-DAG: [[FOR_LIGHT]] |
33 | // CHECK-DAG: [[LIGHT]] |
34 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) |
35 | // CHECK-DAG: [[DISTR_FULL]] |
36 | // CHECK-DAG: [[FULL]] |
37 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) |
38 | // CHECK-DAG: [[DISTR_FULL]] |
39 | // CHECK-DAG: [[FULL]] |
40 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) |
41 | // CHECK-DAG: [[DISTR_FULL]] |
42 | // CHECK-DAG: [[FULL]] |
43 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) |
44 | // CHECK-DAG: [[DISTR_FULL]] |
45 | // CHECK-DAG: [[FULL]] |
46 | #pragma omp target teams distribute parallel for simd |
47 | for (int i = 0; i < 10; ++i) |
48 | ; |
49 | #pragma omp target teams distribute parallel for simd schedule(static) |
50 | for (int i = 0; i < 10; ++i) |
51 | ; |
52 | #pragma omp target teams distribute parallel for simd schedule(static, 1) |
53 | for (int i = 0; i < 10; ++i) |
54 | ; |
55 | #pragma omp target teams distribute parallel for simd schedule(auto) |
56 | for (int i = 0; i < 10; ++i) |
57 | ; |
58 | #pragma omp target teams distribute parallel for simd schedule(runtime) |
59 | for (int i = 0; i < 10; ++i) |
60 | ; |
61 | #pragma omp target teams distribute parallel for simd schedule(dynamic) |
62 | for (int i = 0; i < 10; ++i) |
63 | ; |
64 | #pragma omp target teams distribute parallel for simd schedule(guided) |
65 | for (int i = 0; i < 10; ++i) |
66 | ; |
67 | int a; |
68 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 0, i16 0) |
69 | // CHECK-DAG: [[DISTR_LIGHT]] |
70 | // CHECK-DAG: [[FOR_LIGHT]] |
71 | // CHECK-DAG: [[LIGHT]] |
72 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 0, i16 0) |
73 | // CHECK-DAG: [[DISTR_LIGHT]] |
74 | // CHECK-DAG: [[FOR_LIGHT]] |
75 | // CHECK-DAG: [[LIGHT]] |
76 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 0, i16 0) |
77 | // CHECK-DAG: [[DISTR_LIGHT]] |
78 | // CHECK-DAG: [[FOR_LIGHT]] |
79 | // CHECK-DAG: [[LIGHT]] |
80 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) |
81 | // CHECK-DAG: [[DISTR_FULL]] |
82 | // CHECK-DAG: [[FULL]] |
83 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) |
84 | // CHECK-DAG: [[DISTR_FULL]] |
85 | // CHECK-DAG: [[FULL]] |
86 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) |
87 | // CHECK-DAG: [[DISTR_FULL]] |
88 | // CHECK-DAG: [[FULL]] |
89 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) |
90 | // CHECK-DAG: [[DISTR_FULL]] |
91 | // CHECK-DAG: [[FULL]] |
92 | #pragma omp target teams distribute parallel for lastprivate(a) |
93 | for (int i = 0; i < 10; ++i) |
94 | a = i; |
95 | #pragma omp target teams distribute parallel for schedule(static) |
96 | for (int i = 0; i < 10; ++i) |
97 | ; |
98 | #pragma omp target teams distribute parallel for schedule(static, 1) |
99 | for (int i = 0; i < 10; ++i) |
100 | ; |
101 | #pragma omp target teams distribute parallel for schedule(auto) |
102 | for (int i = 0; i < 10; ++i) |
103 | ; |
104 | #pragma omp target teams distribute parallel for schedule(runtime) |
105 | for (int i = 0; i < 10; ++i) |
106 | ; |
107 | #pragma omp target teams distribute parallel for schedule(dynamic) |
108 | for (int i = 0; i < 10; ++i) |
109 | ; |
110 | #pragma omp target teams distribute parallel for schedule(guided) |
111 | for (int i = 0; i < 10; ++i) |
112 | ; |
113 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 0, i16 0) |
114 | // CHECK-DAG: [[DISTR_LIGHT]] |
115 | // CHECK-DAG: [[FOR_LIGHT]] |
116 | // CHECK-DAG: [[LIGHT]] |
117 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 0, i16 0) |
118 | // CHECK-DAG: [[DISTR_LIGHT]] |
119 | // CHECK-DAG: [[FOR_LIGHT]] |
120 | // CHECK-DAG: [[LIGHT]] |
121 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 0, i16 0) |
122 | // CHECK-DAG: [[DISTR_LIGHT]] |
123 | // CHECK-DAG: [[FOR_LIGHT]] |
124 | // CHECK-DAG: [[LIGHT]] |
125 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) |
126 | // CHECK-DAG: [[DISTR_FULL]] |
127 | // CHECK-DAG: [[FULL]] |
128 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) |
129 | // CHECK-DAG: [[DISTR_FULL]] |
130 | // CHECK-DAG: [[FULL]] |
131 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) |
132 | // CHECK-DAG: [[DISTR_FULL]] |
133 | // CHECK-DAG: [[FULL]] |
134 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) |
135 | // CHECK-DAG: [[DISTR_FULL]] |
136 | // CHECK-DAG: [[FULL]] |
137 | #pragma omp target teams |
138 | { |
139 | int b; |
140 | #pragma omp distribute parallel for simd |
141 | for (int i = 0; i < 10; ++i) |
142 | ; |
143 | ; |
144 | } |
145 | #pragma omp target teams |
146 | { |
147 | int b[] = {2, 3, sizeof(int)}; |
148 | #pragma omp distribute parallel for simd schedule(static) |
149 | for (int i = 0; i < 10; ++i) |
150 | ; |
151 | } |
152 | #pragma omp target teams |
153 | { |
154 | int b; |
155 | #pragma omp distribute parallel for simd schedule(static, 1) |
156 | for (int i = 0; i < 10; ++i) |
157 | ; |
158 | int &c = b; |
159 | } |
160 | #pragma omp target teams |
161 | #pragma omp distribute parallel for simd schedule(auto) |
162 | for (int i = 0; i < 10; ++i) |
163 | ; |
164 | #pragma omp target teams |
165 | #pragma omp distribute parallel for simd schedule(runtime) |
166 | for (int i = 0; i < 10; ++i) |
167 | ; |
168 | #pragma omp target teams |
169 | #pragma omp distribute parallel for simd schedule(dynamic) |
170 | for (int i = 0; i < 10; ++i) |
171 | ; |
172 | #pragma omp target teams |
173 | #pragma omp distribute parallel for simd schedule(guided) |
174 | for (int i = 0; i < 10; ++i) |
175 | ; |
176 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 0, i16 0) |
177 | // CHECK-DAG: [[DISTR_LIGHT]] |
178 | // CHECK-DAG: [[FOR_LIGHT]] |
179 | // CHECK-DAG: [[LIGHT]] |
180 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 0, i16 0) |
181 | // CHECK-DAG: [[DISTR_LIGHT]] |
182 | // CHECK-DAG: [[FOR_LIGHT]] |
183 | // CHECK-DAG: [[LIGHT]] |
184 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 0, i16 0) |
185 | // CHECK-DAG: [[DISTR_LIGHT]] |
186 | // CHECK-DAG: [[FOR_LIGHT]] |
187 | // CHECK-DAG: [[LIGHT]] |
188 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) |
189 | // CHECK-DAG: [[DISTR_FULL]] |
190 | // CHECK-DAG: [[FULL]] |
191 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) |
192 | // CHECK-DAG: [[DISTR_FULL]] |
193 | // CHECK-DAG: [[FULL]] |
194 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) |
195 | // CHECK-DAG: [[DISTR_FULL]] |
196 | // CHECK-DAG: [[FULL]] |
197 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) |
198 | // CHECK-DAG: [[DISTR_FULL]] |
199 | // CHECK-DAG: [[FULL]] |
200 | #pragma omp target teams |
201 | #pragma omp distribute parallel for |
202 | for (int i = 0; i < 10; ++i) |
203 | ; |
204 | #pragma omp target teams |
205 | #pragma omp distribute parallel for schedule(static) |
206 | for (int i = 0; i < 10; ++i) |
207 | ; |
208 | #pragma omp target teams |
209 | #pragma omp distribute parallel for schedule(static, 1) |
210 | for (int i = 0; i < 10; ++i) |
211 | ; |
212 | #pragma omp target teams |
213 | #pragma omp distribute parallel for schedule(auto) |
214 | for (int i = 0; i < 10; ++i) |
215 | ; |
216 | #pragma omp target teams |
217 | #pragma omp distribute parallel for schedule(runtime) |
218 | for (int i = 0; i < 10; ++i) |
219 | ; |
220 | #pragma omp target teams |
221 | #pragma omp distribute parallel for schedule(dynamic) |
222 | for (int i = 0; i < 10; ++i) |
223 | ; |
224 | #pragma omp target teams |
225 | #pragma omp distribute parallel for schedule(guided) |
226 | for (int i = 0; i < 10; ++i) |
227 | ; |
228 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 0, i16 0) |
229 | // CHECK-DAG: [[DISTR_LIGHT]] |
230 | // CHECK-DAG: [[FOR_LIGHT]] |
231 | // CHECK-DAG: [[LIGHT]] |
232 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 0, i16 0) |
233 | // CHECK-DAG: [[DISTR_LIGHT]] |
234 | // CHECK-DAG: [[FOR_LIGHT]] |
235 | // CHECK-DAG: [[LIGHT]] |
236 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 0, i16 0) |
237 | // CHECK-DAG: [[DISTR_LIGHT]] |
238 | // CHECK-DAG: [[FOR_LIGHT]] |
239 | // CHECK-DAG: [[LIGHT]] |
240 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) |
241 | // CHECK-DAG: [[DISTR_FULL]] |
242 | // CHECK-DAG: [[FULL]] |
243 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) |
244 | // CHECK-DAG: [[DISTR_FULL]] |
245 | // CHECK-DAG: [[FULL]] |
246 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) |
247 | // CHECK-DAG: [[DISTR_FULL]] |
248 | // CHECK-DAG: [[FULL]] |
249 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) |
250 | // CHECK-DAG: [[DISTR_FULL]] |
251 | // CHECK-DAG: [[FULL]] |
252 | #pragma omp target |
253 | #pragma omp teams |
254 | #pragma omp distribute parallel for |
255 | for (int i = 0; i < 10; ++i) |
256 | ; |
257 | #pragma omp target |
258 | #pragma omp teams |
259 | #pragma omp distribute parallel for schedule(static) |
260 | for (int i = 0; i < 10; ++i) |
261 | ; |
262 | #pragma omp target |
263 | #pragma omp teams |
264 | #pragma omp distribute parallel for schedule(static, 1) |
265 | for (int i = 0; i < 10; ++i) |
266 | ; |
267 | #pragma omp target |
268 | #pragma omp teams |
269 | #pragma omp distribute parallel for schedule(auto) |
270 | for (int i = 0; i < 10; ++i) |
271 | ; |
272 | #pragma omp target |
273 | #pragma omp teams |
274 | #pragma omp distribute parallel for schedule(runtime) |
275 | for (int i = 0; i < 10; ++i) |
276 | ; |
277 | #pragma omp target |
278 | #pragma omp teams |
279 | #pragma omp distribute parallel for schedule(dynamic) |
280 | for (int i = 0; i < 10; ++i) |
281 | ; |
282 | #pragma omp target |
283 | #pragma omp teams |
284 | #pragma omp distribute parallel for schedule(guided) |
285 | for (int i = 0; i < 10; ++i) |
286 | ; |
287 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 0, i16 0) |
288 | // CHECK-DAG: [[FOR_LIGHT]] |
289 | // CHECK-DAG: [[LIGHT]] |
290 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 0, i16 0) |
291 | // CHECK-DAG: [[FOR_LIGHT]] |
292 | // CHECK-DAG: [[LIGHT]] |
293 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 0, i16 0) |
294 | // CHECK-DAG: [[FOR_LIGHT]] |
295 | // CHECK-DAG: [[LIGHT]] |
296 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) |
297 | // CHECK-DAG: [[FULL]] |
298 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) |
299 | // CHECK-DAG: [[FULL]] |
300 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) |
301 | // CHECK-DAG: [[FULL]] |
302 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) |
303 | // CHECK-DAG: [[FULL]] |
304 | #pragma omp target parallel for |
305 | for (int i = 0; i < 10; ++i) |
306 | ; |
307 | #pragma omp target parallel for schedule(static) |
308 | for (int i = 0; i < 10; ++i) |
309 | ; |
310 | #pragma omp target parallel for schedule(static, 1) |
311 | for (int i = 0; i < 10; ++i) |
312 | ; |
313 | #pragma omp target parallel for schedule(auto) |
314 | for (int i = 0; i < 10; ++i) |
315 | ; |
316 | #pragma omp target parallel for schedule(runtime) |
317 | for (int i = 0; i < 10; ++i) |
318 | ; |
319 | #pragma omp target parallel for schedule(dynamic) |
320 | for (int i = 0; i < 10; ++i) |
321 | ; |
322 | #pragma omp target parallel for schedule(guided) |
323 | for (int i = 0; i < 10; ++i) |
324 | ; |
325 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 0, i16 0) |
326 | // CHECK-DAG: [[FOR_LIGHT]] |
327 | // CHECK-DAG: [[LIGHT]] |
328 | // CHECK-DAG: [[BAR_LIGHT]] |
329 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 0, i16 0) |
330 | // CHECK-DAG: [[FOR_LIGHT]] |
331 | // CHECK-DAG: [[LIGHT]] |
332 | // CHECK-DAG: [[BAR_LIGHT]] |
333 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 0, i16 0) |
334 | // CHECK-DAG: [[FOR_LIGHT]] |
335 | // CHECK-DAG: [[LIGHT]] |
336 | // CHECK-DAG: [[BAR_LIGHT]] |
337 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) |
338 | // CHECK-DAG: [[FULL]] |
339 | // CHECK-DAG: [[BAR_FULL]] |
340 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) |
341 | // CHECK-DAG: [[FULL]] |
342 | // CHECK-DAG: [[BAR_FULL]] |
343 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) |
344 | // CHECK-DAG: [[FULL]] |
345 | // CHECK-DAG: [[BAR_FULL]] |
346 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) |
347 | // CHECK-DAG: [[FULL]] |
348 | // CHECK-DAG: [[BAR_FULL]] |
349 | #pragma omp target parallel |
350 | #pragma omp for simd |
351 | for (int i = 0; i < 10; ++i) |
352 | ; |
353 | #pragma omp target parallel |
354 | #pragma omp for simd schedule(static) |
355 | for (int i = 0; i < 10; ++i) |
356 | ; |
357 | #pragma omp target parallel |
358 | #pragma omp for simd schedule(static, 1) |
359 | for (int i = 0; i < 10; ++i) |
360 | ; |
361 | #pragma omp target parallel |
362 | #pragma omp for simd schedule(auto) |
363 | for (int i = 0; i < 10; ++i) |
364 | ; |
365 | #pragma omp target parallel |
366 | #pragma omp for simd schedule(runtime) |
367 | for (int i = 0; i < 10; ++i) |
368 | ; |
369 | #pragma omp target parallel |
370 | #pragma omp for simd schedule(dynamic) |
371 | for (int i = 0; i < 10; ++i) |
372 | ; |
373 | #pragma omp target parallel |
374 | #pragma omp for simd schedule(guided) |
375 | for (int i = 0; i < 10; ++i) |
376 | ; |
377 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) |
378 | // CHECK-DAG: [[FULL]] |
379 | // CHECK-DAG: [[BAR_FULL]] |
380 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 0, i16 0) |
381 | // CHECK-DAG: [[FOR_LIGHT]] |
382 | // CHECK-DAG: [[LIGHT]] |
383 | // CHECK-DAG: [[BAR_LIGHT]] |
384 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 0, i16 0) |
385 | // CHECK-DAG: [[FOR_LIGHT]] |
386 | // CHECK-DAG: [[LIGHT]] |
387 | // CHECK-DAG: [[BAR_LIGHT]] |
388 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) |
389 | // CHECK-DAG: [[FULL]] |
390 | // CHECK-DAG: [[BAR_FULL]] |
391 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) |
392 | // CHECK-DAG: [[FULL]] |
393 | // CHECK-DAG: [[BAR_FULL]] |
394 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) |
395 | // CHECK-DAG: [[FULL]] |
396 | // CHECK-DAG: [[BAR_FULL]] |
397 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) |
398 | // CHECK-DAG: [[FULL]] |
399 | // CHECK-DAG: [[BAR_FULL]] |
400 | #pragma omp target |
401 | #pragma omp parallel |
402 | #pragma omp for simd ordered |
403 | for (int i = 0; i < 10; ++i) |
404 | ; |
405 | #pragma omp target |
406 | #pragma omp parallel |
407 | #pragma omp for simd schedule(static) |
408 | for (int i = 0; i < 10; ++i) |
409 | ; |
410 | #pragma omp target |
411 | #pragma omp parallel |
412 | #pragma omp for simd schedule(static, 1) |
413 | for (int i = 0; i < 10; ++i) |
414 | ; |
415 | #pragma omp target |
416 | #pragma omp parallel |
417 | #pragma omp for simd schedule(auto) |
418 | for (int i = 0; i < 10; ++i) |
419 | ; |
420 | #pragma omp target |
421 | #pragma omp parallel |
422 | #pragma omp for simd schedule(runtime) |
423 | for (int i = 0; i < 10; ++i) |
424 | ; |
425 | #pragma omp target |
426 | #pragma omp parallel |
427 | #pragma omp for simd schedule(dynamic) |
428 | for (int i = 0; i < 10; ++i) |
429 | ; |
430 | #pragma omp target |
431 | #pragma omp parallel |
432 | #pragma omp for simd schedule(guided) |
433 | for (int i = 0; i < 10; ++i) |
434 | ; |
435 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 0, i16 0) |
436 | // CHECK-DAG: [[FOR_LIGHT]] |
437 | // CHECK-DAG: [[LIGHT]] |
438 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 0, i16 0) |
439 | // CHECK-DAG: [[FOR_LIGHT]] |
440 | // CHECK-DAG: [[LIGHT]] |
441 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 0, i16 0) |
442 | // CHECK-DAG: [[FOR_LIGHT]] |
443 | // CHECK-DAG: [[LIGHT]] |
444 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) |
445 | // CHECK-DAG: [[FULL]] |
446 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) |
447 | // CHECK-DAG: [[FULL]] |
448 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) |
449 | // CHECK-DAG: [[FULL]] |
450 | // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) |
451 | // CHECK-DAG: [[FULL]] |
452 | #pragma omp target |
453 | #pragma omp parallel for |
454 | for (int i = 0; i < 10; ++i) |
455 | ; |
456 | #pragma omp target |
457 | #pragma omp parallel for schedule(static) |
458 | for (int i = 0; i < 10; ++i) |
459 | ; |
460 | #pragma omp target |
461 | #pragma omp parallel for schedule(static, 1) |
462 | for (int i = 0; i < 10; ++i) |
463 | ; |
464 | #pragma omp target |
465 | #pragma omp parallel for schedule(auto) |
466 | for (int i = 0; i < 10; ++i) |
467 | ; |
468 | #pragma omp target |
469 | #pragma omp parallel for schedule(runtime) |
470 | for (int i = 0; i < 10; ++i) |
471 | ; |
472 | #pragma omp target |
473 | #pragma omp parallel for schedule(dynamic) |
474 | for (int i = 0; i < 10; ++i) |
475 | ; |
476 | #pragma omp target |
477 | #pragma omp parallel for schedule(guided) |
478 | for (int i = 0; i < 10; ++i) |
479 | ; |
480 | } |
481 | |
482 | #endif |
483 | |
484 | |