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 
13 void 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     ;
42 int a;
43 // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 0, i16 0)
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
79    {
80      int b;
81 #pragma omp distribute parallel for simd
82   for (int i = 0; i < 10; ++i)
83     ;
84   ;
85    }
86 #pragma omp target teams
87    {
88      int b[] = {2, 3, sizeof(int)};
89 #pragma omp distribute parallel for simd schedule(static)
90   for (int i = 0; i < 10; ++i)
91     ;
92    }
93 #pragma omp target teams
94    {
95      int b;
96 #pragma omp distribute parallel for simd schedule(static, 1)
97   for (int i = 0; i < 10; ++i)
98     ;
99   int &c = b;
100    }
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 
338