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>
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 
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