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 -no-opaque-pointers -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 -no-opaque-pointers -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 -no-opaque-pointers -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 -no-opaque-pointers -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 -no-opaque-pointers -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=CHECK2
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 // CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l16
36 // CHECK1-SAME: () #[[ATTR0:[0-9]+]] {
37 // CHECK1-NEXT: entry:
38 // CHECK1-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4
39 // CHECK1-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 4
40 // CHECK1-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_target_init(%struct.ident_t* @[[GLOB1:[0-9]+]], i8 1, i1 true, i1 true)
41 // CHECK1-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP0]], -1
42 // CHECK1-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
43 // CHECK1: user_code.entry:
44 // CHECK1-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1]])
45 // CHECK1-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4
46 // CHECK1-NEXT: store i32 [[TMP1]], i32* [[DOTTHREADID_TEMP_]], align 4
47 // CHECK1-NEXT: call void @__omp_outlined__(i32* [[DOTTHREADID_TEMP_]], i32* [[DOTZERO_ADDR]]) #[[ATTR5:[0-9]+]]
48 // CHECK1-NEXT: call void @__kmpc_target_deinit(%struct.ident_t* @[[GLOB1]], i8 1, i1 true)
49 // CHECK1-NEXT: ret void
50 // CHECK1: worker.exit:
51 // CHECK1-NEXT: ret void
52 //
53 //
54 // CHECK1-LABEL: define {{[^@]+}}@__omp_outlined__
55 // CHECK1-SAME: (i32* noalias noundef [[DOTGLOBAL_TID_:%.*]], i32* noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR1:[0-9]+]] {
56 // CHECK1-NEXT: entry:
57 // CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8
58 // CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8
59 // CHECK1-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 4
60 // CHECK1-NEXT: [[TMP:%.*]] = alloca i32, align 4
61 // CHECK1-NEXT: [[DOTOMP_LB:%.*]] = alloca i32, align 4
62 // CHECK1-NEXT: [[DOTOMP_UB:%.*]] = alloca i32, align 4
63 // CHECK1-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4
64 // CHECK1-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4
65 // CHECK1-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [1 x i8*], align 8
66 // CHECK1-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 8
67 // CHECK1-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 8
68 // CHECK1-NEXT: [[I:%.*]] = call align 8 i8* @__kmpc_alloc_shared(i64 4)
69 // CHECK1-NEXT: [[I_ON_STACK:%.*]] = bitcast i8* [[I]] to i32*
70 // CHECK1-NEXT: store i32 0, i32* [[DOTOMP_LB]], align 4
71 // CHECK1-NEXT: store i32 9, i32* [[DOTOMP_UB]], align 4
72 // CHECK1-NEXT: store i32 1, i32* [[DOTOMP_STRIDE]], align 4
73 // CHECK1-NEXT: store i32 0, i32* [[DOTOMP_IS_LAST]], align 4
74 // CHECK1-NEXT: [[TMP0:%.*]] = load i32*, i32** [[DOTGLOBAL_TID__ADDR]], align 8
75 // CHECK1-NEXT: [[TMP1:%.*]] = load i32, i32* [[TMP0]], align 4
76 // CHECK1-NEXT: call void @__kmpc_distribute_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)
77 // CHECK1-NEXT: [[TMP2:%.*]] = load i32, i32* [[DOTOMP_UB]], align 4
78 // CHECK1-NEXT: [[CMP:%.*]] = icmp sgt i32 [[TMP2]], 9
79 // CHECK1-NEXT: br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]
80 // CHECK1: cond.true:
81 // CHECK1-NEXT: br label [[COND_END:%.*]]
82 // CHECK1: cond.false:
83 // CHECK1-NEXT: [[TMP3:%.*]] = load i32, i32* [[DOTOMP_UB]], align 4
84 // CHECK1-NEXT: br label [[COND_END]]
85 // CHECK1: cond.end:
86 // CHECK1-NEXT: [[COND:%.*]] = phi i32 [ 9, [[COND_TRUE]] ], [ [[TMP3]], [[COND_FALSE]] ]
87 // CHECK1-NEXT: store i32 [[COND]], i32* [[DOTOMP_UB]], align 4
88 // CHECK1-NEXT: [[TMP4:%.*]] = load i32, i32* [[DOTOMP_LB]], align 4
89 // CHECK1-NEXT: store i32 [[TMP4]], i32* [[DOTOMP_IV]], align 4
90 // CHECK1-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]
91 // CHECK1: omp.inner.for.cond:
92 // CHECK1-NEXT: [[TMP5:%.*]] = load i32, i32* [[DOTOMP_IV]], align 4
93 // CHECK1-NEXT: [[TMP6:%.*]] = load i32, i32* [[DOTOMP_UB]], align 4
94 // CHECK1-NEXT: [[CMP1:%.*]] = icmp sle i32 [[TMP5]], [[TMP6]]
95 // CHECK1-NEXT: br i1 [[CMP1]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]
96 // CHECK1: omp.inner.for.body:
97 // CHECK1-NEXT: [[TMP7:%.*]] = load i32, i32* [[DOTOMP_IV]], align 4
98 // CHECK1-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP7]], 1
99 // CHECK1-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]
100 // CHECK1-NEXT: store i32 [[ADD]], i32* [[I_ON_STACK]], align 4
101 // CHECK1-NEXT: [[TMP8:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[CAPTURED_VARS_ADDRS]], i64 0, i64 0
102 // CHECK1-NEXT: [[TMP9:%.*]] = bitcast i32* [[I_ON_STACK]] to i8*
103 // CHECK1-NEXT: store i8* [[TMP9]], i8** [[TMP8]], align 8
104 // CHECK1-NEXT: [[TMP10:%.*]] = bitcast [1 x i8*]* [[CAPTURED_VARS_ADDRS]] to i8**
105 // 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)
106 // CHECK1-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]
107 // CHECK1: omp.body.continue:
108 // CHECK1-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]
109 // CHECK1: omp.inner.for.inc:
110 // CHECK1-NEXT: [[TMP11:%.*]] = load i32, i32* [[DOTOMP_IV]], align 4
111 // CHECK1-NEXT: [[ADD2:%.*]] = add nsw i32 [[TMP11]], 1
112 // CHECK1-NEXT: store i32 [[ADD2]], i32* [[DOTOMP_IV]], align 4
113 // CHECK1-NEXT: br label [[OMP_INNER_FOR_COND]]
114 // CHECK1: omp.inner.for.end:
115 // CHECK1-NEXT: br label [[OMP_LOOP_EXIT:%.*]]
116 // CHECK1: omp.loop.exit:
117 // CHECK1-NEXT: call void @__kmpc_distribute_static_fini(%struct.ident_t* @[[GLOB2]], i32 [[TMP1]])
118 // CHECK1-NEXT: call void @__kmpc_free_shared(i8* [[I]], i64 4)
119 // CHECK1-NEXT: ret void
120 //
121 //
122 // CHECK1-LABEL: define {{[^@]+}}@__omp_outlined__1
123 // CHECK1-SAME: (i32* noalias noundef [[DOTGLOBAL_TID_:%.*]], i32* noalias noundef [[DOTBOUND_TID_:%.*]], i32* noundef nonnull align 4 dereferenceable(4) [[I:%.*]]) #[[ATTR1]] {
124 // CHECK1-NEXT: entry:
125 // CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8
126 // CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8
127 // CHECK1-NEXT: [[I_ADDR:%.*]] = alloca i32*, align 8
128 // CHECK1-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 8
129 // CHECK1-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 8
130 // CHECK1-NEXT: store i32* [[I]], i32** [[I_ADDR]], align 8
131 // CHECK1-NEXT: [[TMP0:%.*]] = load i32*, i32** [[I_ADDR]], align 8
132 // CHECK1-NEXT: [[TMP1:%.*]] = load i32, i32* [[TMP0]], align 4
133 // CHECK1-NEXT: [[INC:%.*]] = add nsw i32 [[TMP1]], 1
134 // CHECK1-NEXT: store i32 [[INC]], i32* [[TMP0]], align 4
135 // CHECK1-NEXT: ret void
136 //
137 //
138 // CHECK1-LABEL: define {{[^@]+}}@__omp_outlined__1_wrapper
139 // CHECK1-SAME: (i16 noundef zeroext [[TMP0:%.*]], i32 noundef [[TMP1:%.*]]) #[[ATTR3:[0-9]+]] {
140 // CHECK1-NEXT: entry:
141 // CHECK1-NEXT: [[DOTADDR:%.*]] = alloca i16, align 2
142 // CHECK1-NEXT: [[DOTADDR1:%.*]] = alloca i32, align 4
143 // CHECK1-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4
144 // CHECK1-NEXT: [[GLOBAL_ARGS:%.*]] = alloca i8**, align 8
145 // CHECK1-NEXT: store i16 [[TMP0]], i16* [[DOTADDR]], align 2
146 // CHECK1-NEXT: store i32 [[TMP1]], i32* [[DOTADDR1]], align 4
147 // CHECK1-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4
148 // CHECK1-NEXT: call void @__kmpc_get_shared_variables(i8*** [[GLOBAL_ARGS]])
149 // CHECK1-NEXT: [[TMP2:%.*]] = load i8**, i8*** [[GLOBAL_ARGS]], align 8
150 // CHECK1-NEXT: [[TMP3:%.*]] = getelementptr inbounds i8*, i8** [[TMP2]], i64 0
151 // CHECK1-NEXT: [[TMP4:%.*]] = bitcast i8** [[TMP3]] to i32**
152 // CHECK1-NEXT: [[TMP5:%.*]] = load i32*, i32** [[TMP4]], align 8
153 // CHECK1-NEXT: call void @__omp_outlined__1(i32* [[DOTADDR1]], i32* [[DOTZERO_ADDR]], i32* [[TMP5]]) #[[ATTR5]]
154 // CHECK1-NEXT: ret void
155 //
156 //
157 // CHECK2-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l16
158 // CHECK2-SAME: () #[[ATTR0:[0-9]+]] {
159 // CHECK2-NEXT: entry:
160 // CHECK2-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4
161 // CHECK2-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 4
162 // CHECK2-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_target_init(%struct.ident_t* @[[GLOB1:[0-9]+]], i8 1, i1 true, i1 true)
163 // CHECK2-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP0]], -1
164 // CHECK2-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
165 // CHECK2: user_code.entry:
166 // CHECK2-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1]])
167 // CHECK2-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4
168 // CHECK2-NEXT: store i32 [[TMP1]], i32* [[DOTTHREADID_TEMP_]], align 4
169 // CHECK2-NEXT: call void @__omp_outlined__(i32* [[DOTTHREADID_TEMP_]], i32* [[DOTZERO_ADDR]]) #[[ATTR5:[0-9]+]]
170 // CHECK2-NEXT: call void @__kmpc_target_deinit(%struct.ident_t* @[[GLOB1]], i8 1, i1 true)
171 // CHECK2-NEXT: ret void
172 // CHECK2: worker.exit:
173 // CHECK2-NEXT: ret void
174 //
175 //
176 // CHECK2-LABEL: define {{[^@]+}}@__omp_outlined__
177 // CHECK2-SAME: (i32* noalias noundef [[DOTGLOBAL_TID_:%.*]], i32* noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR1:[0-9]+]] {
178 // CHECK2-NEXT: entry:
179 // CHECK2-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 4
180 // CHECK2-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 4
181 // CHECK2-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 4
182 // CHECK2-NEXT: [[TMP:%.*]] = alloca i32, align 4
183 // CHECK2-NEXT: [[DOTOMP_LB:%.*]] = alloca i32, align 4
184 // CHECK2-NEXT: [[DOTOMP_UB:%.*]] = alloca i32, align 4
185 // CHECK2-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4
186 // CHECK2-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4
187 // CHECK2-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [1 x i8*], align 4
188 // CHECK2-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 4
189 // CHECK2-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 4
190 // CHECK2-NEXT: [[I:%.*]] = call align 8 i8* @__kmpc_alloc_shared(i32 4)
191 // CHECK2-NEXT: [[I_ON_STACK:%.*]] = bitcast i8* [[I]] to i32*
192 // CHECK2-NEXT: store i32 0, i32* [[DOTOMP_LB]], align 4
193 // CHECK2-NEXT: store i32 9, i32* [[DOTOMP_UB]], align 4
194 // CHECK2-NEXT: store i32 1, i32* [[DOTOMP_STRIDE]], align 4
195 // CHECK2-NEXT: store i32 0, i32* [[DOTOMP_IS_LAST]], align 4
196 // CHECK2-NEXT: [[TMP0:%.*]] = load i32*, i32** [[DOTGLOBAL_TID__ADDR]], align 4
197 // CHECK2-NEXT: [[TMP1:%.*]] = load i32, i32* [[TMP0]], align 4
198 // CHECK2-NEXT: call void @__kmpc_distribute_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)
199 // CHECK2-NEXT: [[TMP2:%.*]] = load i32, i32* [[DOTOMP_UB]], align 4
200 // CHECK2-NEXT: [[CMP:%.*]] = icmp sgt i32 [[TMP2]], 9
201 // CHECK2-NEXT: br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]
202 // CHECK2: cond.true:
203 // CHECK2-NEXT: br label [[COND_END:%.*]]
204 // CHECK2: cond.false:
205 // CHECK2-NEXT: [[TMP3:%.*]] = load i32, i32* [[DOTOMP_UB]], align 4
206 // CHECK2-NEXT: br label [[COND_END]]
207 // CHECK2: cond.end:
208 // CHECK2-NEXT: [[COND:%.*]] = phi i32 [ 9, [[COND_TRUE]] ], [ [[TMP3]], [[COND_FALSE]] ]
209 // CHECK2-NEXT: store i32 [[COND]], i32* [[DOTOMP_UB]], align 4
210 // CHECK2-NEXT: [[TMP4:%.*]] = load i32, i32* [[DOTOMP_LB]], align 4
211 // CHECK2-NEXT: store i32 [[TMP4]], i32* [[DOTOMP_IV]], align 4
212 // CHECK2-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]
213 // CHECK2: omp.inner.for.cond:
214 // CHECK2-NEXT: [[TMP5:%.*]] = load i32, i32* [[DOTOMP_IV]], align 4
215 // CHECK2-NEXT: [[TMP6:%.*]] = load i32, i32* [[DOTOMP_UB]], align 4
216 // CHECK2-NEXT: [[CMP1:%.*]] = icmp sle i32 [[TMP5]], [[TMP6]]
217 // CHECK2-NEXT: br i1 [[CMP1]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]
218 // CHECK2: omp.inner.for.body:
219 // CHECK2-NEXT: [[TMP7:%.*]] = load i32, i32* [[DOTOMP_IV]], align 4
220 // CHECK2-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP7]], 1
221 // CHECK2-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]
222 // CHECK2-NEXT: store i32 [[ADD]], i32* [[I_ON_STACK]], align 4
223 // CHECK2-NEXT: [[TMP8:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[CAPTURED_VARS_ADDRS]], i32 0, i32 0
224 // CHECK2-NEXT: [[TMP9:%.*]] = bitcast i32* [[I_ON_STACK]] to i8*
225 // CHECK2-NEXT: store i8* [[TMP9]], i8** [[TMP8]], align 4
226 // CHECK2-NEXT: [[TMP10:%.*]] = bitcast [1 x i8*]* [[CAPTURED_VARS_ADDRS]] to i8**
227 // 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)
228 // CHECK2-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]
229 // CHECK2: omp.body.continue:
230 // CHECK2-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]
231 // CHECK2: omp.inner.for.inc:
232 // CHECK2-NEXT: [[TMP11:%.*]] = load i32, i32* [[DOTOMP_IV]], align 4
233 // CHECK2-NEXT: [[ADD2:%.*]] = add nsw i32 [[TMP11]], 1
234 // CHECK2-NEXT: store i32 [[ADD2]], i32* [[DOTOMP_IV]], align 4
235 // CHECK2-NEXT: br label [[OMP_INNER_FOR_COND]]
236 // CHECK2: omp.inner.for.end:
237 // CHECK2-NEXT: br label [[OMP_LOOP_EXIT:%.*]]
238 // CHECK2: omp.loop.exit:
239 // CHECK2-NEXT: call void @__kmpc_distribute_static_fini(%struct.ident_t* @[[GLOB2]], i32 [[TMP1]])
240 // CHECK2-NEXT: call void @__kmpc_free_shared(i8* [[I]], i32 4)
241 // CHECK2-NEXT: ret void
242 //
243 //
244 // CHECK2-LABEL: define {{[^@]+}}@__omp_outlined__1
245 // CHECK2-SAME: (i32* noalias noundef [[DOTGLOBAL_TID_:%.*]], i32* noalias noundef [[DOTBOUND_TID_:%.*]], i32* noundef nonnull align 4 dereferenceable(4) [[I:%.*]]) #[[ATTR1]] {
246 // CHECK2-NEXT: entry:
247 // CHECK2-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 4
248 // CHECK2-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 4
249 // CHECK2-NEXT: [[I_ADDR:%.*]] = alloca i32*, align 4
250 // CHECK2-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 4
251 // CHECK2-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 4
252 // CHECK2-NEXT: store i32* [[I]], i32** [[I_ADDR]], align 4
253 // CHECK2-NEXT: [[TMP0:%.*]] = load i32*, i32** [[I_ADDR]], align 4
254 // CHECK2-NEXT: [[TMP1:%.*]] = load i32, i32* [[TMP0]], align 4
255 // CHECK2-NEXT: [[INC:%.*]] = add nsw i32 [[TMP1]], 1
256 // CHECK2-NEXT: store i32 [[INC]], i32* [[TMP0]], align 4
257 // CHECK2-NEXT: ret void
258 //
259 //
260 // CHECK2-LABEL: define {{[^@]+}}@__omp_outlined__1_wrapper
261 // CHECK2-SAME: (i16 noundef zeroext [[TMP0:%.*]], i32 noundef [[TMP1:%.*]]) #[[ATTR3:[0-9]+]] {
262 // CHECK2-NEXT: entry:
263 // CHECK2-NEXT: [[DOTADDR:%.*]] = alloca i16, align 2
264 // CHECK2-NEXT: [[DOTADDR1:%.*]] = alloca i32, align 4
265 // CHECK2-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4
266 // CHECK2-NEXT: [[GLOBAL_ARGS:%.*]] = alloca i8**, align 4
267 // CHECK2-NEXT: store i16 [[TMP0]], i16* [[DOTADDR]], align 2
268 // CHECK2-NEXT: store i32 [[TMP1]], i32* [[DOTADDR1]], align 4
269 // CHECK2-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4
270 // CHECK2-NEXT: call void @__kmpc_get_shared_variables(i8*** [[GLOBAL_ARGS]])
271 // CHECK2-NEXT: [[TMP2:%.*]] = load i8**, i8*** [[GLOBAL_ARGS]], align 4
272 // CHECK2-NEXT: [[TMP3:%.*]] = getelementptr inbounds i8*, i8** [[TMP2]], i32 0
273 // CHECK2-NEXT: [[TMP4:%.*]] = bitcast i8** [[TMP3]] to i32**
274 // CHECK2-NEXT: [[TMP5:%.*]] = load i32*, i32** [[TMP4]], align 4
275 // CHECK2-NEXT: call void @__omp_outlined__1(i32* [[DOTADDR1]], i32* [[DOTZERO_ADDR]], i32* [[TMP5]]) #[[ATTR5]]
276 // CHECK2-NEXT: ret void
277 //
278