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
8 #ifndef HEADER
9 #define HEADER
10
11 // CHECK-NOT: @__omp_offloading_{{.+}}_exec_mode = weak constant i8 1
12
foo()13 void foo() {
14 // CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
15 // CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
16 // CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
17 // CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
18 // CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
19 // CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
20 // CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
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 i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
44 // CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
45 // CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
46 // CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
47 // CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
48 // CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
49 // CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
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 i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
72 // CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
73 // CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
74 // CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
75 // CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
76 // CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
77 // CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
78 #pragma omp target teams
79 #pragma omp distribute parallel for simd
80 for (int i = 0; i < 10; ++i)
81 ;
82 #pragma omp target teams
83 #pragma omp distribute parallel for simd schedule(static)
84 for (int i = 0; i < 10; ++i)
85 ;
86 #pragma omp target teams
87 #pragma omp distribute parallel for simd schedule(static, 1)
88 for (int i = 0; i < 10; ++i)
89 ;
90 #pragma omp target teams
91 #pragma omp distribute parallel for simd schedule(auto)
92 for (int i = 0; i < 10; ++i)
93 ;
94 #pragma omp target teams
95 #pragma omp distribute parallel for simd schedule(runtime)
96 for (int i = 0; i < 10; ++i)
97 ;
98 #pragma omp target teams
99 #pragma omp distribute parallel for simd schedule(dynamic)
100 for (int i = 0; i < 10; ++i)
101 ;
102 #pragma omp target teams
103 #pragma omp distribute parallel for simd schedule(guided)
104 for (int i = 0; i < 10; ++i)
105 ;
106 // CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
107 // CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
108 // CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
109 // CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
110 // CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
111 // CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
112 // CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
113 #pragma omp target teams
114 #pragma omp distribute parallel for
115 for (int i = 0; i < 10; ++i)
116 ;
117 #pragma omp target teams
118 #pragma omp distribute parallel for schedule(static)
119 for (int i = 0; i < 10; ++i)
120 ;
121 #pragma omp target teams
122 #pragma omp distribute parallel for schedule(static, 1)
123 for (int i = 0; i < 10; ++i)
124 ;
125 #pragma omp target teams
126 #pragma omp distribute parallel for schedule(auto)
127 for (int i = 0; i < 10; ++i)
128 ;
129 #pragma omp target teams
130 #pragma omp distribute parallel for schedule(runtime)
131 for (int i = 0; i < 10; ++i)
132 ;
133 #pragma omp target teams
134 #pragma omp distribute parallel for schedule(dynamic)
135 for (int i = 0; i < 10; ++i)
136 ;
137 #pragma omp target teams
138 #pragma omp distribute parallel for schedule(guided)
139 for (int i = 0; i < 10; ++i)
140 ;
141 // CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
142 // CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
143 // CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
144 // CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
145 // CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
146 // CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
147 // CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
148 #pragma omp target
149 #pragma omp teams
150 #pragma omp distribute parallel for
151 for (int i = 0; i < 10; ++i)
152 ;
153 #pragma omp target
154 #pragma omp teams
155 #pragma omp distribute parallel for schedule(static)
156 for (int i = 0; i < 10; ++i)
157 ;
158 #pragma omp target
159 #pragma omp teams
160 #pragma omp distribute parallel for schedule(static, 1)
161 for (int i = 0; i < 10; ++i)
162 ;
163 #pragma omp target
164 #pragma omp teams
165 #pragma omp distribute parallel for schedule(auto)
166 for (int i = 0; i < 10; ++i)
167 ;
168 #pragma omp target
169 #pragma omp teams
170 #pragma omp distribute parallel for schedule(runtime)
171 for (int i = 0; i < 10; ++i)
172 ;
173 #pragma omp target
174 #pragma omp teams
175 #pragma omp distribute parallel for schedule(dynamic)
176 for (int i = 0; i < 10; ++i)
177 ;
178 #pragma omp target
179 #pragma omp teams
180 #pragma omp distribute parallel for schedule(guided)
181 for (int i = 0; i < 10; ++i)
182 ;
183 // CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
184 // CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
185 // CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
186 // CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
187 // CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
188 // CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
189 // CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
190 #pragma omp target parallel for
191 for (int i = 0; i < 10; ++i)
192 ;
193 #pragma omp target parallel for schedule(static)
194 for (int i = 0; i < 10; ++i)
195 ;
196 #pragma omp target parallel for schedule(static, 1)
197 for (int i = 0; i < 10; ++i)
198 ;
199 #pragma omp target parallel for schedule(auto)
200 for (int i = 0; i < 10; ++i)
201 ;
202 #pragma omp target parallel for schedule(runtime)
203 for (int i = 0; i < 10; ++i)
204 ;
205 #pragma omp target parallel for schedule(dynamic)
206 for (int i = 0; i < 10; ++i)
207 ;
208 #pragma omp target parallel for schedule(guided)
209 for (int i = 0; i < 10; ++i)
210 ;
211 // CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
212 // CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
213 // CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
214 // CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
215 // CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
216 // CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
217 // CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
218 #pragma omp target parallel
219 #pragma omp for simd
220 for (int i = 0; i < 10; ++i)
221 ;
222 #pragma omp target parallel
223 #pragma omp for simd schedule(static)
224 for (int i = 0; i < 10; ++i)
225 ;
226 #pragma omp target parallel
227 #pragma omp for simd schedule(static, 1)
228 for (int i = 0; i < 10; ++i)
229 ;
230 #pragma omp target parallel
231 #pragma omp for simd schedule(auto)
232 for (int i = 0; i < 10; ++i)
233 ;
234 #pragma omp target parallel
235 #pragma omp for simd schedule(runtime)
236 for (int i = 0; i < 10; ++i)
237 ;
238 #pragma omp target parallel
239 #pragma omp for simd schedule(dynamic)
240 for (int i = 0; i < 10; ++i)
241 ;
242 #pragma omp target parallel
243 #pragma omp for simd schedule(guided)
244 for (int i = 0; i < 10; ++i)
245 ;
246 // CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
247 // CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
248 // CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
249 // CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
250 // CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
251 // CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
252 // CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
253 #pragma omp target
254 #pragma omp parallel
255 #pragma omp for simd ordered
256 for (int i = 0; i < 10; ++i)
257 ;
258 #pragma omp target
259 #pragma omp parallel
260 #pragma omp for simd schedule(static)
261 for (int i = 0; i < 10; ++i)
262 ;
263 #pragma omp target
264 #pragma omp parallel
265 #pragma omp for simd schedule(static, 1)
266 for (int i = 0; i < 10; ++i)
267 ;
268 #pragma omp target
269 #pragma omp parallel
270 #pragma omp for simd schedule(auto)
271 for (int i = 0; i < 10; ++i)
272 ;
273 #pragma omp target
274 #pragma omp parallel
275 #pragma omp for simd schedule(runtime)
276 for (int i = 0; i < 10; ++i)
277 ;
278 #pragma omp target
279 #pragma omp parallel
280 #pragma omp for simd schedule(dynamic)
281 for (int i = 0; i < 10; ++i)
282 ;
283 #pragma omp target
284 #pragma omp parallel
285 #pragma omp for simd schedule(guided)
286 for (int i = 0; i < 10; ++i)
287 ;
288 // CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
289 // CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
290 // CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
291 // CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
292 // CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
293 // CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
294 // CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
295 #pragma omp target
296 #pragma omp parallel for
297 for (int i = 0; i < 10; ++i)
298 ;
299 #pragma omp target
300 #pragma omp parallel for schedule(static)
301 for (int i = 0; i < 10; ++i)
302 ;
303 #pragma omp target
304 #pragma omp parallel for schedule(static, 1)
305 for (int i = 0; i < 10; ++i)
306 ;
307 #pragma omp target
308 #pragma omp parallel for schedule(auto)
309 for (int i = 0; i < 10; ++i)
310 ;
311 #pragma omp target
312 #pragma omp parallel for schedule(runtime)
313 for (int i = 0; i < 10; ++i)
314 ;
315 #pragma omp target
316 #pragma omp parallel for schedule(dynamic)
317 for (int i = 0; i < 10; ++i)
318 ;
319 #pragma omp target
320 #pragma omp parallel for schedule(guided)
321 for (int i = 0; i < 10; ++i)
322 ;
323 }
324
325 #endif
326
327