Gitiles
Code Review Sign In
gerrit-public.fairphone.software / toolchain / llvm-project / 835140a4277e8f0ab6b7436a04c23f1c3a7b51c4 / . / clang / test / OpenMP / nvptx_SPMD_codegen.cpp
blob: cf68ee7dd8dc6dad126cbdc9dfc7a2813d5bac51 [file] [log] [blame]
Alexey Bataev8d8e1232018-08-29 18:32:21 +0000[diff] [blame]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
Alexey Bataev8d8e1232018-08-29 18:32:21 +0000[diff] [blame]12
13void foo() {
14// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 0, i16 0)
15// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 0, i16 0)
16// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 0, i16 0)
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 ;
42int a;
Alexey Bataev8bcc69c2018-11-09 20:03:19 +0000[diff] [blame]43// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 0, i16 0)
Alexey Bataev8d8e1232018-08-29 18:32:21 +0000[diff] [blame]44// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 0, i16 0)
45// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 0, i16 0)
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 0, i16 0)
72// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 0, i16 0)
73// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 0, i16 0)
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
Alexey Bataev8bcc69c2018-11-09 20:03:19 +0000[diff] [blame]79 {
80 int b;
Alexey Bataev8d8e1232018-08-29 18:32:21 +0000[diff] [blame]81#pragma omp distribute parallel for simd
82 for (int i = 0; i < 10; ++i)
83 ;
Alexey Bataev8bcc69c2018-11-09 20:03:19 +0000[diff] [blame]84 ;
85 }
Alexey Bataev8d8e1232018-08-29 18:32:21 +0000[diff] [blame]86#pragma omp target teams
Alexey Bataev8bcc69c2018-11-09 20:03:19 +0000[diff] [blame]87 {
88 int b[] = {2, 3, sizeof(int)};
Alexey Bataev8d8e1232018-08-29 18:32:21 +0000[diff] [blame]89#pragma omp distribute parallel for simd schedule(static)
90 for (int i = 0; i < 10; ++i)
91 ;
Alexey Bataev8bcc69c2018-11-09 20:03:19 +0000[diff] [blame]92 }
Alexey Bataev8d8e1232018-08-29 18:32:21 +0000[diff] [blame]93#pragma omp target teams
Alexey Bataev8bcc69c2018-11-09 20:03:19 +0000[diff] [blame]94 {
95 int b;
Alexey Bataev8d8e1232018-08-29 18:32:21 +0000[diff] [blame]96#pragma omp distribute parallel for simd schedule(static, 1)
97 for (int i = 0; i < 10; ++i)
98 ;
Alexey Bataev8bcc69c2018-11-09 20:03:19 +0000[diff] [blame]99 int &c = b;
100 }
Alexey Bataev8d8e1232018-08-29 18:32:21 +0000[diff] [blame]101#pragma omp target teams
102#pragma omp distribute parallel for simd schedule(auto)
103 for (int i = 0; i < 10; ++i)
104 ;
105#pragma omp target teams
106#pragma omp distribute parallel for simd schedule(runtime)
107 for (int i = 0; i < 10; ++i)
108 ;
109#pragma omp target teams
110#pragma omp distribute parallel for simd schedule(dynamic)
111 for (int i = 0; i < 10; ++i)
112 ;
113#pragma omp target teams
114#pragma omp distribute parallel for simd schedule(guided)
115 for (int i = 0; i < 10; ++i)
116 ;
117// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 0, i16 0)
118// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 0, i16 0)
119// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 0, i16 0)
120// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
121// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
122// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
123// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
124#pragma omp target teams
125#pragma omp distribute parallel for
126 for (int i = 0; i < 10; ++i)
127 ;
128#pragma omp target teams
129#pragma omp distribute parallel for schedule(static)
130 for (int i = 0; i < 10; ++i)
131 ;
132#pragma omp target teams
133#pragma omp distribute parallel for schedule(static, 1)
134 for (int i = 0; i < 10; ++i)
135 ;
136#pragma omp target teams
137#pragma omp distribute parallel for schedule(auto)
138 for (int i = 0; i < 10; ++i)
139 ;
140#pragma omp target teams
141#pragma omp distribute parallel for schedule(runtime)
142 for (int i = 0; i < 10; ++i)
143 ;
144#pragma omp target teams
145#pragma omp distribute parallel for schedule(dynamic)
146 for (int i = 0; i < 10; ++i)
147 ;
148#pragma omp target teams
149#pragma omp distribute parallel for schedule(guided)
150 for (int i = 0; i < 10; ++i)
151 ;
152// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 0, i16 0)
153// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 0, i16 0)
154// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 0, i16 0)
155// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
156// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
157// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
158// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
159#pragma omp target
160#pragma omp teams
161#pragma omp distribute parallel for
162 for (int i = 0; i < 10; ++i)
163 ;
164#pragma omp target
165#pragma omp teams
166#pragma omp distribute parallel for schedule(static)
167 for (int i = 0; i < 10; ++i)
168 ;
169#pragma omp target
170#pragma omp teams
171#pragma omp distribute parallel for schedule(static, 1)
172 for (int i = 0; i < 10; ++i)
173 ;
174#pragma omp target
175#pragma omp teams
176#pragma omp distribute parallel for schedule(auto)
177 for (int i = 0; i < 10; ++i)
178 ;
179#pragma omp target
180#pragma omp teams
181#pragma omp distribute parallel for schedule(runtime)
182 for (int i = 0; i < 10; ++i)
183 ;
184#pragma omp target
185#pragma omp teams
186#pragma omp distribute parallel for schedule(dynamic)
187 for (int i = 0; i < 10; ++i)
188 ;
189#pragma omp target
190#pragma omp teams
191#pragma omp distribute parallel for schedule(guided)
192 for (int i = 0; i < 10; ++i)
193 ;
194// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 0, i16 0)
195// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 0, i16 0)
196// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 0, i16 0)
197// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
198// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
199// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
200// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
201#pragma omp target parallel for
202 for (int i = 0; i < 10; ++i)
203 ;
204#pragma omp target parallel for schedule(static)
205 for (int i = 0; i < 10; ++i)
206 ;
207#pragma omp target parallel for schedule(static, 1)
208 for (int i = 0; i < 10; ++i)
209 ;
210#pragma omp target parallel for schedule(auto)
211 for (int i = 0; i < 10; ++i)
212 ;
213#pragma omp target parallel for schedule(runtime)
214 for (int i = 0; i < 10; ++i)
215 ;
216#pragma omp target parallel for schedule(dynamic)
217 for (int i = 0; i < 10; ++i)
218 ;
219#pragma omp target parallel for schedule(guided)
220 for (int i = 0; i < 10; ++i)
221 ;
222// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 0, i16 0)
223// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 0, i16 0)
224// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 0, i16 0)
225// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
226// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
227// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
228// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
229#pragma omp target parallel
230#pragma omp for simd
231 for (int i = 0; i < 10; ++i)
232 ;
233#pragma omp target parallel
234#pragma omp for simd schedule(static)
235 for (int i = 0; i < 10; ++i)
236 ;
237#pragma omp target parallel
238#pragma omp for simd schedule(static, 1)
239 for (int i = 0; i < 10; ++i)
240 ;
241#pragma omp target parallel
242#pragma omp for simd schedule(auto)
243 for (int i = 0; i < 10; ++i)
244 ;
245#pragma omp target parallel
246#pragma omp for simd schedule(runtime)
247 for (int i = 0; i < 10; ++i)
248 ;
249#pragma omp target parallel
250#pragma omp for simd schedule(dynamic)
251 for (int i = 0; i < 10; ++i)
252 ;
253#pragma omp target parallel
254#pragma omp for simd schedule(guided)
255 for (int i = 0; i < 10; ++i)
256 ;
257// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
258// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 0, i16 0)
259// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 0, i16 0)
260// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
261// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
262// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
263// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
264#pragma omp target
265#pragma omp parallel
266#pragma omp for simd ordered
267 for (int i = 0; i < 10; ++i)
268 ;
269#pragma omp target
270#pragma omp parallel
271#pragma omp for simd schedule(static)
272 for (int i = 0; i < 10; ++i)
273 ;
274#pragma omp target
275#pragma omp parallel
276#pragma omp for simd schedule(static, 1)
277 for (int i = 0; i < 10; ++i)
278 ;
279#pragma omp target
280#pragma omp parallel
281#pragma omp for simd schedule(auto)
282 for (int i = 0; i < 10; ++i)
283 ;
284#pragma omp target
285#pragma omp parallel
286#pragma omp for simd schedule(runtime)
287 for (int i = 0; i < 10; ++i)
288 ;
289#pragma omp target
290#pragma omp parallel
291#pragma omp for simd schedule(dynamic)
292 for (int i = 0; i < 10; ++i)
293 ;
294#pragma omp target
295#pragma omp parallel
296#pragma omp for simd schedule(guided)
297 for (int i = 0; i < 10; ++i)
298 ;
299// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 0, i16 0)
300// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 0, i16 0)
301// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 0, i16 0)
302// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
303// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
304// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
305// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
306#pragma omp target
307#pragma omp parallel for
308 for (int i = 0; i < 10; ++i)
309 ;
310#pragma omp target
311#pragma omp parallel for schedule(static)
312 for (int i = 0; i < 10; ++i)
313 ;
314#pragma omp target
315#pragma omp parallel for schedule(static, 1)
316 for (int i = 0; i < 10; ++i)
317 ;
318#pragma omp target
319#pragma omp parallel for schedule(auto)
320 for (int i = 0; i < 10; ++i)
321 ;
322#pragma omp target
323#pragma omp parallel for schedule(runtime)
324 for (int i = 0; i < 10; ++i)
325 ;
326#pragma omp target
327#pragma omp parallel for schedule(dynamic)
328 for (int i = 0; i < 10; ++i)
329 ;
330#pragma omp target
331#pragma omp parallel for schedule(guided)
332 for (int i = 0; i < 10; ++i)
333 ;
334}
335
336#endif
337