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
11 // CHECK-NOT: @__omp_offloading_{{.+}}_exec_mode = weak constant i8 1
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)
24 #pragma omp target teams distribute parallel for simd schedule(static)
25 for (int i = 0; i < 10; ++i)
27 #pragma omp target teams distribute parallel for simd schedule(static, 1)
28 for (int i = 0; i < 10; ++i)
30 #pragma omp target teams distribute parallel for simd schedule(auto)
31 for (int i = 0; i < 10; ++i)
33 #pragma omp target teams distribute parallel for simd schedule(runtime)
34 for (int i = 0; i < 10; ++i)
36 #pragma omp target teams distribute parallel for simd schedule(dynamic)
37 for (int i = 0; i < 10; ++i)
39 #pragma omp target teams distribute parallel for simd schedule(guided)
40 for (int i = 0; i < 10; ++i)
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)
53 #pragma omp target teams distribute parallel for schedule(static)
54 for (int i = 0; i < 10; ++i)
56 #pragma omp target teams distribute parallel for schedule(static, 1)
57 for (int i = 0; i < 10; ++i)
59 #pragma omp target teams distribute parallel for schedule(auto)
60 for (int i = 0; i < 10; ++i)
62 #pragma omp target teams distribute parallel for schedule(runtime)
63 for (int i = 0; i < 10; ++i)
65 #pragma omp target teams distribute parallel for schedule(dynamic)
66 for (int i = 0; i < 10; ++i)
68 #pragma omp target teams distribute parallel for schedule(guided)
69 for (int i = 0; i < 10; ++i)
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)
82 #pragma omp target teams
83 #pragma omp distribute parallel for simd schedule(static)
84 for (int i = 0; i < 10; ++i)
86 #pragma omp target teams
87 #pragma omp distribute parallel for simd schedule(static, 1)
88 for (int i = 0; i < 10; ++i)
90 #pragma omp target teams
91 #pragma omp distribute parallel for simd schedule(auto)
92 for (int i = 0; i < 10; ++i)
94 #pragma omp target teams
95 #pragma omp distribute parallel for simd schedule(runtime)
96 for (int i = 0; i < 10; ++i)
98 #pragma omp target teams
99 #pragma omp distribute parallel for simd schedule(dynamic)
100 for (int i = 0; i < 10; ++i)
102 #pragma omp target teams
103 #pragma omp distribute parallel for simd schedule(guided)
104 for (int i = 0; i < 10; ++i)
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)
117 #pragma omp target teams
118 #pragma omp distribute parallel for schedule(static)
119 for (int i = 0; i < 10; ++i)
121 #pragma omp target teams
122 #pragma omp distribute parallel for schedule(static, 1)
123 for (int i = 0; i < 10; ++i)
125 #pragma omp target teams
126 #pragma omp distribute parallel for schedule(auto)
127 for (int i = 0; i < 10; ++i)
129 #pragma omp target teams
130 #pragma omp distribute parallel for schedule(runtime)
131 for (int i = 0; i < 10; ++i)
133 #pragma omp target teams
134 #pragma omp distribute parallel for schedule(dynamic)
135 for (int i = 0; i < 10; ++i)
137 #pragma omp target teams
138 #pragma omp distribute parallel for schedule(guided)
139 for (int i = 0; i < 10; ++i)
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 {{.+}})
150 #pragma omp distribute parallel for
151 for (int i = 0; i < 10; ++i)
155 #pragma omp distribute parallel for schedule(static)
156 for (int i = 0; i < 10; ++i)
160 #pragma omp distribute parallel for schedule(static, 1)
161 for (int i = 0; i < 10; ++i)
165 #pragma omp distribute parallel for schedule(auto)
166 for (int i = 0; i < 10; ++i)
170 #pragma omp distribute parallel for schedule(runtime)
171 for (int i = 0; i < 10; ++i)
175 #pragma omp distribute parallel for schedule(dynamic)
176 for (int i = 0; i < 10; ++i)
180 #pragma omp distribute parallel for schedule(guided)
181 for (int i = 0; i < 10; ++i)
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)
193 #pragma omp target parallel for schedule(static)
194 for (int i = 0; i < 10; ++i)
196 #pragma omp target parallel for schedule(static, 1)
197 for (int i = 0; i < 10; ++i)
199 #pragma omp target parallel for schedule(auto)
200 for (int i = 0; i < 10; ++i)
202 #pragma omp target parallel for schedule(runtime)
203 for (int i = 0; i < 10; ++i)
205 #pragma omp target parallel for schedule(dynamic)
206 for (int i = 0; i < 10; ++i)
208 #pragma omp target parallel for schedule(guided)
209 for (int i = 0; i < 10; ++i)
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
220 for (int i = 0; i < 10; ++i)
222 #pragma omp target parallel
223 #pragma omp for simd schedule(static)
224 for (int i = 0; i < 10; ++i)
226 #pragma omp target parallel
227 #pragma omp for simd schedule(static, 1)
228 for (int i = 0; i < 10; ++i)
230 #pragma omp target parallel
231 #pragma omp for simd schedule(auto)
232 for (int i = 0; i < 10; ++i)
234 #pragma omp target parallel
235 #pragma omp for simd schedule(runtime)
236 for (int i = 0; i < 10; ++i)
238 #pragma omp target parallel
239 #pragma omp for simd schedule(dynamic)
240 for (int i = 0; i < 10; ++i)
242 #pragma omp target parallel
243 #pragma omp for simd schedule(guided)
244 for (int i = 0; i < 10; ++i)
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 {{.+}})
255 #pragma omp for simd ordered
256 for (int i = 0; i < 10; ++i)
260 #pragma omp for simd schedule(static)
261 for (int i = 0; i < 10; ++i)
265 #pragma omp for simd schedule(static, 1)
266 for (int i = 0; i < 10; ++i)
270 #pragma omp for simd schedule(auto)
271 for (int i = 0; i < 10; ++i)
275 #pragma omp for simd schedule(runtime)
276 for (int i = 0; i < 10; ++i)
280 #pragma omp for simd schedule(dynamic)
281 for (int i = 0; i < 10; ++i)
285 #pragma omp for simd schedule(guided)
286 for (int i = 0; i < 10; ++i)
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 {{.+}})
296 #pragma omp parallel for
297 for (int i = 0; i < 10; ++i)
300 #pragma omp parallel for schedule(static)
301 for (int i = 0; i < 10; ++i)
304 #pragma omp parallel for schedule(static, 1)
305 for (int i = 0; i < 10; ++i)
308 #pragma omp parallel for schedule(auto)
309 for (int i = 0; i < 10; ++i)
312 #pragma omp parallel for schedule(runtime)
313 for (int i = 0; i < 10; ++i)
316 #pragma omp parallel for schedule(dynamic)
317 for (int i = 0; i < 10; ++i)
320 #pragma omp parallel for schedule(guided)
321 for (int i = 0; i < 10; ++i)