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 --check-prefix CHECK --check-prefix CHECK-64
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 --check-prefix CHECK --check-prefix CHECK-32
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 --check-prefix CHECK --check-prefix CHECK-32
7 // expected-no-diagnostics
8 #ifndef HEADER
9 #define HEADER
10 
11 template<typename tx>
12 tx ftemplate(int n) {
13   tx a = 0;
14   short aa = 0;
15   tx b[10];
16 
17   #pragma omp target if(0)
18   {
19     #pragma omp parallel
20     {
21       int a = 41;
22     }
23     a += 1;
24   }
25 
26   #pragma omp target
27   {
28     #pragma omp parallel
29     {
30       int a = 42;
31     }
32     #pragma omp parallel if(0)
33     {
34       int a = 43;
35     }
36     #pragma omp parallel if(1)
37     {
38       int a = 44;
39     }
40     a += 1;
41   }
42 
43   #pragma omp target if(n>40)
44   {
45     #pragma omp parallel if(n>1000)
46     {
47       int a = 45;
48     }
49     a += 1;
50     aa += 1;
51     b[2] += 1;
52   }
53 
54   #pragma omp target
55   {
56     #pragma omp parallel
57     {
58     #pragma omp critical
59     ++a;
60     }
61   }
62   return a;
63 }
64 
65 int bar(int n){
66   int a = 0;
67 
68   a += ftemplate<int>(n);
69 
70   return a;
71 }
72 
73 // CHECK-NOT: define {{.*}}void {{@__omp_offloading_.+template.+l17}}_worker()
74 
75 // CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+template.+l26}}_worker()
76 // CHECK-DAG: [[OMP_EXEC_STATUS:%.+]] = alloca i8,
77 // CHECK-DAG: [[OMP_WORK_FN:%.+]] = alloca i8*,
78 // CHECK: store i8* null, i8** [[OMP_WORK_FN]],
79 // CHECK: store i8 0, i8* [[OMP_EXEC_STATUS]],
80 // CHECK: br label {{%?}}[[AWAIT_WORK:.+]]
81 //
82 // CHECK: [[AWAIT_WORK]]
83 // CHECK: call void @llvm.nvvm.barrier0()
84 // CHECK: [[KPR:%.+]] = call i1 @__kmpc_kernel_parallel(i8** [[OMP_WORK_FN]]
85 // CHECK: [[KPRB:%.+]] = zext i1 [[KPR]] to i8
86 // store i8 [[KPRB]], i8* [[OMP_EXEC_STATUS]], align 1
87 // CHECK: [[WORK:%.+]] = load i8*, i8** [[OMP_WORK_FN]],
88 // CHECK: [[SHOULD_EXIT:%.+]] = icmp eq i8* [[WORK]], null
89 // CHECK: br i1 [[SHOULD_EXIT]], label {{%?}}[[EXIT:.+]], label {{%?}}[[SEL_WORKERS:.+]]
90 //
91 // CHECK: [[SEL_WORKERS]]
92 // CHECK: [[ST:%.+]] = load i8, i8* [[OMP_EXEC_STATUS]]
93 // CHECK: [[IS_ACTIVE:%.+]] = icmp ne i8 [[ST]], 0
94 // CHECK: br i1 [[IS_ACTIVE]], label {{%?}}[[EXEC_PARALLEL:.+]], label {{%?}}[[BAR_PARALLEL:.+]]
95 //
96 // CHECK: [[EXEC_PARALLEL]]
97 // CHECK: [[WF1:%.+]] = load i8*, i8** [[OMP_WORK_FN]],
98 // CHECK: [[WM1:%.+]] = icmp eq i8* [[WF1]], bitcast (void (i16, i32)* [[PARALLEL_FN1:@.+]]_wrapper to i8*)
99 // CHECK: br i1 [[WM1]], label {{%?}}[[EXEC_PFN1:.+]], label {{%?}}[[CHECK_NEXT1:.+]]
100 //
101 // CHECK: [[EXEC_PFN1]]
102 // CHECK: call void [[PARALLEL_FN1]]_wrapper(
103 // CHECK: br label {{%?}}[[TERM_PARALLEL:.+]]
104 //
105 // CHECK: [[CHECK_NEXT1]]
106 // CHECK: [[WF2:%.+]] = load i8*, i8** [[OMP_WORK_FN]],
107 // CHECK: [[WM2:%.+]] = icmp eq i8* [[WF2]], bitcast (void (i16, i32)* [[PARALLEL_FN2:@.+]]_wrapper to i8*)
108 // CHECK: br i1 [[WM2]], label {{%?}}[[EXEC_PFN2:.+]], label {{%?}}[[CHECK_NEXT2:.+]]
109 //
110 // CHECK: [[EXEC_PFN2]]
111 // CHECK: call void [[PARALLEL_FN2]]_wrapper(
112 // CHECK: br label {{%?}}[[TERM_PARALLEL:.+]]
113 //
114 // CHECK: [[CHECK_NEXT2]]
115 // CHECK: br label {{%?}}[[TERM_PARALLEL:.+]]
116 //
117 // CHECK: [[TERM_PARALLEL]]
118 // CHECK: call void @__kmpc_kernel_end_parallel()
119 // CHECK: br label {{%?}}[[BAR_PARALLEL]]
120 //
121 // CHECK: [[BAR_PARALLEL]]
122 // CHECK: call void @llvm.nvvm.barrier0()
123 // CHECK: br label {{%?}}[[AWAIT_WORK]]
124 //
125 // CHECK: [[EXIT]]
126 // CHECK: ret void
127 
128 // CHECK: define {{.*}}void [[T6:@__omp_offloading_.+template.+l26]](i[[SZ:32|64]]
129 // Create local storage for each capture.
130 // CHECK:  [[LOCAL_A:%.+]] = alloca i[[SZ]],
131 // CHECK-DAG:  store i[[SZ]] [[ARG_A:%.+]], i[[SZ]]* [[LOCAL_A]]
132 // Store captures in the context.
133 // CHECK-64-DAG:[[REF_A:%.+]] = bitcast i[[SZ]]* [[LOCAL_A]] to i32*
134 //
135 // CHECK-DAG: [[TID:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
136 // CHECK-DAG: [[NTH:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x()
137 // CHECK-DAG: [[WS:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize()
138 // CHECK-DAG: [[TH_LIMIT:%.+]] = sub nuw i32 [[NTH]], [[WS]]
139 // CHECK: [[IS_WORKER:%.+]] = icmp ult i32 [[TID]], [[TH_LIMIT]]
140 // CHECK: br i1 [[IS_WORKER]], label {{%?}}[[WORKER:.+]], label {{%?}}[[CHECK_MASTER:.+]]
141 //
142 // CHECK: [[WORKER]]
143 // CHECK: {{call|invoke}} void [[T6]]_worker()
144 // CHECK: br label {{%?}}[[EXIT:.+]]
145 //
146 // CHECK: [[CHECK_MASTER]]
147 // CHECK-DAG: [[CMTID:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
148 // CHECK-DAG: [[CMNTH:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x()
149 // CHECK-DAG: [[CMWS:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize()
150 // CHECK: [[IS_MASTER:%.+]] = icmp eq i32 [[CMTID]],
151 // CHECK: br i1 [[IS_MASTER]], label {{%?}}[[MASTER:.+]], label {{%?}}[[EXIT]]
152 //
153 // CHECK: [[MASTER]]
154 // CHECK-DAG: [[MNTH:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x()
155 // CHECK-DAG: [[MWS:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize()
156 // CHECK: [[MTMP1:%.+]] = sub nuw i32 [[MNTH]], [[MWS]]
157 // CHECK: call void @__kmpc_kernel_init(i32 [[MTMP1]]
158 // CHECK: call void @__kmpc_kernel_prepare_parallel(i8* bitcast (void (i16, i32)* [[PARALLEL_FN1]]_wrapper to i8*),
159 // CHECK: call void @llvm.nvvm.barrier0()
160 // CHECK: call void @llvm.nvvm.barrier0()
161 // CHECK: call void @__kmpc_serialized_parallel(
162 // CHECK: {{call|invoke}} void [[PARALLEL_FN3:@.+]](
163 // CHECK: call void @__kmpc_end_serialized_parallel(
164 // CHECK: call void @__kmpc_kernel_prepare_parallel(i8* bitcast (void (i16, i32)* [[PARALLEL_FN2]]_wrapper to i8*),
165 // CHECK: call void @llvm.nvvm.barrier0()
166 // CHECK: call void @llvm.nvvm.barrier0()
167 // CHECK-64-DAG: load i32, i32* [[REF_A]]
168 // CHECK-32-DAG: load i32, i32* [[LOCAL_A]]
169 // CHECK: br label {{%?}}[[TERMINATE:.+]]
170 //
171 // CHECK: [[TERMINATE]]
172 // CHECK: call void @__kmpc_kernel_deinit(
173 // CHECK: call void @llvm.nvvm.barrier0()
174 // CHECK: br label {{%?}}[[EXIT]]
175 //
176 // CHECK: [[EXIT]]
177 // CHECK: ret void
178 
179 // CHECK-DAG: define internal void [[PARALLEL_FN1]](
180 // CHECK: [[A:%.+]] = alloca i[[SZ:32|64]],
181 // CHECK: store i[[SZ]] 42, i[[SZ]]* %a,
182 // CHECK: ret void
183 
184 // CHECK-DAG: define internal void [[PARALLEL_FN3]](
185 // CHECK: [[A:%.+]] = alloca i[[SZ:32|64]],
186 // CHECK: store i[[SZ]] 43, i[[SZ]]* %a,
187 // CHECK: ret void
188 
189 // CHECK-DAG: define internal void [[PARALLEL_FN2]](
190 // CHECK: [[A:%.+]] = alloca i[[SZ:32|64]],
191 // CHECK: store i[[SZ]] 44, i[[SZ]]* %a,
192 // CHECK: ret void
193 
194 // CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+template.+l43}}_worker()
195 // CHECK-DAG: [[OMP_EXEC_STATUS:%.+]] = alloca i8,
196 // CHECK-DAG: [[OMP_WORK_FN:%.+]] = alloca i8*,
197 // CHECK: store i8* null, i8** [[OMP_WORK_FN]],
198 // CHECK: store i8 0, i8* [[OMP_EXEC_STATUS]],
199 // CHECK: br label {{%?}}[[AWAIT_WORK:.+]]
200 //
201 // CHECK: [[AWAIT_WORK]]
202 // CHECK: call void @llvm.nvvm.barrier0()
203 // CHECK: [[KPR:%.+]] = call i1 @__kmpc_kernel_parallel(i8** [[OMP_WORK_FN]],
204 // CHECK: [[KPRB:%.+]] = zext i1 [[KPR]] to i8
205 // store i8 [[KPRB]], i8* [[OMP_EXEC_STATUS]], align 1
206 // CHECK: [[WORK:%.+]] = load i8*, i8** [[OMP_WORK_FN]],
207 // CHECK: [[SHOULD_EXIT:%.+]] = icmp eq i8* [[WORK]], null
208 // CHECK: br i1 [[SHOULD_EXIT]], label {{%?}}[[EXIT:.+]], label {{%?}}[[SEL_WORKERS:.+]]
209 //
210 // CHECK: [[SEL_WORKERS]]
211 // CHECK: [[ST:%.+]] = load i8, i8* [[OMP_EXEC_STATUS]]
212 // CHECK: [[IS_ACTIVE:%.+]] = icmp ne i8 [[ST]], 0
213 // CHECK: br i1 [[IS_ACTIVE]], label {{%?}}[[EXEC_PARALLEL:.+]], label {{%?}}[[BAR_PARALLEL:.+]]
214 //
215 // CHECK: [[EXEC_PARALLEL]]
216 // CHECK: [[WF:%.+]] = load i8*, i8** [[OMP_WORK_FN]],
217 // CHECK: [[WM:%.+]] = icmp eq i8* [[WF]], bitcast (void (i16, i32)* [[PARALLEL_FN4:@.+]]_wrapper to i8*)
218 // CHECK: br i1 [[WM]], label {{%?}}[[EXEC_PFN:.+]], label {{%?}}[[CHECK_NEXT:.+]]
219 //
220 // CHECK: [[EXEC_PFN]]
221 // CHECK: call void [[PARALLEL_FN4]]_wrapper(
222 // CHECK: br label {{%?}}[[TERM_PARALLEL:.+]]
223 //
224 // CHECK: [[CHECK_NEXT]]
225 // CHECK: br label {{%?}}[[TERM_PARALLEL:.+]]
226 //
227 // CHECK: [[TERM_PARALLEL]]
228 // CHECK: call void @__kmpc_kernel_end_parallel()
229 // CHECK: br label {{%?}}[[BAR_PARALLEL]]
230 //
231 // CHECK: [[BAR_PARALLEL]]
232 // CHECK: call void @llvm.nvvm.barrier0()
233 // CHECK: br label {{%?}}[[AWAIT_WORK]]
234 //
235 // CHECK: [[EXIT]]
236 // CHECK: ret void
237 
238 // CHECK: define {{.*}}void [[T6:@__omp_offloading_.+template.+l43]](i[[SZ:32|64]]
239 // Create local storage for each capture.
240 // CHECK:  [[LOCAL_N:%.+]] = alloca i[[SZ]],
241 // CHECK:  [[LOCAL_A:%.+]] = alloca i[[SZ]],
242 // CHECK:  [[LOCAL_AA:%.+]] = alloca i[[SZ]],
243 // CHECK:  [[LOCAL_B:%.+]] = alloca [10 x i32]*
244 // CHECK-DAG:  store i[[SZ]] [[ARG_N:%.+]], i[[SZ]]* [[LOCAL_N]]
245 // CHECK-DAG:  store i[[SZ]] [[ARG_A:%.+]], i[[SZ]]* [[LOCAL_A]]
246 // CHECK-DAG:  store i[[SZ]] [[ARG_AA:%.+]], i[[SZ]]* [[LOCAL_AA]]
247 // CHECK-DAG:   store [10 x i32]* [[ARG_B:%.+]], [10 x i32]** [[LOCAL_B]]
248 // Store captures in the context.
249 // CHECK-64-DAG:[[REF_N:%.+]] = bitcast i[[SZ]]* [[LOCAL_N]] to i32*
250 // CHECK-64-DAG:[[REF_A:%.+]] = bitcast i[[SZ]]* [[LOCAL_A]] to i32*
251 // CHECK-DAG:   [[REF_AA:%.+]] = bitcast i[[SZ]]* [[LOCAL_AA]] to i16*
252 // CHECK-DAG:   [[REF_B:%.+]] = load [10 x i32]*, [10 x i32]** [[LOCAL_B]],
253 //
254 // CHECK-DAG: [[TID:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
255 // CHECK-DAG: [[NTH:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x()
256 // CHECK-DAG: [[WS:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize()
257 // CHECK-DAG: [[TH_LIMIT:%.+]] = sub nuw i32 [[NTH]], [[WS]]
258 // CHECK: [[IS_WORKER:%.+]] = icmp ult i32 [[TID]], [[TH_LIMIT]]
259 // CHECK: br i1 [[IS_WORKER]], label {{%?}}[[WORKER:.+]], label {{%?}}[[CHECK_MASTER:.+]]
260 //
261 // CHECK: [[WORKER]]
262 // CHECK: {{call|invoke}} void [[T6]]_worker()
263 // CHECK: br label {{%?}}[[EXIT:.+]]
264 //
265 // CHECK: [[CHECK_MASTER]]
266 // CHECK-DAG: [[CMTID:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
267 // CHECK-DAG: [[CMNTH:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x()
268 // CHECK-DAG: [[CMWS:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize()
269 // CHECK: [[IS_MASTER:%.+]] = icmp eq i32 [[CMTID]],
270 // CHECK: br i1 [[IS_MASTER]], label {{%?}}[[MASTER:.+]], label {{%?}}[[EXIT]]
271 //
272 // CHECK: [[MASTER]]
273 // CHECK-DAG: [[MNTH:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x()
274 // CHECK-DAG: [[MWS:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize()
275 // CHECK: [[MTMP1:%.+]] = sub nuw i32 [[MNTH]], [[MWS]]
276 // CHECK: call void @__kmpc_kernel_init(i32 [[MTMP1]]
277 // CHECK-64: [[N:%.+]] = load i32, i32* [[REF_N]],
278 // CHECK-32: [[N:%.+]] = load i32, i32* [[LOCAL_N]],
279 // CHECK: [[CMP:%.+]] = icmp sgt i32 [[N]], 1000
280 // CHECK: br i1 [[CMP]], label {{%?}}[[IF_THEN:.+]], label {{%?}}[[IF_ELSE:.+]]
281 //
282 // CHECK: [[IF_THEN]]
283 // CHECK: call void @__kmpc_kernel_prepare_parallel(i8* bitcast (void (i16, i32)* [[PARALLEL_FN4]]_wrapper to i8*),
284 // CHECK: call void @llvm.nvvm.barrier0()
285 // CHECK: call void @llvm.nvvm.barrier0()
286 // CHECK: br label {{%?}}[[IF_END:.+]]
287 //
288 // CHECK: [[IF_ELSE]]
289 // CHECK: call void @__kmpc_serialized_parallel(
290 // CHECK: {{call|invoke}} void [[PARALLEL_FN4]](
291 // CHECK: call void @__kmpc_end_serialized_parallel(
292 // br label [[IF_END]]
293 //
294 // CHECK: [[IF_END]]
295 // CHECK-64-DAG: load i32, i32* [[REF_A]]
296 // CHECK-32-DAG: load i32, i32* [[LOCAL_A]]
297 // CHECK-DAG:    load i16, i16* [[REF_AA]]
298 // CHECK-DAG:    getelementptr inbounds [10 x i32], [10 x i32]* [[REF_B]], i[[SZ]] 0, i[[SZ]] 2
299 //
300 // CHECK: br label {{%?}}[[TERMINATE:.+]]
301 //
302 // CHECK: [[TERMINATE]]
303 // CHECK: call void @__kmpc_kernel_deinit(
304 // CHECK: call void @llvm.nvvm.barrier0()
305 // CHECK: br label {{%?}}[[EXIT]]
306 //
307 // CHECK: [[EXIT]]
308 // CHECK: ret void
309 
310 // CHECK: define internal void [[PARALLEL_FN4]](
311 // CHECK: [[A:%.+]] = alloca i[[SZ:32|64]],
312 // CHECK: store i[[SZ]] 45, i[[SZ]]* %a,
313 // CHECK: ret void
314 
315 // CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+template.+l54}}_worker()
316 // CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+template.+l54}}(
317 
318 // CHECK-LABEL: define internal void @{{.+}}(i32* noalias %{{.+}}, i32* noalias %{{.+}}, i32* dereferenceable{{.*}})
319 // CHECK:  [[CC:%.+]] = alloca i32,
320 // CHECK:  [[TID:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
321 // CHECK:  [[NUM_THREADS:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x()
322 // CHECK:  store i32 0, i32* [[CC]],
323 // CHECK:  br label
324 
325 // CHECK:  [[CC_VAL:%.+]] = load i32, i32* [[CC]],
326 // CHECK:  [[RES:%.+]] = icmp slt i32 [[CC_VAL]], [[NUM_THREADS]]
327 // CHECK:  br i1 [[RES]], label
328 
329 // CHECK:  [[CC_VAL:%.+]] = load i32, i32* [[CC]],
330 // CHECK:  [[RES:%.+]] = icmp eq i32 [[TID]], [[CC_VAL]]
331 // CHECK:  br i1 [[RES]], label
332 
333 // CHECK:  call void @llvm.nvvm.barrier0()
334 // CHECK:  [[NEW_CC_VAL:%.+]] = add nsw i32 [[CC_VAL]], 1
335 // CHECK:  store i32 [[NEW_CC_VAL]], i32* [[CC]],
336 // CHECK:  br label
337 
338 #endif
339