1 // NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --function-signature --include-generated-funcs --replace-value-regex "__omp_offloading_[0-9a-z]+_[0-9a-z]+" "reduction_size[.].+[.]" "pl_cond[.].+[.|,]" --prefix-filecheck-ir-name _
2 // Test target codegen - host bc file has to be created first.
3 // 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
4 // 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=CHECK1
5 // 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
6 // 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=CHECK2
7 // 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=CHECK3
8 // expected-no-diagnostics
9 #ifndef HEADER
10 #define HEADER
11
12 template<typename tx>
ftemplate(int n)13 tx ftemplate(int n) {
14 int i;
15
16 #pragma omp target teams distribute
17 for (i = 0; i < 10; ++i)
18 {
19 #pragma omp parallel
20 ++i;
21 }
22
23 return i;
24 }
25
bar(int n)26 int bar(int n){
27 int a = 0;
28
29 a += ftemplate<char>(n);
30
31 return a;
32 }
33
34 #endif
35 // CHECK4-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l19_worker
36 // CHECK4-SAME: () #[[ATTR0:[0-9]+]] {
37 // CHECK4-NEXT: entry:
38 // CHECK4-NEXT: [[WORK_FN:%.*]] = alloca i8*, align 4
39 // CHECK4-NEXT: [[EXEC_STATUS:%.*]] = alloca i8, align 1
40 // CHECK4-NEXT: store i8* null, i8** [[WORK_FN]], align 4
41 // CHECK4-NEXT: store i8 0, i8* [[EXEC_STATUS]], align 1
42 // CHECK4-NEXT: br label [[DOTAWAIT_WORK:%.*]]
43 // CHECK4: .await.work:
44 // CHECK4-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0)
45 // CHECK4-NEXT: [[TMP0:%.*]] = call i1 @__kmpc_kernel_parallel(i8** [[WORK_FN]])
46 // CHECK4-NEXT: [[TMP1:%.*]] = zext i1 [[TMP0]] to i8
47 // CHECK4-NEXT: store i8 [[TMP1]], i8* [[EXEC_STATUS]], align 1
48 // CHECK4-NEXT: [[TMP2:%.*]] = load i8*, i8** [[WORK_FN]], align 4
49 // CHECK4-NEXT: [[SHOULD_TERMINATE:%.*]] = icmp eq i8* [[TMP2]], null
50 // CHECK4-NEXT: br i1 [[SHOULD_TERMINATE]], label [[DOTEXIT:%.*]], label [[DOTSELECT_WORKERS:%.*]]
51 // CHECK4: .select.workers:
52 // CHECK4-NEXT: [[TMP3:%.*]] = load i8, i8* [[EXEC_STATUS]], align 1
53 // CHECK4-NEXT: [[IS_ACTIVE:%.*]] = icmp ne i8 [[TMP3]], 0
54 // CHECK4-NEXT: br i1 [[IS_ACTIVE]], label [[DOTEXECUTE_PARALLEL:%.*]], label [[DOTBARRIER_PARALLEL:%.*]]
55 // CHECK4: .execute.parallel:
56 // CHECK4-NEXT: [[TMP4:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB2:[0-9]+]])
57 // CHECK4-NEXT: [[TMP5:%.*]] = load i8*, i8** [[WORK_FN]], align 4
58 // CHECK4-NEXT: [[WORK_MATCH:%.*]] = icmp eq i8* [[TMP5]], bitcast (void (i16, i32)* @__omp_outlined__1_wrapper to i8*)
59 // CHECK4-NEXT: br i1 [[WORK_MATCH]], label [[DOTEXECUTE_FN:%.*]], label [[DOTCHECK_NEXT:%.*]]
60 // CHECK4: .execute.fn:
61 // CHECK4-NEXT: call void @__omp_outlined__1_wrapper(i16 0, i32 [[TMP4]]) #[[ATTR3:[0-9]+]]
62 // CHECK4-NEXT: br label [[DOTTERMINATE_PARALLEL:%.*]]
63 // CHECK4: .check.next:
64 // CHECK4-NEXT: [[TMP6:%.*]] = bitcast i8* [[TMP2]] to void (i16, i32)*
65 // CHECK4-NEXT: call void [[TMP6]](i16 0, i32 [[TMP4]])
66 // CHECK4-NEXT: br label [[DOTTERMINATE_PARALLEL]]
67 // CHECK4: .terminate.parallel:
68 // CHECK4-NEXT: call void @__kmpc_kernel_end_parallel()
69 // CHECK4-NEXT: br label [[DOTBARRIER_PARALLEL]]
70 // CHECK4: .barrier.parallel:
71 // CHECK4-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0)
72 // CHECK4-NEXT: br label [[DOTAWAIT_WORK]]
73 // CHECK4: .exit:
74 // CHECK4-NEXT: ret void
75 // CHECK4-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l19
76 // CHECK4-SAME: () #[[ATTR1:[0-9]+]] {
77 // CHECK4-NEXT: entry:
78 // CHECK4-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4
79 // CHECK4-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 4
80 // CHECK4-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4
81 // CHECK4-NEXT: [[NVPTX_TID:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
82 // CHECK4-NEXT: [[NVPTX_NUM_THREADS:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x()
83 // CHECK4-NEXT: [[NVPTX_WARP_SIZE:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize()
84 // CHECK4-NEXT: [[THREAD_LIMIT:%.*]] = sub nuw i32 [[NVPTX_NUM_THREADS]], [[NVPTX_WARP_SIZE]]
85 // CHECK4-NEXT: [[TMP0:%.*]] = icmp ult i32 [[NVPTX_TID]], [[THREAD_LIMIT]]
86 // CHECK4-NEXT: br i1 [[TMP0]], label [[DOTWORKER:%.*]], label [[DOTMASTERCHECK:%.*]]
87 // CHECK4: .worker:
88 // CHECK4-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l19_worker() #[[ATTR3]]
89 // CHECK4-NEXT: br label [[DOTEXIT:%.*]]
90 // CHECK4: .mastercheck:
91 // CHECK4-NEXT: [[NVPTX_TID1:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
92 // CHECK4-NEXT: [[NVPTX_NUM_THREADS2:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x()
93 // CHECK4-NEXT: [[NVPTX_WARP_SIZE3:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize()
94 // CHECK4-NEXT: [[TMP1:%.*]] = sub nuw i32 [[NVPTX_WARP_SIZE3]], 1
95 // CHECK4-NEXT: [[TMP2:%.*]] = sub nuw i32 [[NVPTX_NUM_THREADS2]], 1
96 // CHECK4-NEXT: [[TMP3:%.*]] = xor i32 [[TMP1]], -1
97 // CHECK4-NEXT: [[MASTER_TID:%.*]] = and i32 [[TMP2]], [[TMP3]]
98 // CHECK4-NEXT: [[TMP4:%.*]] = icmp eq i32 [[NVPTX_TID1]], [[MASTER_TID]]
99 // CHECK4-NEXT: br i1 [[TMP4]], label [[DOTMASTER:%.*]], label [[DOTEXIT]]
100 // CHECK4: .master:
101 // CHECK4-NEXT: [[NVPTX_NUM_THREADS4:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x()
102 // CHECK4-NEXT: [[NVPTX_WARP_SIZE5:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize()
103 // CHECK4-NEXT: [[THREAD_LIMIT6:%.*]] = sub nuw i32 [[NVPTX_NUM_THREADS4]], [[NVPTX_WARP_SIZE5]]
104 // CHECK4-NEXT: call void @__kmpc_kernel_init(i32 [[THREAD_LIMIT6]], i16 1)
105 // CHECK4-NEXT: call void @__kmpc_data_sharing_init_stack()
106 // CHECK4-NEXT: [[TMP5:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB2]])
107 // CHECK4-NEXT: store i32 [[TMP5]], i32* [[DOTTHREADID_TEMP_]], align 4
108 // CHECK4-NEXT: call void @__omp_outlined__(i32* [[DOTTHREADID_TEMP_]], i32* [[DOTZERO_ADDR]]) #[[ATTR3]]
109 // CHECK4-NEXT: br label [[DOTTERMINATION_NOTIFIER:%.*]]
110 // CHECK4: .termination.notifier:
111 // CHECK4-NEXT: call void @__kmpc_kernel_deinit(i16 1)
112 // CHECK4-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0)
113 // CHECK4-NEXT: br label [[DOTEXIT]]
114 // CHECK4: .exit:
115 // CHECK4-NEXT: ret void
116 // CHECK4-LABEL: define {{[^@]+}}@__omp_outlined__
117 // CHECK4-SAME: (i32* noalias [[DOTGLOBAL_TID_:%.*]], i32* noalias [[DOTBOUND_TID_:%.*]]) #[[ATTR1]] {
118 // CHECK4-NEXT: entry:
119 // CHECK4-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 4
120 // CHECK4-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 4
121 // CHECK4-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 4
122 // CHECK4-NEXT: [[TMP:%.*]] = alloca i32, align 4
123 // CHECK4-NEXT: [[DOTOMP_LB:%.*]] = alloca i32, align 4
124 // CHECK4-NEXT: [[DOTOMP_UB:%.*]] = alloca i32, align 4
125 // CHECK4-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4
126 // CHECK4-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4
127 // CHECK4-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [1 x i8*], align 4
128 // CHECK4-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 4
129 // CHECK4-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 4
130 // CHECK4-NEXT: [[TMP0:%.*]] = load i16, i16* @"_openmp_static_kernel$is_shared", align 2
131 // CHECK4-NEXT: [[TMP1:%.*]] = load i32, i32* @"_openmp_static_kernel$size", align 4
132 // CHECK4-NEXT: call void @__kmpc_get_team_static_memory(i16 0, i8* addrspacecast (i8 addrspace(3)* getelementptr inbounds (%"union._shared_openmp_static_memory_type_$_", %"union._shared_openmp_static_memory_type_$_" addrspace(3)* @"_openmp_shared_static_glob_rd_$_", i32 0, i32 0, i32 0) to i8*), i32 [[TMP1]], i16 [[TMP0]], i8** addrspacecast (i8* addrspace(3)* @"_openmp_kernel_static_glob_rd$ptr" to i8**))
133 // CHECK4-NEXT: [[TMP2:%.*]] = load i8*, i8* addrspace(3)* @"_openmp_kernel_static_glob_rd$ptr", align 4
134 // CHECK4-NEXT: [[TMP3:%.*]] = getelementptr inbounds i8, i8* [[TMP2]], i32 0
135 // CHECK4-NEXT: [[TMP4:%.*]] = bitcast i8* [[TMP3]] to %struct._globalized_locals_ty*
136 // CHECK4-NEXT: [[I:%.*]] = getelementptr inbounds [[STRUCT__GLOBALIZED_LOCALS_TY:%.*]], %struct._globalized_locals_ty* [[TMP4]], i32 0, i32 0
137 // CHECK4-NEXT: store i32 0, i32* [[DOTOMP_LB]], align 4
138 // CHECK4-NEXT: store i32 9, i32* [[DOTOMP_UB]], align 4
139 // CHECK4-NEXT: store i32 1, i32* [[DOTOMP_STRIDE]], align 4
140 // CHECK4-NEXT: store i32 0, i32* [[DOTOMP_IS_LAST]], align 4
141 // CHECK4-NEXT: [[TMP5:%.*]] = load i32*, i32** [[DOTGLOBAL_TID__ADDR]], align 4
142 // CHECK4-NEXT: [[TMP6:%.*]] = load i32, i32* [[TMP5]], align 4
143 // CHECK4-NEXT: call void @__kmpc_for_static_init_4(%struct.ident_t* @[[GLOB1:[0-9]+]], i32 [[TMP6]], i32 92, i32* [[DOTOMP_IS_LAST]], i32* [[DOTOMP_LB]], i32* [[DOTOMP_UB]], i32* [[DOTOMP_STRIDE]], i32 1, i32 1)
144 // CHECK4-NEXT: [[TMP7:%.*]] = load i32, i32* [[DOTOMP_UB]], align 4
145 // CHECK4-NEXT: [[CMP:%.*]] = icmp sgt i32 [[TMP7]], 9
146 // CHECK4-NEXT: br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]
147 // CHECK4: cond.true:
148 // CHECK4-NEXT: br label [[COND_END:%.*]]
149 // CHECK4: cond.false:
150 // CHECK4-NEXT: [[TMP8:%.*]] = load i32, i32* [[DOTOMP_UB]], align 4
151 // CHECK4-NEXT: br label [[COND_END]]
152 // CHECK4: cond.end:
153 // CHECK4-NEXT: [[COND:%.*]] = phi i32 [ 9, [[COND_TRUE]] ], [ [[TMP8]], [[COND_FALSE]] ]
154 // CHECK4-NEXT: store i32 [[COND]], i32* [[DOTOMP_UB]], align 4
155 // CHECK4-NEXT: [[TMP9:%.*]] = load i32, i32* [[DOTOMP_LB]], align 4
156 // CHECK4-NEXT: store i32 [[TMP9]], i32* [[DOTOMP_IV]], align 4
157 // CHECK4-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]
158 // CHECK4: omp.inner.for.cond:
159 // CHECK4-NEXT: [[TMP10:%.*]] = load i32, i32* [[DOTOMP_IV]], align 4
160 // CHECK4-NEXT: [[TMP11:%.*]] = load i32, i32* [[DOTOMP_UB]], align 4
161 // CHECK4-NEXT: [[CMP1:%.*]] = icmp sle i32 [[TMP10]], [[TMP11]]
162 // CHECK4-NEXT: br i1 [[CMP1]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]
163 // CHECK4: omp.inner.for.body:
164 // CHECK4-NEXT: [[TMP12:%.*]] = load i32, i32* [[DOTOMP_IV]], align 4
165 // CHECK4-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP12]], 1
166 // CHECK4-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]
167 // CHECK4-NEXT: store i32 [[ADD]], i32* [[I]], align 4
168 // CHECK4-NEXT: [[TMP13:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[CAPTURED_VARS_ADDRS]], i32 0, i32 0
169 // CHECK4-NEXT: [[TMP14:%.*]] = bitcast i32* [[I]] to i8*
170 // CHECK4-NEXT: store i8* [[TMP14]], i8** [[TMP13]], align 4
171 // CHECK4-NEXT: [[TMP15:%.*]] = bitcast [1 x i8*]* [[CAPTURED_VARS_ADDRS]] to i8**
172 // CHECK4-NEXT: call void @__kmpc_parallel_51(%struct.ident_t* @[[GLOB2]], i32 [[TMP6]], i32 1, i32 -1, i32 -1, i8* bitcast (void (i32*, i32*, i32*)* @__omp_outlined__1 to i8*), i8* bitcast (void (i16, i32)* @__omp_outlined__1_wrapper to i8*), i8** [[TMP15]], i32 1)
173 // CHECK4-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]
174 // CHECK4: omp.body.continue:
175 // CHECK4-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]
176 // CHECK4: omp.inner.for.inc:
177 // CHECK4-NEXT: [[TMP16:%.*]] = load i32, i32* [[DOTOMP_IV]], align 4
178 // CHECK4-NEXT: [[ADD2:%.*]] = add nsw i32 [[TMP16]], 1
179 // CHECK4-NEXT: store i32 [[ADD2]], i32* [[DOTOMP_IV]], align 4
180 // CHECK4-NEXT: br label [[OMP_INNER_FOR_COND]]
181 // CHECK4: omp.inner.for.end:
182 // CHECK4-NEXT: br label [[OMP_LOOP_EXIT:%.*]]
183 // CHECK4: omp.loop.exit:
184 // CHECK4-NEXT: call void @__kmpc_for_static_fini(%struct.ident_t* @[[GLOB1]], i32 [[TMP6]])
185 // CHECK4-NEXT: [[TMP17:%.*]] = load i16, i16* @"_openmp_static_kernel$is_shared", align 2
186 // CHECK4-NEXT: call void @__kmpc_restore_team_static_memory(i16 0, i16 [[TMP17]])
187 // CHECK4-NEXT: ret void
188 // CHECK4-LABEL: define {{[^@]+}}@__omp_outlined__1
189 // CHECK4-SAME: (i32* noalias [[DOTGLOBAL_TID_:%.*]], i32* noalias [[DOTBOUND_TID_:%.*]], i32* nonnull align 4 dereferenceable(4) [[I:%.*]]) #[[ATTR1]] {
190 // CHECK4-NEXT: entry:
191 // CHECK4-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 4
192 // CHECK4-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 4
193 // CHECK4-NEXT: [[I_ADDR:%.*]] = alloca i32*, align 4
194 // CHECK4-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 4
195 // CHECK4-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 4
196 // CHECK4-NEXT: store i32* [[I]], i32** [[I_ADDR]], align 4
197 // CHECK4-NEXT: [[TMP0:%.*]] = load i32*, i32** [[I_ADDR]], align 4
198 // CHECK4-NEXT: [[TMP1:%.*]] = load i32, i32* [[TMP0]], align 4
199 // CHECK4-NEXT: [[INC:%.*]] = add nsw i32 [[TMP1]], 1
200 // CHECK4-NEXT: store i32 [[INC]], i32* [[TMP0]], align 4
201 // CHECK4-NEXT: ret void
202 // CHECK4-LABEL: define {{[^@]+}}@__omp_outlined__1_wrapper
203 // CHECK4-SAME: (i16 zeroext [[TMP0:%.*]], i32 [[TMP1:%.*]]) #[[ATTR0]] {
204 // CHECK4-NEXT: entry:
205 // CHECK4-NEXT: [[DOTADDR:%.*]] = alloca i16, align 2
206 // CHECK4-NEXT: [[DOTADDR1:%.*]] = alloca i32, align 4
207 // CHECK4-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4
208 // CHECK4-NEXT: [[GLOBAL_ARGS:%.*]] = alloca i8**, align 4
209 // CHECK4-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4
210 // CHECK4-NEXT: store i16 [[TMP0]], i16* [[DOTADDR]], align 2
211 // CHECK4-NEXT: store i32 [[TMP1]], i32* [[DOTADDR1]], align 4
212 // CHECK4-NEXT: call void @__kmpc_get_shared_variables(i8*** [[GLOBAL_ARGS]])
213 // CHECK4-NEXT: [[TMP2:%.*]] = load i8**, i8*** [[GLOBAL_ARGS]], align 4
214 // CHECK4-NEXT: [[TMP3:%.*]] = getelementptr inbounds i8*, i8** [[TMP2]], i32 0
215 // CHECK4-NEXT: [[TMP4:%.*]] = bitcast i8** [[TMP3]] to i32**
216 // CHECK4-NEXT: [[TMP5:%.*]] = load i32*, i32** [[TMP4]], align 4
217 // CHECK4-NEXT: call void @__omp_outlined__1(i32* [[DOTADDR1]], i32* [[DOTZERO_ADDR]], i32* [[TMP5]]) #[[ATTR3]]
218 // CHECK4-NEXT: ret void
219 // CHECK5-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l19_worker
220 // CHECK5-SAME: () #[[ATTR0:[0-9]+]] {
221 // CHECK5-NEXT: entry:
222 // CHECK5-NEXT: [[WORK_FN:%.*]] = alloca i8*, align 4
223 // CHECK5-NEXT: [[EXEC_STATUS:%.*]] = alloca i8, align 1
224 // CHECK5-NEXT: store i8* null, i8** [[WORK_FN]], align 4
225 // CHECK5-NEXT: store i8 0, i8* [[EXEC_STATUS]], align 1
226 // CHECK5-NEXT: br label [[DOTAWAIT_WORK:%.*]]
227 // CHECK5: .await.work:
228 // CHECK5-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0)
229 // CHECK5-NEXT: [[TMP0:%.*]] = call i1 @__kmpc_kernel_parallel(i8** [[WORK_FN]])
230 // CHECK5-NEXT: [[TMP1:%.*]] = zext i1 [[TMP0]] to i8
231 // CHECK5-NEXT: store i8 [[TMP1]], i8* [[EXEC_STATUS]], align 1
232 // CHECK5-NEXT: [[TMP2:%.*]] = load i8*, i8** [[WORK_FN]], align 4
233 // CHECK5-NEXT: [[SHOULD_TERMINATE:%.*]] = icmp eq i8* [[TMP2]], null
234 // CHECK5-NEXT: br i1 [[SHOULD_TERMINATE]], label [[DOTEXIT:%.*]], label [[DOTSELECT_WORKERS:%.*]]
235 // CHECK5: .select.workers:
236 // CHECK5-NEXT: [[TMP3:%.*]] = load i8, i8* [[EXEC_STATUS]], align 1
237 // CHECK5-NEXT: [[IS_ACTIVE:%.*]] = icmp ne i8 [[TMP3]], 0
238 // CHECK5-NEXT: br i1 [[IS_ACTIVE]], label [[DOTEXECUTE_PARALLEL:%.*]], label [[DOTBARRIER_PARALLEL:%.*]]
239 // CHECK5: .execute.parallel:
240 // CHECK5-NEXT: [[TMP4:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB2:[0-9]+]])
241 // CHECK5-NEXT: [[TMP5:%.*]] = load i8*, i8** [[WORK_FN]], align 4
242 // CHECK5-NEXT: [[WORK_MATCH:%.*]] = icmp eq i8* [[TMP5]], bitcast (void (i16, i32)* @__omp_outlined__1_wrapper to i8*)
243 // CHECK5-NEXT: br i1 [[WORK_MATCH]], label [[DOTEXECUTE_FN:%.*]], label [[DOTCHECK_NEXT:%.*]]
244 // CHECK5: .execute.fn:
245 // CHECK5-NEXT: call void @__omp_outlined__1_wrapper(i16 0, i32 [[TMP4]]) #[[ATTR3:[0-9]+]]
246 // CHECK5-NEXT: br label [[DOTTERMINATE_PARALLEL:%.*]]
247 // CHECK5: .check.next:
248 // CHECK5-NEXT: [[TMP6:%.*]] = bitcast i8* [[TMP2]] to void (i16, i32)*
249 // CHECK5-NEXT: call void [[TMP6]](i16 0, i32 [[TMP4]])
250 // CHECK5-NEXT: br label [[DOTTERMINATE_PARALLEL]]
251 // CHECK5: .terminate.parallel:
252 // CHECK5-NEXT: call void @__kmpc_kernel_end_parallel()
253 // CHECK5-NEXT: br label [[DOTBARRIER_PARALLEL]]
254 // CHECK5: .barrier.parallel:
255 // CHECK5-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0)
256 // CHECK5-NEXT: br label [[DOTAWAIT_WORK]]
257 // CHECK5: .exit:
258 // CHECK5-NEXT: ret void
259 // CHECK5-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l19
260 // CHECK5-SAME: () #[[ATTR1:[0-9]+]] {
261 // CHECK5-NEXT: entry:
262 // CHECK5-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4
263 // CHECK5-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 4
264 // CHECK5-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4
265 // CHECK5-NEXT: [[NVPTX_TID:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
266 // CHECK5-NEXT: [[NVPTX_NUM_THREADS:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x()
267 // CHECK5-NEXT: [[NVPTX_WARP_SIZE:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize()
268 // CHECK5-NEXT: [[THREAD_LIMIT:%.*]] = sub nuw i32 [[NVPTX_NUM_THREADS]], [[NVPTX_WARP_SIZE]]
269 // CHECK5-NEXT: [[TMP0:%.*]] = icmp ult i32 [[NVPTX_TID]], [[THREAD_LIMIT]]
270 // CHECK5-NEXT: br i1 [[TMP0]], label [[DOTWORKER:%.*]], label [[DOTMASTERCHECK:%.*]]
271 // CHECK5: .worker:
272 // CHECK5-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l19_worker() #[[ATTR3]]
273 // CHECK5-NEXT: br label [[DOTEXIT:%.*]]
274 // CHECK5: .mastercheck:
275 // CHECK5-NEXT: [[NVPTX_TID1:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
276 // CHECK5-NEXT: [[NVPTX_NUM_THREADS2:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x()
277 // CHECK5-NEXT: [[NVPTX_WARP_SIZE3:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize()
278 // CHECK5-NEXT: [[TMP1:%.*]] = sub nuw i32 [[NVPTX_WARP_SIZE3]], 1
279 // CHECK5-NEXT: [[TMP2:%.*]] = sub nuw i32 [[NVPTX_NUM_THREADS2]], 1
280 // CHECK5-NEXT: [[TMP3:%.*]] = xor i32 [[TMP1]], -1
281 // CHECK5-NEXT: [[MASTER_TID:%.*]] = and i32 [[TMP2]], [[TMP3]]
282 // CHECK5-NEXT: [[TMP4:%.*]] = icmp eq i32 [[NVPTX_TID1]], [[MASTER_TID]]
283 // CHECK5-NEXT: br i1 [[TMP4]], label [[DOTMASTER:%.*]], label [[DOTEXIT]]
284 // CHECK5: .master:
285 // CHECK5-NEXT: [[NVPTX_NUM_THREADS4:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x()
286 // CHECK5-NEXT: [[NVPTX_WARP_SIZE5:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize()
287 // CHECK5-NEXT: [[THREAD_LIMIT6:%.*]] = sub nuw i32 [[NVPTX_NUM_THREADS4]], [[NVPTX_WARP_SIZE5]]
288 // CHECK5-NEXT: call void @__kmpc_kernel_init(i32 [[THREAD_LIMIT6]], i16 1)
289 // CHECK5-NEXT: call void @__kmpc_data_sharing_init_stack()
290 // CHECK5-NEXT: [[TMP5:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB2]])
291 // CHECK5-NEXT: store i32 [[TMP5]], i32* [[DOTTHREADID_TEMP_]], align 4
292 // CHECK5-NEXT: call void @__omp_outlined__(i32* [[DOTTHREADID_TEMP_]], i32* [[DOTZERO_ADDR]]) #[[ATTR3]]
293 // CHECK5-NEXT: br label [[DOTTERMINATION_NOTIFIER:%.*]]
294 // CHECK5: .termination.notifier:
295 // CHECK5-NEXT: call void @__kmpc_kernel_deinit(i16 1)
296 // CHECK5-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0)
297 // CHECK5-NEXT: br label [[DOTEXIT]]
298 // CHECK5: .exit:
299 // CHECK5-NEXT: ret void
300 // CHECK5-LABEL: define {{[^@]+}}@__omp_outlined__
301 // CHECK5-SAME: (i32* noalias [[DOTGLOBAL_TID_:%.*]], i32* noalias [[DOTBOUND_TID_:%.*]]) #[[ATTR1]] {
302 // CHECK5-NEXT: entry:
303 // CHECK5-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 4
304 // CHECK5-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 4
305 // CHECK5-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 4
306 // CHECK5-NEXT: [[TMP:%.*]] = alloca i32, align 4
307 // CHECK5-NEXT: [[DOTOMP_LB:%.*]] = alloca i32, align 4
308 // CHECK5-NEXT: [[DOTOMP_UB:%.*]] = alloca i32, align 4
309 // CHECK5-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4
310 // CHECK5-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4
311 // CHECK5-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [1 x i8*], align 4
312 // CHECK5-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 4
313 // CHECK5-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 4
314 // CHECK5-NEXT: [[TMP0:%.*]] = call i8* @__kmpc_data_sharing_push_stack(i32 4, i16 1)
315 // CHECK5-NEXT: [[TMP1:%.*]] = bitcast i8* [[TMP0]] to %struct._globalized_locals_ty*
316 // CHECK5-NEXT: [[I:%.*]] = getelementptr inbounds [[STRUCT__GLOBALIZED_LOCALS_TY:%.*]], %struct._globalized_locals_ty* [[TMP1]], i32 0, i32 0
317 // CHECK5-NEXT: store i32 0, i32* [[DOTOMP_LB]], align 4
318 // CHECK5-NEXT: store i32 9, i32* [[DOTOMP_UB]], align 4
319 // CHECK5-NEXT: store i32 1, i32* [[DOTOMP_STRIDE]], align 4
320 // CHECK5-NEXT: store i32 0, i32* [[DOTOMP_IS_LAST]], align 4
321 // CHECK5-NEXT: [[TMP2:%.*]] = load i32*, i32** [[DOTGLOBAL_TID__ADDR]], align 4
322 // CHECK5-NEXT: [[TMP3:%.*]] = load i32, i32* [[TMP2]], align 4
323 // CHECK5-NEXT: call void @__kmpc_for_static_init_4(%struct.ident_t* @[[GLOB1:[0-9]+]], i32 [[TMP3]], i32 92, i32* [[DOTOMP_IS_LAST]], i32* [[DOTOMP_LB]], i32* [[DOTOMP_UB]], i32* [[DOTOMP_STRIDE]], i32 1, i32 1)
324 // CHECK5-NEXT: [[TMP4:%.*]] = load i32, i32* [[DOTOMP_UB]], align 4
325 // CHECK5-NEXT: [[CMP:%.*]] = icmp sgt i32 [[TMP4]], 9
326 // CHECK5-NEXT: br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]
327 // CHECK5: cond.true:
328 // CHECK5-NEXT: br label [[COND_END:%.*]]
329 // CHECK5: cond.false:
330 // CHECK5-NEXT: [[TMP5:%.*]] = load i32, i32* [[DOTOMP_UB]], align 4
331 // CHECK5-NEXT: br label [[COND_END]]
332 // CHECK5: cond.end:
333 // CHECK5-NEXT: [[COND:%.*]] = phi i32 [ 9, [[COND_TRUE]] ], [ [[TMP5]], [[COND_FALSE]] ]
334 // CHECK5-NEXT: store i32 [[COND]], i32* [[DOTOMP_UB]], align 4
335 // CHECK5-NEXT: [[TMP6:%.*]] = load i32, i32* [[DOTOMP_LB]], align 4
336 // CHECK5-NEXT: store i32 [[TMP6]], i32* [[DOTOMP_IV]], align 4
337 // CHECK5-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]
338 // CHECK5: omp.inner.for.cond:
339 // CHECK5-NEXT: [[TMP7:%.*]] = load i32, i32* [[DOTOMP_IV]], align 4
340 // CHECK5-NEXT: [[TMP8:%.*]] = load i32, i32* [[DOTOMP_UB]], align 4
341 // CHECK5-NEXT: [[CMP1:%.*]] = icmp sle i32 [[TMP7]], [[TMP8]]
342 // CHECK5-NEXT: br i1 [[CMP1]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]
343 // CHECK5: omp.inner.for.body:
344 // CHECK5-NEXT: [[TMP9:%.*]] = load i32, i32* [[DOTOMP_IV]], align 4
345 // CHECK5-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP9]], 1
346 // CHECK5-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]
347 // CHECK5-NEXT: store i32 [[ADD]], i32* [[I]], align 4
348 // CHECK5-NEXT: [[TMP10:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[CAPTURED_VARS_ADDRS]], i32 0, i32 0
349 // CHECK5-NEXT: [[TMP11:%.*]] = bitcast i32* [[I]] to i8*
350 // CHECK5-NEXT: store i8* [[TMP11]], i8** [[TMP10]], align 4
351 // CHECK5-NEXT: [[TMP12:%.*]] = bitcast [1 x i8*]* [[CAPTURED_VARS_ADDRS]] to i8**
352 // CHECK5-NEXT: call void @__kmpc_parallel_51(%struct.ident_t* @[[GLOB2]], i32 [[TMP3]], i32 1, i32 -1, i32 -1, i8* bitcast (void (i32*, i32*, i32*)* @__omp_outlined__1 to i8*), i8* bitcast (void (i16, i32)* @__omp_outlined__1_wrapper to i8*), i8** [[TMP12]], i32 1)
353 // CHECK5-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]
354 // CHECK5: omp.body.continue:
355 // CHECK5-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]
356 // CHECK5: omp.inner.for.inc:
357 // CHECK5-NEXT: [[TMP13:%.*]] = load i32, i32* [[DOTOMP_IV]], align 4
358 // CHECK5-NEXT: [[ADD2:%.*]] = add nsw i32 [[TMP13]], 1
359 // CHECK5-NEXT: store i32 [[ADD2]], i32* [[DOTOMP_IV]], align 4
360 // CHECK5-NEXT: br label [[OMP_INNER_FOR_COND]]
361 // CHECK5: omp.inner.for.end:
362 // CHECK5-NEXT: br label [[OMP_LOOP_EXIT:%.*]]
363 // CHECK5: omp.loop.exit:
364 // CHECK5-NEXT: call void @__kmpc_for_static_fini(%struct.ident_t* @[[GLOB1]], i32 [[TMP3]])
365 // CHECK5-NEXT: call void @__kmpc_data_sharing_pop_stack(i8* [[TMP0]])
366 // CHECK5-NEXT: ret void
367 // CHECK5-LABEL: define {{[^@]+}}@__omp_outlined__1
368 // CHECK5-SAME: (i32* noalias [[DOTGLOBAL_TID_:%.*]], i32* noalias [[DOTBOUND_TID_:%.*]], i32* nonnull align 4 dereferenceable(4) [[I:%.*]]) #[[ATTR1]] {
369 // CHECK5-NEXT: entry:
370 // CHECK5-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 4
371 // CHECK5-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 4
372 // CHECK5-NEXT: [[I_ADDR:%.*]] = alloca i32*, align 4
373 // CHECK5-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 4
374 // CHECK5-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 4
375 // CHECK5-NEXT: store i32* [[I]], i32** [[I_ADDR]], align 4
376 // CHECK5-NEXT: [[TMP0:%.*]] = load i32*, i32** [[I_ADDR]], align 4
377 // CHECK5-NEXT: [[TMP1:%.*]] = load i32, i32* [[TMP0]], align 4
378 // CHECK5-NEXT: [[INC:%.*]] = add nsw i32 [[TMP1]], 1
379 // CHECK5-NEXT: store i32 [[INC]], i32* [[TMP0]], align 4
380 // CHECK5-NEXT: ret void
381 // CHECK5-LABEL: define {{[^@]+}}@__omp_outlined__1_wrapper
382 // CHECK5-SAME: (i16 zeroext [[TMP0:%.*]], i32 [[TMP1:%.*]]) #[[ATTR0]] {
383 // CHECK5-NEXT: entry:
384 // CHECK5-NEXT: [[DOTADDR:%.*]] = alloca i16, align 2
385 // CHECK5-NEXT: [[DOTADDR1:%.*]] = alloca i32, align 4
386 // CHECK5-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4
387 // CHECK5-NEXT: [[GLOBAL_ARGS:%.*]] = alloca i8**, align 4
388 // CHECK5-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4
389 // CHECK5-NEXT: store i16 [[TMP0]], i16* [[DOTADDR]], align 2
390 // CHECK5-NEXT: store i32 [[TMP1]], i32* [[DOTADDR1]], align 4
391 // CHECK5-NEXT: call void @__kmpc_get_shared_variables(i8*** [[GLOBAL_ARGS]])
392 // CHECK5-NEXT: [[TMP2:%.*]] = load i8**, i8*** [[GLOBAL_ARGS]], align 4
393 // CHECK5-NEXT: [[TMP3:%.*]] = getelementptr inbounds i8*, i8** [[TMP2]], i32 0
394 // CHECK5-NEXT: [[TMP4:%.*]] = bitcast i8** [[TMP3]] to i32**
395 // CHECK5-NEXT: [[TMP5:%.*]] = load i32*, i32** [[TMP4]], align 4
396 // CHECK5-NEXT: call void @__omp_outlined__1(i32* [[DOTADDR1]], i32* [[DOTZERO_ADDR]], i32* [[TMP5]]) #[[ATTR3]]
397 // CHECK5-NEXT: ret void
398 // CHECK6-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l19_worker
399 // CHECK6-SAME: () #[[ATTR0:[0-9]+]] {
400 // CHECK6-NEXT: entry:
401 // CHECK6-NEXT: [[WORK_FN:%.*]] = alloca i8*, align 4
402 // CHECK6-NEXT: [[EXEC_STATUS:%.*]] = alloca i8, align 1
403 // CHECK6-NEXT: store i8* null, i8** [[WORK_FN]], align 4
404 // CHECK6-NEXT: store i8 0, i8* [[EXEC_STATUS]], align 1
405 // CHECK6-NEXT: br label [[DOTAWAIT_WORK:%.*]]
406 // CHECK6: .await.work:
407 // CHECK6-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0)
408 // CHECK6-NEXT: [[TMP0:%.*]] = call i1 @__kmpc_kernel_parallel(i8** [[WORK_FN]])
409 // CHECK6-NEXT: [[TMP1:%.*]] = zext i1 [[TMP0]] to i8
410 // CHECK6-NEXT: store i8 [[TMP1]], i8* [[EXEC_STATUS]], align 1
411 // CHECK6-NEXT: [[TMP2:%.*]] = load i8*, i8** [[WORK_FN]], align 4
412 // CHECK6-NEXT: [[SHOULD_TERMINATE:%.*]] = icmp eq i8* [[TMP2]], null
413 // CHECK6-NEXT: br i1 [[SHOULD_TERMINATE]], label [[DOTEXIT:%.*]], label [[DOTSELECT_WORKERS:%.*]]
414 // CHECK6: .select.workers:
415 // CHECK6-NEXT: [[TMP3:%.*]] = load i8, i8* [[EXEC_STATUS]], align 1
416 // CHECK6-NEXT: [[IS_ACTIVE:%.*]] = icmp ne i8 [[TMP3]], 0
417 // CHECK6-NEXT: br i1 [[IS_ACTIVE]], label [[DOTEXECUTE_PARALLEL:%.*]], label [[DOTBARRIER_PARALLEL:%.*]]
418 // CHECK6: .execute.parallel:
419 // CHECK6-NEXT: [[TMP4:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB2:[0-9]+]])
420 // CHECK6-NEXT: [[TMP5:%.*]] = load i8*, i8** [[WORK_FN]], align 4
421 // CHECK6-NEXT: [[WORK_MATCH:%.*]] = icmp eq i8* [[TMP5]], bitcast (void (i16, i32)* @__omp_outlined__1_wrapper to i8*)
422 // CHECK6-NEXT: br i1 [[WORK_MATCH]], label [[DOTEXECUTE_FN:%.*]], label [[DOTCHECK_NEXT:%.*]]
423 // CHECK6: .execute.fn:
424 // CHECK6-NEXT: call void @__omp_outlined__1_wrapper(i16 0, i32 [[TMP4]]) #[[ATTR3:[0-9]+]]
425 // CHECK6-NEXT: br label [[DOTTERMINATE_PARALLEL:%.*]]
426 // CHECK6: .check.next:
427 // CHECK6-NEXT: [[TMP6:%.*]] = bitcast i8* [[TMP2]] to void (i16, i32)*
428 // CHECK6-NEXT: call void [[TMP6]](i16 0, i32 [[TMP4]])
429 // CHECK6-NEXT: br label [[DOTTERMINATE_PARALLEL]]
430 // CHECK6: .terminate.parallel:
431 // CHECK6-NEXT: call void @__kmpc_kernel_end_parallel()
432 // CHECK6-NEXT: br label [[DOTBARRIER_PARALLEL]]
433 // CHECK6: .barrier.parallel:
434 // CHECK6-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0)
435 // CHECK6-NEXT: br label [[DOTAWAIT_WORK]]
436 // CHECK6: .exit:
437 // CHECK6-NEXT: ret void
438 // CHECK6-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l19
439 // CHECK6-SAME: () #[[ATTR1:[0-9]+]] {
440 // CHECK6-NEXT: entry:
441 // CHECK6-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4
442 // CHECK6-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 4
443 // CHECK6-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4
444 // CHECK6-NEXT: [[NVPTX_TID:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
445 // CHECK6-NEXT: [[NVPTX_NUM_THREADS:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x()
446 // CHECK6-NEXT: [[NVPTX_WARP_SIZE:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize()
447 // CHECK6-NEXT: [[THREAD_LIMIT:%.*]] = sub nuw i32 [[NVPTX_NUM_THREADS]], [[NVPTX_WARP_SIZE]]
448 // CHECK6-NEXT: [[TMP0:%.*]] = icmp ult i32 [[NVPTX_TID]], [[THREAD_LIMIT]]
449 // CHECK6-NEXT: br i1 [[TMP0]], label [[DOTWORKER:%.*]], label [[DOTMASTERCHECK:%.*]]
450 // CHECK6: .worker:
451 // CHECK6-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l19_worker() #[[ATTR3]]
452 // CHECK6-NEXT: br label [[DOTEXIT:%.*]]
453 // CHECK6: .mastercheck:
454 // CHECK6-NEXT: [[NVPTX_TID1:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
455 // CHECK6-NEXT: [[NVPTX_NUM_THREADS2:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x()
456 // CHECK6-NEXT: [[NVPTX_WARP_SIZE3:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize()
457 // CHECK6-NEXT: [[TMP1:%.*]] = sub nuw i32 [[NVPTX_WARP_SIZE3]], 1
458 // CHECK6-NEXT: [[TMP2:%.*]] = sub nuw i32 [[NVPTX_NUM_THREADS2]], 1
459 // CHECK6-NEXT: [[TMP3:%.*]] = xor i32 [[TMP1]], -1
460 // CHECK6-NEXT: [[MASTER_TID:%.*]] = and i32 [[TMP2]], [[TMP3]]
461 // CHECK6-NEXT: [[TMP4:%.*]] = icmp eq i32 [[NVPTX_TID1]], [[MASTER_TID]]
462 // CHECK6-NEXT: br i1 [[TMP4]], label [[DOTMASTER:%.*]], label [[DOTEXIT]]
463 // CHECK6: .master:
464 // CHECK6-NEXT: [[NVPTX_NUM_THREADS4:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x()
465 // CHECK6-NEXT: [[NVPTX_WARP_SIZE5:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize()
466 // CHECK6-NEXT: [[THREAD_LIMIT6:%.*]] = sub nuw i32 [[NVPTX_NUM_THREADS4]], [[NVPTX_WARP_SIZE5]]
467 // CHECK6-NEXT: call void @__kmpc_kernel_init(i32 [[THREAD_LIMIT6]], i16 1)
468 // CHECK6-NEXT: call void @__kmpc_data_sharing_init_stack()
469 // CHECK6-NEXT: [[TMP5:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB2]])
470 // CHECK6-NEXT: store i32 [[TMP5]], i32* [[DOTTHREADID_TEMP_]], align 4
471 // CHECK6-NEXT: call void @__omp_outlined__(i32* [[DOTTHREADID_TEMP_]], i32* [[DOTZERO_ADDR]]) #[[ATTR3]]
472 // CHECK6-NEXT: br label [[DOTTERMINATION_NOTIFIER:%.*]]
473 // CHECK6: .termination.notifier:
474 // CHECK6-NEXT: call void @__kmpc_kernel_deinit(i16 1)
475 // CHECK6-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0)
476 // CHECK6-NEXT: br label [[DOTEXIT]]
477 // CHECK6: .exit:
478 // CHECK6-NEXT: ret void
479 // CHECK6-LABEL: define {{[^@]+}}@__omp_outlined__
480 // CHECK6-SAME: (i32* noalias [[DOTGLOBAL_TID_:%.*]], i32* noalias [[DOTBOUND_TID_:%.*]]) #[[ATTR1]] {
481 // CHECK6-NEXT: entry:
482 // CHECK6-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 4
483 // CHECK6-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 4
484 // CHECK6-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 4
485 // CHECK6-NEXT: [[TMP:%.*]] = alloca i32, align 4
486 // CHECK6-NEXT: [[DOTOMP_LB:%.*]] = alloca i32, align 4
487 // CHECK6-NEXT: [[DOTOMP_UB:%.*]] = alloca i32, align 4
488 // CHECK6-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4
489 // CHECK6-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4
490 // CHECK6-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [1 x i8*], align 4
491 // CHECK6-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 4
492 // CHECK6-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 4
493 // CHECK6-NEXT: [[TMP0:%.*]] = call i8* @__kmpc_data_sharing_push_stack(i32 4, i16 1)
494 // CHECK6-NEXT: [[TMP1:%.*]] = bitcast i8* [[TMP0]] to %struct._globalized_locals_ty*
495 // CHECK6-NEXT: [[I:%.*]] = getelementptr inbounds [[STRUCT__GLOBALIZED_LOCALS_TY:%.*]], %struct._globalized_locals_ty* [[TMP1]], i32 0, i32 0
496 // CHECK6-NEXT: store i32 0, i32* [[DOTOMP_LB]], align 4
497 // CHECK6-NEXT: store i32 9, i32* [[DOTOMP_UB]], align 4
498 // CHECK6-NEXT: store i32 1, i32* [[DOTOMP_STRIDE]], align 4
499 // CHECK6-NEXT: store i32 0, i32* [[DOTOMP_IS_LAST]], align 4
500 // CHECK6-NEXT: [[TMP2:%.*]] = load i32*, i32** [[DOTGLOBAL_TID__ADDR]], align 4
501 // CHECK6-NEXT: [[TMP3:%.*]] = load i32, i32* [[TMP2]], align 4
502 // CHECK6-NEXT: call void @__kmpc_for_static_init_4(%struct.ident_t* @[[GLOB1:[0-9]+]], i32 [[TMP3]], i32 92, i32* [[DOTOMP_IS_LAST]], i32* [[DOTOMP_LB]], i32* [[DOTOMP_UB]], i32* [[DOTOMP_STRIDE]], i32 1, i32 1)
503 // CHECK6-NEXT: [[TMP4:%.*]] = load i32, i32* [[DOTOMP_UB]], align 4
504 // CHECK6-NEXT: [[CMP:%.*]] = icmp sgt i32 [[TMP4]], 9
505 // CHECK6-NEXT: br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]
506 // CHECK6: cond.true:
507 // CHECK6-NEXT: br label [[COND_END:%.*]]
508 // CHECK6: cond.false:
509 // CHECK6-NEXT: [[TMP5:%.*]] = load i32, i32* [[DOTOMP_UB]], align 4
510 // CHECK6-NEXT: br label [[COND_END]]
511 // CHECK6: cond.end:
512 // CHECK6-NEXT: [[COND:%.*]] = phi i32 [ 9, [[COND_TRUE]] ], [ [[TMP5]], [[COND_FALSE]] ]
513 // CHECK6-NEXT: store i32 [[COND]], i32* [[DOTOMP_UB]], align 4
514 // CHECK6-NEXT: [[TMP6:%.*]] = load i32, i32* [[DOTOMP_LB]], align 4
515 // CHECK6-NEXT: store i32 [[TMP6]], i32* [[DOTOMP_IV]], align 4
516 // CHECK6-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]
517 // CHECK6: omp.inner.for.cond:
518 // CHECK6-NEXT: [[TMP7:%.*]] = load i32, i32* [[DOTOMP_IV]], align 4
519 // CHECK6-NEXT: [[TMP8:%.*]] = load i32, i32* [[DOTOMP_UB]], align 4
520 // CHECK6-NEXT: [[CMP1:%.*]] = icmp sle i32 [[TMP7]], [[TMP8]]
521 // CHECK6-NEXT: br i1 [[CMP1]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]
522 // CHECK6: omp.inner.for.body:
523 // CHECK6-NEXT: [[TMP9:%.*]] = load i32, i32* [[DOTOMP_IV]], align 4
524 // CHECK6-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP9]], 1
525 // CHECK6-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]
526 // CHECK6-NEXT: store i32 [[ADD]], i32* [[I]], align 4
527 // CHECK6-NEXT: [[TMP10:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[CAPTURED_VARS_ADDRS]], i32 0, i32 0
528 // CHECK6-NEXT: [[TMP11:%.*]] = bitcast i32* [[I]] to i8*
529 // CHECK6-NEXT: store i8* [[TMP11]], i8** [[TMP10]], align 4
530 // CHECK6-NEXT: [[TMP12:%.*]] = bitcast [1 x i8*]* [[CAPTURED_VARS_ADDRS]] to i8**
531 // CHECK6-NEXT: call void @__kmpc_parallel_51(%struct.ident_t* @[[GLOB2]], i32 [[TMP3]], i32 1, i32 -1, i32 -1, i8* bitcast (void (i32*, i32*, i32*)* @__omp_outlined__1 to i8*), i8* bitcast (void (i16, i32)* @__omp_outlined__1_wrapper to i8*), i8** [[TMP12]], i32 1)
532 // CHECK6-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]
533 // CHECK6: omp.body.continue:
534 // CHECK6-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]
535 // CHECK6: omp.inner.for.inc:
536 // CHECK6-NEXT: [[TMP13:%.*]] = load i32, i32* [[DOTOMP_IV]], align 4
537 // CHECK6-NEXT: [[ADD2:%.*]] = add nsw i32 [[TMP13]], 1
538 // CHECK6-NEXT: store i32 [[ADD2]], i32* [[DOTOMP_IV]], align 4
539 // CHECK6-NEXT: br label [[OMP_INNER_FOR_COND]]
540 // CHECK6: omp.inner.for.end:
541 // CHECK6-NEXT: br label [[OMP_LOOP_EXIT:%.*]]
542 // CHECK6: omp.loop.exit:
543 // CHECK6-NEXT: call void @__kmpc_for_static_fini(%struct.ident_t* @[[GLOB1]], i32 [[TMP3]])
544 // CHECK6-NEXT: call void @__kmpc_data_sharing_pop_stack(i8* [[TMP0]])
545 // CHECK6-NEXT: ret void
546 // CHECK6-LABEL: define {{[^@]+}}@__omp_outlined__1
547 // CHECK6-SAME: (i32* noalias [[DOTGLOBAL_TID_:%.*]], i32* noalias [[DOTBOUND_TID_:%.*]], i32* nonnull align 4 dereferenceable(4) [[I:%.*]]) #[[ATTR1]] {
548 // CHECK6-NEXT: entry:
549 // CHECK6-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 4
550 // CHECK6-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 4
551 // CHECK6-NEXT: [[I_ADDR:%.*]] = alloca i32*, align 4
552 // CHECK6-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 4
553 // CHECK6-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 4
554 // CHECK6-NEXT: store i32* [[I]], i32** [[I_ADDR]], align 4
555 // CHECK6-NEXT: [[TMP0:%.*]] = load i32*, i32** [[I_ADDR]], align 4
556 // CHECK6-NEXT: [[TMP1:%.*]] = load i32, i32* [[TMP0]], align 4
557 // CHECK6-NEXT: [[INC:%.*]] = add nsw i32 [[TMP1]], 1
558 // CHECK6-NEXT: store i32 [[INC]], i32* [[TMP0]], align 4
559 // CHECK6-NEXT: ret void
560 // CHECK6-LABEL: define {{[^@]+}}@__omp_outlined__1_wrapper
561 // CHECK6-SAME: (i16 zeroext [[TMP0:%.*]], i32 [[TMP1:%.*]]) #[[ATTR0]] {
562 // CHECK6-NEXT: entry:
563 // CHECK6-NEXT: [[DOTADDR:%.*]] = alloca i16, align 2
564 // CHECK6-NEXT: [[DOTADDR1:%.*]] = alloca i32, align 4
565 // CHECK6-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4
566 // CHECK6-NEXT: [[GLOBAL_ARGS:%.*]] = alloca i8**, align 4
567 // CHECK6-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4
568 // CHECK6-NEXT: store i16 [[TMP0]], i16* [[DOTADDR]], align 2
569 // CHECK6-NEXT: store i32 [[TMP1]], i32* [[DOTADDR1]], align 4
570 // CHECK6-NEXT: call void @__kmpc_get_shared_variables(i8*** [[GLOBAL_ARGS]])
571 // CHECK6-NEXT: [[TMP2:%.*]] = load i8**, i8*** [[GLOBAL_ARGS]], align 4
572 // CHECK6-NEXT: [[TMP3:%.*]] = getelementptr inbounds i8*, i8** [[TMP2]], i32 0
573 // CHECK6-NEXT: [[TMP4:%.*]] = bitcast i8** [[TMP3]] to i32**
574 // CHECK6-NEXT: [[TMP5:%.*]] = load i32*, i32** [[TMP4]], align 4
575 // CHECK6-NEXT: call void @__omp_outlined__1(i32* [[DOTADDR1]], i32* [[DOTZERO_ADDR]], i32* [[TMP5]]) #[[ATTR3]]
576 // CHECK6-NEXT: ret void
577 // CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l16
578 // CHECK1-SAME: () #[[ATTR0:[0-9]+]] {
579 // CHECK1-NEXT: entry:
580 // CHECK1-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4
581 // CHECK1-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 4
582 // CHECK1-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4
583 // CHECK1-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_target_init(%struct.ident_t* @[[GLOB1:[0-9]+]], i1 false, i1 true, i1 true)
584 // CHECK1-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP0]], -1
585 // CHECK1-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
586 // CHECK1: user_code.entry:
587 // CHECK1-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1]])
588 // CHECK1-NEXT: store i32 [[TMP1]], i32* [[DOTTHREADID_TEMP_]], align 4
589 // CHECK1-NEXT: call void @__omp_outlined__(i32* [[DOTTHREADID_TEMP_]], i32* [[DOTZERO_ADDR]]) #[[ATTR1:[0-9]+]]
590 // CHECK1-NEXT: call void @__kmpc_target_deinit(%struct.ident_t* @[[GLOB1]], i1 false, i1 true)
591 // CHECK1-NEXT: ret void
592 // CHECK1: worker.exit:
593 // CHECK1-NEXT: ret void
594 //
595 //
596 // CHECK1-LABEL: define {{[^@]+}}@__omp_outlined__
597 // CHECK1-SAME: (i32* noalias [[DOTGLOBAL_TID_:%.*]], i32* noalias [[DOTBOUND_TID_:%.*]]) #[[ATTR0]] {
598 // CHECK1-NEXT: entry:
599 // CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8
600 // CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8
601 // CHECK1-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 4
602 // CHECK1-NEXT: [[TMP:%.*]] = alloca i32, align 4
603 // CHECK1-NEXT: [[DOTOMP_LB:%.*]] = alloca i32, align 4
604 // CHECK1-NEXT: [[DOTOMP_UB:%.*]] = alloca i32, align 4
605 // CHECK1-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4
606 // CHECK1-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4
607 // CHECK1-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [1 x i8*], align 8
608 // CHECK1-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 8
609 // CHECK1-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 8
610 // CHECK1-NEXT: [[I:%.*]] = call i8* @__kmpc_alloc_shared(i64 4)
611 // CHECK1-NEXT: [[I_ON_STACK:%.*]] = bitcast i8* [[I]] to i32*
612 // CHECK1-NEXT: store i32 0, i32* [[DOTOMP_LB]], align 4
613 // CHECK1-NEXT: store i32 9, i32* [[DOTOMP_UB]], align 4
614 // CHECK1-NEXT: store i32 1, i32* [[DOTOMP_STRIDE]], align 4
615 // CHECK1-NEXT: store i32 0, i32* [[DOTOMP_IS_LAST]], align 4
616 // CHECK1-NEXT: [[TMP0:%.*]] = load i32*, i32** [[DOTGLOBAL_TID__ADDR]], align 8
617 // CHECK1-NEXT: [[TMP1:%.*]] = load i32, i32* [[TMP0]], align 4
618 // CHECK1-NEXT: call void @__kmpc_for_static_init_4(%struct.ident_t* @[[GLOB2:[0-9]+]], i32 [[TMP1]], i32 92, i32* [[DOTOMP_IS_LAST]], i32* [[DOTOMP_LB]], i32* [[DOTOMP_UB]], i32* [[DOTOMP_STRIDE]], i32 1, i32 1)
619 // CHECK1-NEXT: [[TMP2:%.*]] = load i32, i32* [[DOTOMP_UB]], align 4
620 // CHECK1-NEXT: [[CMP:%.*]] = icmp sgt i32 [[TMP2]], 9
621 // CHECK1-NEXT: br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]
622 // CHECK1: cond.true:
623 // CHECK1-NEXT: br label [[COND_END:%.*]]
624 // CHECK1: cond.false:
625 // CHECK1-NEXT: [[TMP3:%.*]] = load i32, i32* [[DOTOMP_UB]], align 4
626 // CHECK1-NEXT: br label [[COND_END]]
627 // CHECK1: cond.end:
628 // CHECK1-NEXT: [[COND:%.*]] = phi i32 [ 9, [[COND_TRUE]] ], [ [[TMP3]], [[COND_FALSE]] ]
629 // CHECK1-NEXT: store i32 [[COND]], i32* [[DOTOMP_UB]], align 4
630 // CHECK1-NEXT: [[TMP4:%.*]] = load i32, i32* [[DOTOMP_LB]], align 4
631 // CHECK1-NEXT: store i32 [[TMP4]], i32* [[DOTOMP_IV]], align 4
632 // CHECK1-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]
633 // CHECK1: omp.inner.for.cond:
634 // CHECK1-NEXT: [[TMP5:%.*]] = load i32, i32* [[DOTOMP_IV]], align 4
635 // CHECK1-NEXT: [[TMP6:%.*]] = load i32, i32* [[DOTOMP_UB]], align 4
636 // CHECK1-NEXT: [[CMP1:%.*]] = icmp sle i32 [[TMP5]], [[TMP6]]
637 // CHECK1-NEXT: br i1 [[CMP1]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]
638 // CHECK1: omp.inner.for.body:
639 // CHECK1-NEXT: [[TMP7:%.*]] = load i32, i32* [[DOTOMP_IV]], align 4
640 // CHECK1-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP7]], 1
641 // CHECK1-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]
642 // CHECK1-NEXT: store i32 [[ADD]], i32* [[I_ON_STACK]], align 4
643 // CHECK1-NEXT: [[TMP8:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[CAPTURED_VARS_ADDRS]], i64 0, i64 0
644 // CHECK1-NEXT: [[TMP9:%.*]] = bitcast i32* [[I_ON_STACK]] to i8*
645 // CHECK1-NEXT: store i8* [[TMP9]], i8** [[TMP8]], align 8
646 // CHECK1-NEXT: [[TMP10:%.*]] = bitcast [1 x i8*]* [[CAPTURED_VARS_ADDRS]] to i8**
647 // CHECK1-NEXT: call void @__kmpc_parallel_51(%struct.ident_t* @[[GLOB1]], i32 [[TMP1]], i32 1, i32 -1, i32 -1, i8* bitcast (void (i32*, i32*, i32*)* @__omp_outlined__1 to i8*), i8* bitcast (void (i16, i32)* @__omp_outlined__1_wrapper to i8*), i8** [[TMP10]], i64 1)
648 // CHECK1-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]
649 // CHECK1: omp.body.continue:
650 // CHECK1-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]
651 // CHECK1: omp.inner.for.inc:
652 // CHECK1-NEXT: [[TMP11:%.*]] = load i32, i32* [[DOTOMP_IV]], align 4
653 // CHECK1-NEXT: [[ADD2:%.*]] = add nsw i32 [[TMP11]], 1
654 // CHECK1-NEXT: store i32 [[ADD2]], i32* [[DOTOMP_IV]], align 4
655 // CHECK1-NEXT: br label [[OMP_INNER_FOR_COND]]
656 // CHECK1: omp.inner.for.end:
657 // CHECK1-NEXT: br label [[OMP_LOOP_EXIT:%.*]]
658 // CHECK1: omp.loop.exit:
659 // CHECK1-NEXT: call void @__kmpc_for_static_fini(%struct.ident_t* @[[GLOB2]], i32 [[TMP1]])
660 // CHECK1-NEXT: call void @__kmpc_free_shared(i8* [[I]], i64 4)
661 // CHECK1-NEXT: ret void
662 //
663 //
664 // CHECK1-LABEL: define {{[^@]+}}@__omp_outlined__1
665 // CHECK1-SAME: (i32* noalias [[DOTGLOBAL_TID_:%.*]], i32* noalias [[DOTBOUND_TID_:%.*]], i32* nonnull align 4 dereferenceable(4) [[I:%.*]]) #[[ATTR0]] {
666 // CHECK1-NEXT: entry:
667 // CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8
668 // CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8
669 // CHECK1-NEXT: [[I_ADDR:%.*]] = alloca i32*, align 8
670 // CHECK1-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 8
671 // CHECK1-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 8
672 // CHECK1-NEXT: store i32* [[I]], i32** [[I_ADDR]], align 8
673 // CHECK1-NEXT: [[TMP0:%.*]] = load i32*, i32** [[I_ADDR]], align 8
674 // CHECK1-NEXT: [[TMP1:%.*]] = load i32, i32* [[TMP0]], align 4
675 // CHECK1-NEXT: [[INC:%.*]] = add nsw i32 [[TMP1]], 1
676 // CHECK1-NEXT: store i32 [[INC]], i32* [[TMP0]], align 4
677 // CHECK1-NEXT: ret void
678 //
679 //
680 // CHECK1-LABEL: define {{[^@]+}}@__omp_outlined__1_wrapper
681 // CHECK1-SAME: (i16 zeroext [[TMP0:%.*]], i32 [[TMP1:%.*]]) #[[ATTR2:[0-9]+]] {
682 // CHECK1-NEXT: entry:
683 // CHECK1-NEXT: [[DOTADDR:%.*]] = alloca i16, align 2
684 // CHECK1-NEXT: [[DOTADDR1:%.*]] = alloca i32, align 4
685 // CHECK1-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4
686 // CHECK1-NEXT: [[GLOBAL_ARGS:%.*]] = alloca i8**, align 8
687 // CHECK1-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4
688 // CHECK1-NEXT: store i16 [[TMP0]], i16* [[DOTADDR]], align 2
689 // CHECK1-NEXT: store i32 [[TMP1]], i32* [[DOTADDR1]], align 4
690 // CHECK1-NEXT: call void @__kmpc_get_shared_variables(i8*** [[GLOBAL_ARGS]])
691 // CHECK1-NEXT: [[TMP2:%.*]] = load i8**, i8*** [[GLOBAL_ARGS]], align 8
692 // CHECK1-NEXT: [[TMP3:%.*]] = getelementptr inbounds i8*, i8** [[TMP2]], i64 0
693 // CHECK1-NEXT: [[TMP4:%.*]] = bitcast i8** [[TMP3]] to i32**
694 // CHECK1-NEXT: [[TMP5:%.*]] = load i32*, i32** [[TMP4]], align 8
695 // CHECK1-NEXT: call void @__omp_outlined__1(i32* [[DOTADDR1]], i32* [[DOTZERO_ADDR]], i32* [[TMP5]]) #[[ATTR1]]
696 // CHECK1-NEXT: ret void
697 //
698 //
699 // CHECK2-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l16
700 // CHECK2-SAME: () #[[ATTR0:[0-9]+]] {
701 // CHECK2-NEXT: entry:
702 // CHECK2-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4
703 // CHECK2-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 4
704 // CHECK2-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4
705 // CHECK2-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_target_init(%struct.ident_t* @[[GLOB1:[0-9]+]], i1 false, i1 true, i1 true)
706 // CHECK2-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP0]], -1
707 // CHECK2-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
708 // CHECK2: user_code.entry:
709 // CHECK2-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1]])
710 // CHECK2-NEXT: store i32 [[TMP1]], i32* [[DOTTHREADID_TEMP_]], align 4
711 // CHECK2-NEXT: call void @__omp_outlined__(i32* [[DOTTHREADID_TEMP_]], i32* [[DOTZERO_ADDR]]) #[[ATTR1:[0-9]+]]
712 // CHECK2-NEXT: call void @__kmpc_target_deinit(%struct.ident_t* @[[GLOB1]], i1 false, i1 true)
713 // CHECK2-NEXT: ret void
714 // CHECK2: worker.exit:
715 // CHECK2-NEXT: ret void
716 //
717 //
718 // CHECK2-LABEL: define {{[^@]+}}@__omp_outlined__
719 // CHECK2-SAME: (i32* noalias [[DOTGLOBAL_TID_:%.*]], i32* noalias [[DOTBOUND_TID_:%.*]]) #[[ATTR0]] {
720 // CHECK2-NEXT: entry:
721 // CHECK2-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 4
722 // CHECK2-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 4
723 // CHECK2-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 4
724 // CHECK2-NEXT: [[TMP:%.*]] = alloca i32, align 4
725 // CHECK2-NEXT: [[DOTOMP_LB:%.*]] = alloca i32, align 4
726 // CHECK2-NEXT: [[DOTOMP_UB:%.*]] = alloca i32, align 4
727 // CHECK2-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4
728 // CHECK2-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4
729 // CHECK2-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [1 x i8*], align 4
730 // CHECK2-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 4
731 // CHECK2-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 4
732 // CHECK2-NEXT: [[I:%.*]] = call i8* @__kmpc_alloc_shared(i32 4)
733 // CHECK2-NEXT: [[I_ON_STACK:%.*]] = bitcast i8* [[I]] to i32*
734 // CHECK2-NEXT: store i32 0, i32* [[DOTOMP_LB]], align 4
735 // CHECK2-NEXT: store i32 9, i32* [[DOTOMP_UB]], align 4
736 // CHECK2-NEXT: store i32 1, i32* [[DOTOMP_STRIDE]], align 4
737 // CHECK2-NEXT: store i32 0, i32* [[DOTOMP_IS_LAST]], align 4
738 // CHECK2-NEXT: [[TMP0:%.*]] = load i32*, i32** [[DOTGLOBAL_TID__ADDR]], align 4
739 // CHECK2-NEXT: [[TMP1:%.*]] = load i32, i32* [[TMP0]], align 4
740 // CHECK2-NEXT: call void @__kmpc_for_static_init_4(%struct.ident_t* @[[GLOB2:[0-9]+]], i32 [[TMP1]], i32 92, i32* [[DOTOMP_IS_LAST]], i32* [[DOTOMP_LB]], i32* [[DOTOMP_UB]], i32* [[DOTOMP_STRIDE]], i32 1, i32 1)
741 // CHECK2-NEXT: [[TMP2:%.*]] = load i32, i32* [[DOTOMP_UB]], align 4
742 // CHECK2-NEXT: [[CMP:%.*]] = icmp sgt i32 [[TMP2]], 9
743 // CHECK2-NEXT: br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]
744 // CHECK2: cond.true:
745 // CHECK2-NEXT: br label [[COND_END:%.*]]
746 // CHECK2: cond.false:
747 // CHECK2-NEXT: [[TMP3:%.*]] = load i32, i32* [[DOTOMP_UB]], align 4
748 // CHECK2-NEXT: br label [[COND_END]]
749 // CHECK2: cond.end:
750 // CHECK2-NEXT: [[COND:%.*]] = phi i32 [ 9, [[COND_TRUE]] ], [ [[TMP3]], [[COND_FALSE]] ]
751 // CHECK2-NEXT: store i32 [[COND]], i32* [[DOTOMP_UB]], align 4
752 // CHECK2-NEXT: [[TMP4:%.*]] = load i32, i32* [[DOTOMP_LB]], align 4
753 // CHECK2-NEXT: store i32 [[TMP4]], i32* [[DOTOMP_IV]], align 4
754 // CHECK2-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]
755 // CHECK2: omp.inner.for.cond:
756 // CHECK2-NEXT: [[TMP5:%.*]] = load i32, i32* [[DOTOMP_IV]], align 4
757 // CHECK2-NEXT: [[TMP6:%.*]] = load i32, i32* [[DOTOMP_UB]], align 4
758 // CHECK2-NEXT: [[CMP1:%.*]] = icmp sle i32 [[TMP5]], [[TMP6]]
759 // CHECK2-NEXT: br i1 [[CMP1]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]
760 // CHECK2: omp.inner.for.body:
761 // CHECK2-NEXT: [[TMP7:%.*]] = load i32, i32* [[DOTOMP_IV]], align 4
762 // CHECK2-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP7]], 1
763 // CHECK2-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]
764 // CHECK2-NEXT: store i32 [[ADD]], i32* [[I_ON_STACK]], align 4
765 // CHECK2-NEXT: [[TMP8:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[CAPTURED_VARS_ADDRS]], i32 0, i32 0
766 // CHECK2-NEXT: [[TMP9:%.*]] = bitcast i32* [[I_ON_STACK]] to i8*
767 // CHECK2-NEXT: store i8* [[TMP9]], i8** [[TMP8]], align 4
768 // CHECK2-NEXT: [[TMP10:%.*]] = bitcast [1 x i8*]* [[CAPTURED_VARS_ADDRS]] to i8**
769 // CHECK2-NEXT: call void @__kmpc_parallel_51(%struct.ident_t* @[[GLOB1]], i32 [[TMP1]], i32 1, i32 -1, i32 -1, i8* bitcast (void (i32*, i32*, i32*)* @__omp_outlined__1 to i8*), i8* bitcast (void (i16, i32)* @__omp_outlined__1_wrapper to i8*), i8** [[TMP10]], i32 1)
770 // CHECK2-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]
771 // CHECK2: omp.body.continue:
772 // CHECK2-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]
773 // CHECK2: omp.inner.for.inc:
774 // CHECK2-NEXT: [[TMP11:%.*]] = load i32, i32* [[DOTOMP_IV]], align 4
775 // CHECK2-NEXT: [[ADD2:%.*]] = add nsw i32 [[TMP11]], 1
776 // CHECK2-NEXT: store i32 [[ADD2]], i32* [[DOTOMP_IV]], align 4
777 // CHECK2-NEXT: br label [[OMP_INNER_FOR_COND]]
778 // CHECK2: omp.inner.for.end:
779 // CHECK2-NEXT: br label [[OMP_LOOP_EXIT:%.*]]
780 // CHECK2: omp.loop.exit:
781 // CHECK2-NEXT: call void @__kmpc_for_static_fini(%struct.ident_t* @[[GLOB2]], i32 [[TMP1]])
782 // CHECK2-NEXT: call void @__kmpc_free_shared(i8* [[I]], i32 4)
783 // CHECK2-NEXT: ret void
784 //
785 //
786 // CHECK2-LABEL: define {{[^@]+}}@__omp_outlined__1
787 // CHECK2-SAME: (i32* noalias [[DOTGLOBAL_TID_:%.*]], i32* noalias [[DOTBOUND_TID_:%.*]], i32* nonnull align 4 dereferenceable(4) [[I:%.*]]) #[[ATTR0]] {
788 // CHECK2-NEXT: entry:
789 // CHECK2-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 4
790 // CHECK2-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 4
791 // CHECK2-NEXT: [[I_ADDR:%.*]] = alloca i32*, align 4
792 // CHECK2-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 4
793 // CHECK2-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 4
794 // CHECK2-NEXT: store i32* [[I]], i32** [[I_ADDR]], align 4
795 // CHECK2-NEXT: [[TMP0:%.*]] = load i32*, i32** [[I_ADDR]], align 4
796 // CHECK2-NEXT: [[TMP1:%.*]] = load i32, i32* [[TMP0]], align 4
797 // CHECK2-NEXT: [[INC:%.*]] = add nsw i32 [[TMP1]], 1
798 // CHECK2-NEXT: store i32 [[INC]], i32* [[TMP0]], align 4
799 // CHECK2-NEXT: ret void
800 //
801 //
802 // CHECK2-LABEL: define {{[^@]+}}@__omp_outlined__1_wrapper
803 // CHECK2-SAME: (i16 zeroext [[TMP0:%.*]], i32 [[TMP1:%.*]]) #[[ATTR2:[0-9]+]] {
804 // CHECK2-NEXT: entry:
805 // CHECK2-NEXT: [[DOTADDR:%.*]] = alloca i16, align 2
806 // CHECK2-NEXT: [[DOTADDR1:%.*]] = alloca i32, align 4
807 // CHECK2-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4
808 // CHECK2-NEXT: [[GLOBAL_ARGS:%.*]] = alloca i8**, align 4
809 // CHECK2-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4
810 // CHECK2-NEXT: store i16 [[TMP0]], i16* [[DOTADDR]], align 2
811 // CHECK2-NEXT: store i32 [[TMP1]], i32* [[DOTADDR1]], align 4
812 // CHECK2-NEXT: call void @__kmpc_get_shared_variables(i8*** [[GLOBAL_ARGS]])
813 // CHECK2-NEXT: [[TMP2:%.*]] = load i8**, i8*** [[GLOBAL_ARGS]], align 4
814 // CHECK2-NEXT: [[TMP3:%.*]] = getelementptr inbounds i8*, i8** [[TMP2]], i32 0
815 // CHECK2-NEXT: [[TMP4:%.*]] = bitcast i8** [[TMP3]] to i32**
816 // CHECK2-NEXT: [[TMP5:%.*]] = load i32*, i32** [[TMP4]], align 4
817 // CHECK2-NEXT: call void @__omp_outlined__1(i32* [[DOTADDR1]], i32* [[DOTZERO_ADDR]], i32* [[TMP5]]) #[[ATTR1]]
818 // CHECK2-NEXT: ret void
819 //
820 //
821 // CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l16
822 // CHECK3-SAME: () #[[ATTR0:[0-9]+]] {
823 // CHECK3-NEXT: entry:
824 // CHECK3-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4
825 // CHECK3-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 4
826 // CHECK3-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4
827 // CHECK3-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_target_init(%struct.ident_t* @[[GLOB1:[0-9]+]], i1 false, i1 true, i1 true)
828 // CHECK3-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP0]], -1
829 // CHECK3-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
830 // CHECK3: user_code.entry:
831 // CHECK3-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1]])
832 // CHECK3-NEXT: store i32 [[TMP1]], i32* [[DOTTHREADID_TEMP_]], align 4
833 // CHECK3-NEXT: call void @__omp_outlined__(i32* [[DOTTHREADID_TEMP_]], i32* [[DOTZERO_ADDR]]) #[[ATTR1:[0-9]+]]
834 // CHECK3-NEXT: call void @__kmpc_target_deinit(%struct.ident_t* @[[GLOB1]], i1 false, i1 true)
835 // CHECK3-NEXT: ret void
836 // CHECK3: worker.exit:
837 // CHECK3-NEXT: ret void
838 //
839 //
840 // CHECK3-LABEL: define {{[^@]+}}@__omp_outlined__
841 // CHECK3-SAME: (i32* noalias [[DOTGLOBAL_TID_:%.*]], i32* noalias [[DOTBOUND_TID_:%.*]]) #[[ATTR0]] {
842 // CHECK3-NEXT: entry:
843 // CHECK3-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 4
844 // CHECK3-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 4
845 // CHECK3-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 4
846 // CHECK3-NEXT: [[TMP:%.*]] = alloca i32, align 4
847 // CHECK3-NEXT: [[DOTOMP_LB:%.*]] = alloca i32, align 4
848 // CHECK3-NEXT: [[DOTOMP_UB:%.*]] = alloca i32, align 4
849 // CHECK3-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4
850 // CHECK3-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4
851 // CHECK3-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [1 x i8*], align 4
852 // CHECK3-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 4
853 // CHECK3-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 4
854 // CHECK3-NEXT: [[I:%.*]] = call i8* @__kmpc_alloc_shared(i32 4)
855 // CHECK3-NEXT: [[I_ON_STACK:%.*]] = bitcast i8* [[I]] to i32*
856 // CHECK3-NEXT: store i32 0, i32* [[DOTOMP_LB]], align 4
857 // CHECK3-NEXT: store i32 9, i32* [[DOTOMP_UB]], align 4
858 // CHECK3-NEXT: store i32 1, i32* [[DOTOMP_STRIDE]], align 4
859 // CHECK3-NEXT: store i32 0, i32* [[DOTOMP_IS_LAST]], align 4
860 // CHECK3-NEXT: [[TMP0:%.*]] = load i32*, i32** [[DOTGLOBAL_TID__ADDR]], align 4
861 // CHECK3-NEXT: [[TMP1:%.*]] = load i32, i32* [[TMP0]], align 4
862 // CHECK3-NEXT: call void @__kmpc_for_static_init_4(%struct.ident_t* @[[GLOB2:[0-9]+]], i32 [[TMP1]], i32 92, i32* [[DOTOMP_IS_LAST]], i32* [[DOTOMP_LB]], i32* [[DOTOMP_UB]], i32* [[DOTOMP_STRIDE]], i32 1, i32 1)
863 // CHECK3-NEXT: [[TMP2:%.*]] = load i32, i32* [[DOTOMP_UB]], align 4
864 // CHECK3-NEXT: [[CMP:%.*]] = icmp sgt i32 [[TMP2]], 9
865 // CHECK3-NEXT: br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]
866 // CHECK3: cond.true:
867 // CHECK3-NEXT: br label [[COND_END:%.*]]
868 // CHECK3: cond.false:
869 // CHECK3-NEXT: [[TMP3:%.*]] = load i32, i32* [[DOTOMP_UB]], align 4
870 // CHECK3-NEXT: br label [[COND_END]]
871 // CHECK3: cond.end:
872 // CHECK3-NEXT: [[COND:%.*]] = phi i32 [ 9, [[COND_TRUE]] ], [ [[TMP3]], [[COND_FALSE]] ]
873 // CHECK3-NEXT: store i32 [[COND]], i32* [[DOTOMP_UB]], align 4
874 // CHECK3-NEXT: [[TMP4:%.*]] = load i32, i32* [[DOTOMP_LB]], align 4
875 // CHECK3-NEXT: store i32 [[TMP4]], i32* [[DOTOMP_IV]], align 4
876 // CHECK3-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]
877 // CHECK3: omp.inner.for.cond:
878 // CHECK3-NEXT: [[TMP5:%.*]] = load i32, i32* [[DOTOMP_IV]], align 4
879 // CHECK3-NEXT: [[TMP6:%.*]] = load i32, i32* [[DOTOMP_UB]], align 4
880 // CHECK3-NEXT: [[CMP1:%.*]] = icmp sle i32 [[TMP5]], [[TMP6]]
881 // CHECK3-NEXT: br i1 [[CMP1]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]
882 // CHECK3: omp.inner.for.body:
883 // CHECK3-NEXT: [[TMP7:%.*]] = load i32, i32* [[DOTOMP_IV]], align 4
884 // CHECK3-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP7]], 1
885 // CHECK3-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]
886 // CHECK3-NEXT: store i32 [[ADD]], i32* [[I_ON_STACK]], align 4
887 // CHECK3-NEXT: [[TMP8:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[CAPTURED_VARS_ADDRS]], i32 0, i32 0
888 // CHECK3-NEXT: [[TMP9:%.*]] = bitcast i32* [[I_ON_STACK]] to i8*
889 // CHECK3-NEXT: store i8* [[TMP9]], i8** [[TMP8]], align 4
890 // CHECK3-NEXT: [[TMP10:%.*]] = bitcast [1 x i8*]* [[CAPTURED_VARS_ADDRS]] to i8**
891 // CHECK3-NEXT: call void @__kmpc_parallel_51(%struct.ident_t* @[[GLOB1]], i32 [[TMP1]], i32 1, i32 -1, i32 -1, i8* bitcast (void (i32*, i32*, i32*)* @__omp_outlined__1 to i8*), i8* bitcast (void (i16, i32)* @__omp_outlined__1_wrapper to i8*), i8** [[TMP10]], i32 1)
892 // CHECK3-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]
893 // CHECK3: omp.body.continue:
894 // CHECK3-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]
895 // CHECK3: omp.inner.for.inc:
896 // CHECK3-NEXT: [[TMP11:%.*]] = load i32, i32* [[DOTOMP_IV]], align 4
897 // CHECK3-NEXT: [[ADD2:%.*]] = add nsw i32 [[TMP11]], 1
898 // CHECK3-NEXT: store i32 [[ADD2]], i32* [[DOTOMP_IV]], align 4
899 // CHECK3-NEXT: br label [[OMP_INNER_FOR_COND]]
900 // CHECK3: omp.inner.for.end:
901 // CHECK3-NEXT: br label [[OMP_LOOP_EXIT:%.*]]
902 // CHECK3: omp.loop.exit:
903 // CHECK3-NEXT: call void @__kmpc_for_static_fini(%struct.ident_t* @[[GLOB2]], i32 [[TMP1]])
904 // CHECK3-NEXT: call void @__kmpc_free_shared(i8* [[I]], i32 4)
905 // CHECK3-NEXT: ret void
906 //
907 //
908 // CHECK3-LABEL: define {{[^@]+}}@__omp_outlined__1
909 // CHECK3-SAME: (i32* noalias [[DOTGLOBAL_TID_:%.*]], i32* noalias [[DOTBOUND_TID_:%.*]], i32* nonnull align 4 dereferenceable(4) [[I:%.*]]) #[[ATTR0]] {
910 // CHECK3-NEXT: entry:
911 // CHECK3-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 4
912 // CHECK3-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 4
913 // CHECK3-NEXT: [[I_ADDR:%.*]] = alloca i32*, align 4
914 // CHECK3-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 4
915 // CHECK3-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 4
916 // CHECK3-NEXT: store i32* [[I]], i32** [[I_ADDR]], align 4
917 // CHECK3-NEXT: [[TMP0:%.*]] = load i32*, i32** [[I_ADDR]], align 4
918 // CHECK3-NEXT: [[TMP1:%.*]] = load i32, i32* [[TMP0]], align 4
919 // CHECK3-NEXT: [[INC:%.*]] = add nsw i32 [[TMP1]], 1
920 // CHECK3-NEXT: store i32 [[INC]], i32* [[TMP0]], align 4
921 // CHECK3-NEXT: ret void
922 //
923 //
924 // CHECK3-LABEL: define {{[^@]+}}@__omp_outlined__1_wrapper
925 // CHECK3-SAME: (i16 zeroext [[TMP0:%.*]], i32 [[TMP1:%.*]]) #[[ATTR2:[0-9]+]] {
926 // CHECK3-NEXT: entry:
927 // CHECK3-NEXT: [[DOTADDR:%.*]] = alloca i16, align 2
928 // CHECK3-NEXT: [[DOTADDR1:%.*]] = alloca i32, align 4
929 // CHECK3-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4
930 // CHECK3-NEXT: [[GLOBAL_ARGS:%.*]] = alloca i8**, align 4
931 // CHECK3-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4
932 // CHECK3-NEXT: store i16 [[TMP0]], i16* [[DOTADDR]], align 2
933 // CHECK3-NEXT: store i32 [[TMP1]], i32* [[DOTADDR1]], align 4
934 // CHECK3-NEXT: call void @__kmpc_get_shared_variables(i8*** [[GLOBAL_ARGS]])
935 // CHECK3-NEXT: [[TMP2:%.*]] = load i8**, i8*** [[GLOBAL_ARGS]], align 4
936 // CHECK3-NEXT: [[TMP3:%.*]] = getelementptr inbounds i8*, i8** [[TMP2]], i32 0
937 // CHECK3-NEXT: [[TMP4:%.*]] = bitcast i8** [[TMP3]] to i32**
938 // CHECK3-NEXT: [[TMP5:%.*]] = load i32*, i32** [[TMP4]], align 4
939 // CHECK3-NEXT: call void @__omp_outlined__1(i32* [[DOTADDR1]], i32* [[DOTZERO_ADDR]], i32* [[TMP5]]) #[[ATTR1]]
940 // CHECK3-NEXT: ret void
941 //
942