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