xref: /llvm-project/clang/test/OpenMP/distribute_simd_reduction_codegen.cpp (revision 94473f4db6a6f5f12d7c4081455b5b596094eac5) 
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 // RUN: %clang_cc1 -DCHECK -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix=CHECK1
3 // RUN: %clang_cc1 -DCHECK -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
4 // RUN: %clang_cc1 -DCHECK -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix=CHECK1
5 // RUN: %clang_cc1 -DCHECK -verify -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix=CHECK3
6 // RUN: %clang_cc1 -DCHECK -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s
7 // RUN: %clang_cc1 -DCHECK -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix=CHECK3
8 
9 // RUN: %clang_cc1 -DCHECK -verify -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix=CHECK5
10 // RUN: %clang_cc1 -DCHECK -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
11 // RUN: %clang_cc1 -DCHECK -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix=CHECK5
12 // RUN: %clang_cc1 -DCHECK -verify -fopenmp-simd -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix=CHECK7
13 // RUN: %clang_cc1 -DCHECK -fopenmp-simd -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s
14 // RUN: %clang_cc1 -DCHECK -fopenmp-simd -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix=CHECK7
15 
16 // RUN: %clang_cc1 -DLAMBDA -verify -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix=CHECK9
17 // RUN: %clang_cc1 -DLAMBDA -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
18 // RUN: %clang_cc1 -DLAMBDA -fopenmp -x c++  -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix=CHECK9
19 
20 // RUN: %clang_cc1 -DLAMBDA -verify -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix=CHECK11
21 // RUN: %clang_cc1 -DLAMBDA -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
22 // RUN: %clang_cc1 -DLAMBDA -fopenmp-simd -x c++  -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix=CHECK11
23 
24 // expected-no-diagnostics
25 #ifndef HEADER
26 #define HEADER
27 
28 template <typename T>
29 T tmain() {
30   T t_var = T();
31   T vec[] = {1, 2};
32 #pragma omp target
33 #pragma omp teams
34 #pragma omp distribute simd reduction(+: t_var)
35   for (int i = 0; i < 2; ++i) {
36     t_var += (T) i;
37   }
38   return T();
39 }
40 
41 int main() {
42   static int sivar;
43 #ifdef LAMBDA
44   [&]() {
45 #pragma omp target
46 #pragma omp teams
47 #pragma omp distribute simd reduction(+: sivar)
48   for (int i = 0; i < 2; ++i) {
49 
50     // Skip global and bound tid vars
51 
52 
53     sivar += i;
54 
55     [&]() {
56 
57       sivar += 4;
58 
59     }();
60   }
61   }();
62   return 0;
63 #else
64 #pragma omp target
65 #pragma omp teams
66 #pragma omp distribute simd reduction(+: sivar)
67   for (int i = 0; i < 2; ++i) {
68     sivar += i;
69   }
70   return tmain<int>();
71 #endif
72 }
73 
74 
75 
76 // Skip global and bound tid vars
77 
78 
79 
80 
81 // Skip global and bound tid vars
82 
83 
84 #endif
85 // CHECK1-LABEL: define {{[^@]+}}@main
86 // CHECK1-SAME: () #[[ATTR0:[0-9]+]] {
87 // CHECK1-NEXT:  entry:
88 // CHECK1-NEXT:    [[RETVAL:%.*]] = alloca i32, align 4
89 // CHECK1-NEXT:    [[SIVAR_CASTED:%.*]] = alloca i64, align 8
90 // CHECK1-NEXT:    [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x ptr], align 8
91 // CHECK1-NEXT:    [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x ptr], align 8
92 // CHECK1-NEXT:    [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [1 x ptr], align 8
93 // CHECK1-NEXT:    [[TMP:%.*]] = alloca i32, align 4
94 // CHECK1-NEXT:    [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8
95 // CHECK1-NEXT:    store i32 0, ptr [[RETVAL]], align 4
96 // CHECK1-NEXT:    [[TMP0:%.*]] = load i32, ptr @_ZZ4mainE5sivar, align 4
97 // CHECK1-NEXT:    store i32 [[TMP0]], ptr [[SIVAR_CASTED]], align 4
98 // CHECK1-NEXT:    [[TMP1:%.*]] = load i64, ptr [[SIVAR_CASTED]], align 8
99 // CHECK1-NEXT:    [[TMP2:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
100 // CHECK1-NEXT:    store i64 [[TMP1]], ptr [[TMP2]], align 8
101 // CHECK1-NEXT:    [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0
102 // CHECK1-NEXT:    store i64 [[TMP1]], ptr [[TMP3]], align 8
103 // CHECK1-NEXT:    [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 0
104 // CHECK1-NEXT:    store ptr null, ptr [[TMP4]], align 8
105 // CHECK1-NEXT:    [[TMP5:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
106 // CHECK1-NEXT:    [[TMP6:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0
107 // CHECK1-NEXT:    [[TMP7:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0
108 // CHECK1-NEXT:    store i32 3, ptr [[TMP7]], align 4
109 // CHECK1-NEXT:    [[TMP8:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1
110 // CHECK1-NEXT:    store i32 1, ptr [[TMP8]], align 4
111 // CHECK1-NEXT:    [[TMP9:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2
112 // CHECK1-NEXT:    store ptr [[TMP5]], ptr [[TMP9]], align 8
113 // CHECK1-NEXT:    [[TMP10:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3
114 // CHECK1-NEXT:    store ptr [[TMP6]], ptr [[TMP10]], align 8
115 // CHECK1-NEXT:    [[TMP11:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4
116 // CHECK1-NEXT:    store ptr @.offload_sizes, ptr [[TMP11]], align 8
117 // CHECK1-NEXT:    [[TMP12:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5
118 // CHECK1-NEXT:    store ptr @.offload_maptypes, ptr [[TMP12]], align 8
119 // CHECK1-NEXT:    [[TMP13:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6
120 // CHECK1-NEXT:    store ptr null, ptr [[TMP13]], align 8
121 // CHECK1-NEXT:    [[TMP14:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7
122 // CHECK1-NEXT:    store ptr null, ptr [[TMP14]], align 8
123 // CHECK1-NEXT:    [[TMP15:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8
124 // CHECK1-NEXT:    store i64 2, ptr [[TMP15]], align 8
125 // CHECK1-NEXT:    [[TMP16:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9
126 // CHECK1-NEXT:    store i64 0, ptr [[TMP16]], align 8
127 // CHECK1-NEXT:    [[TMP17:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10
128 // CHECK1-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP17]], align 4
129 // CHECK1-NEXT:    [[TMP18:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11
130 // CHECK1-NEXT:    store [3 x i32] [i32 1, i32 0, i32 0], ptr [[TMP18]], align 4
131 // CHECK1-NEXT:    [[TMP19:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12
132 // CHECK1-NEXT:    store i32 0, ptr [[TMP19]], align 4
133 // CHECK1-NEXT:    [[TMP20:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB2:[0-9]+]], i64 -1, i32 0, i32 1, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l64.region_id, ptr [[KERNEL_ARGS]])
134 // CHECK1-NEXT:    [[TMP21:%.*]] = icmp ne i32 [[TMP20]], 0
135 // CHECK1-NEXT:    br i1 [[TMP21]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]
136 // CHECK1:       omp_offload.failed:
137 // CHECK1-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l64(i64 [[TMP1]]) #[[ATTR2:[0-9]+]]
138 // CHECK1-NEXT:    br label [[OMP_OFFLOAD_CONT]]
139 // CHECK1:       omp_offload.cont:
140 // CHECK1-NEXT:    [[CALL:%.*]] = call noundef signext i32 @_Z5tmainIiET_v()
141 // CHECK1-NEXT:    ret i32 [[CALL]]
142 //
143 //
144 // CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l64
145 // CHECK1-SAME: (i64 noundef [[SIVAR:%.*]]) #[[ATTR1:[0-9]+]] {
146 // CHECK1-NEXT:  entry:
147 // CHECK1-NEXT:    [[SIVAR_ADDR:%.*]] = alloca i64, align 8
148 // CHECK1-NEXT:    store i64 [[SIVAR]], ptr [[SIVAR_ADDR]], align 8
149 // CHECK1-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB2]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l64.omp_outlined, ptr [[SIVAR_ADDR]])
150 // CHECK1-NEXT:    ret void
151 //
152 //
153 // CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l64.omp_outlined
154 // CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[SIVAR:%.*]]) #[[ATTR1]] {
155 // CHECK1-NEXT:  entry:
156 // CHECK1-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8
157 // CHECK1-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8
158 // CHECK1-NEXT:    [[SIVAR_ADDR:%.*]] = alloca ptr, align 8
159 // CHECK1-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 4
160 // CHECK1-NEXT:    [[TMP:%.*]] = alloca i32, align 4
161 // CHECK1-NEXT:    [[DOTOMP_LB:%.*]] = alloca i32, align 4
162 // CHECK1-NEXT:    [[DOTOMP_UB:%.*]] = alloca i32, align 4
163 // CHECK1-NEXT:    [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4
164 // CHECK1-NEXT:    [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4
165 // CHECK1-NEXT:    [[SIVAR1:%.*]] = alloca i32, align 4
166 // CHECK1-NEXT:    [[I:%.*]] = alloca i32, align 4
167 // CHECK1-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8
168 // CHECK1-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8
169 // CHECK1-NEXT:    store ptr [[SIVAR]], ptr [[SIVAR_ADDR]], align 8
170 // CHECK1-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[SIVAR_ADDR]], align 8
171 // CHECK1-NEXT:    store i32 0, ptr [[DOTOMP_LB]], align 4
172 // CHECK1-NEXT:    store i32 1, ptr [[DOTOMP_UB]], align 4
173 // CHECK1-NEXT:    store i32 1, ptr [[DOTOMP_STRIDE]], align 4
174 // CHECK1-NEXT:    store i32 0, ptr [[DOTOMP_IS_LAST]], align 4
175 // CHECK1-NEXT:    store i32 0, ptr [[SIVAR1]], align 4
176 // CHECK1-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8
177 // CHECK1-NEXT:    [[TMP2:%.*]] = load i32, ptr [[TMP1]], align 4
178 // CHECK1-NEXT:    call void @__kmpc_for_static_init_4(ptr @[[GLOB1:[0-9]+]], i32 [[TMP2]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)
179 // CHECK1-NEXT:    [[TMP3:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4
180 // CHECK1-NEXT:    [[CMP:%.*]] = icmp sgt i32 [[TMP3]], 1
181 // CHECK1-NEXT:    br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]
182 // CHECK1:       cond.true:
183 // CHECK1-NEXT:    br label [[COND_END:%.*]]
184 // CHECK1:       cond.false:
185 // CHECK1-NEXT:    [[TMP4:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4
186 // CHECK1-NEXT:    br label [[COND_END]]
187 // CHECK1:       cond.end:
188 // CHECK1-NEXT:    [[COND:%.*]] = phi i32 [ 1, [[COND_TRUE]] ], [ [[TMP4]], [[COND_FALSE]] ]
189 // CHECK1-NEXT:    store i32 [[COND]], ptr [[DOTOMP_UB]], align 4
190 // CHECK1-NEXT:    [[TMP5:%.*]] = load i32, ptr [[DOTOMP_LB]], align 4
191 // CHECK1-NEXT:    store i32 [[TMP5]], ptr [[DOTOMP_IV]], align 4
192 // CHECK1-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]
193 // CHECK1:       omp.inner.for.cond:
194 // CHECK1-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP5:![0-9]+]]
195 // CHECK1-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4, !llvm.access.group [[ACC_GRP5]]
196 // CHECK1-NEXT:    [[CMP2:%.*]] = icmp sle i32 [[TMP6]], [[TMP7]]
197 // CHECK1-NEXT:    br i1 [[CMP2]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]
198 // CHECK1:       omp.inner.for.body:
199 // CHECK1-NEXT:    [[TMP8:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP5]]
200 // CHECK1-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP8]], 1
201 // CHECK1-NEXT:    [[ADD:%.*]] = add nsw i32 0, [[MUL]]
202 // CHECK1-NEXT:    store i32 [[ADD]], ptr [[I]], align 4, !llvm.access.group [[ACC_GRP5]]
203 // CHECK1-NEXT:    [[TMP9:%.*]] = load i32, ptr [[I]], align 4, !llvm.access.group [[ACC_GRP5]]
204 // CHECK1-NEXT:    [[TMP10:%.*]] = load i32, ptr [[SIVAR1]], align 4, !llvm.access.group [[ACC_GRP5]]
205 // CHECK1-NEXT:    [[ADD3:%.*]] = add nsw i32 [[TMP10]], [[TMP9]]
206 // CHECK1-NEXT:    store i32 [[ADD3]], ptr [[SIVAR1]], align 4, !llvm.access.group [[ACC_GRP5]]
207 // CHECK1-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]
208 // CHECK1:       omp.body.continue:
209 // CHECK1-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]
210 // CHECK1:       omp.inner.for.inc:
211 // CHECK1-NEXT:    [[TMP11:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP5]]
212 // CHECK1-NEXT:    [[ADD4:%.*]] = add nsw i32 [[TMP11]], 1
213 // CHECK1-NEXT:    store i32 [[ADD4]], ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP5]]
214 // CHECK1-NEXT:    br label [[OMP_INNER_FOR_COND]], !llvm.loop [[LOOP6:![0-9]+]]
215 // CHECK1:       omp.inner.for.end:
216 // CHECK1-NEXT:    br label [[OMP_LOOP_EXIT:%.*]]
217 // CHECK1:       omp.loop.exit:
218 // CHECK1-NEXT:    call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP2]])
219 // CHECK1-NEXT:    [[TMP12:%.*]] = load i32, ptr [[DOTOMP_IS_LAST]], align 4
220 // CHECK1-NEXT:    [[TMP13:%.*]] = icmp ne i32 [[TMP12]], 0
221 // CHECK1-NEXT:    br i1 [[TMP13]], label [[DOTOMP_FINAL_THEN:%.*]], label [[DOTOMP_FINAL_DONE:%.*]]
222 // CHECK1:       .omp.final.then:
223 // CHECK1-NEXT:    store i32 2, ptr [[I]], align 4
224 // CHECK1-NEXT:    br label [[DOTOMP_FINAL_DONE]]
225 // CHECK1:       .omp.final.done:
226 // CHECK1-NEXT:    [[TMP14:%.*]] = load i32, ptr [[TMP0]], align 4
227 // CHECK1-NEXT:    [[TMP15:%.*]] = load i32, ptr [[SIVAR1]], align 4
228 // CHECK1-NEXT:    [[ADD5:%.*]] = add nsw i32 [[TMP14]], [[TMP15]]
229 // CHECK1-NEXT:    store i32 [[ADD5]], ptr [[TMP0]], align 4
230 // CHECK1-NEXT:    ret void
231 //
232 //
233 // CHECK1-LABEL: define {{[^@]+}}@_Z5tmainIiET_v
234 // CHECK1-SAME: () #[[ATTR3:[0-9]+]] comdat {
235 // CHECK1-NEXT:  entry:
236 // CHECK1-NEXT:    [[T_VAR:%.*]] = alloca i32, align 4
237 // CHECK1-NEXT:    [[VEC:%.*]] = alloca [2 x i32], align 4
238 // CHECK1-NEXT:    [[T_VAR_CASTED:%.*]] = alloca i64, align 8
239 // CHECK1-NEXT:    [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x ptr], align 8
240 // CHECK1-NEXT:    [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x ptr], align 8
241 // CHECK1-NEXT:    [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [1 x ptr], align 8
242 // CHECK1-NEXT:    [[TMP:%.*]] = alloca i32, align 4
243 // CHECK1-NEXT:    [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8
244 // CHECK1-NEXT:    store i32 0, ptr [[T_VAR]], align 4
245 // CHECK1-NEXT:    call void @llvm.memcpy.p0.p0.i64(ptr align 4 [[VEC]], ptr align 4 @__const._Z5tmainIiET_v.vec, i64 8, i1 false)
246 // CHECK1-NEXT:    [[TMP0:%.*]] = load i32, ptr [[T_VAR]], align 4
247 // CHECK1-NEXT:    store i32 [[TMP0]], ptr [[T_VAR_CASTED]], align 4
248 // CHECK1-NEXT:    [[TMP1:%.*]] = load i64, ptr [[T_VAR_CASTED]], align 8
249 // CHECK1-NEXT:    [[TMP2:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
250 // CHECK1-NEXT:    store i64 [[TMP1]], ptr [[TMP2]], align 8
251 // CHECK1-NEXT:    [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0
252 // CHECK1-NEXT:    store i64 [[TMP1]], ptr [[TMP3]], align 8
253 // CHECK1-NEXT:    [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 0
254 // CHECK1-NEXT:    store ptr null, ptr [[TMP4]], align 8
255 // CHECK1-NEXT:    [[TMP5:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
256 // CHECK1-NEXT:    [[TMP6:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0
257 // CHECK1-NEXT:    [[TMP7:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0
258 // CHECK1-NEXT:    store i32 3, ptr [[TMP7]], align 4
259 // CHECK1-NEXT:    [[TMP8:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1
260 // CHECK1-NEXT:    store i32 1, ptr [[TMP8]], align 4
261 // CHECK1-NEXT:    [[TMP9:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2
262 // CHECK1-NEXT:    store ptr [[TMP5]], ptr [[TMP9]], align 8
263 // CHECK1-NEXT:    [[TMP10:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3
264 // CHECK1-NEXT:    store ptr [[TMP6]], ptr [[TMP10]], align 8
265 // CHECK1-NEXT:    [[TMP11:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4
266 // CHECK1-NEXT:    store ptr @.offload_sizes.1, ptr [[TMP11]], align 8
267 // CHECK1-NEXT:    [[TMP12:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5
268 // CHECK1-NEXT:    store ptr @.offload_maptypes.2, ptr [[TMP12]], align 8
269 // CHECK1-NEXT:    [[TMP13:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6
270 // CHECK1-NEXT:    store ptr null, ptr [[TMP13]], align 8
271 // CHECK1-NEXT:    [[TMP14:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7
272 // CHECK1-NEXT:    store ptr null, ptr [[TMP14]], align 8
273 // CHECK1-NEXT:    [[TMP15:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8
274 // CHECK1-NEXT:    store i64 2, ptr [[TMP15]], align 8
275 // CHECK1-NEXT:    [[TMP16:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9
276 // CHECK1-NEXT:    store i64 0, ptr [[TMP16]], align 8
277 // CHECK1-NEXT:    [[TMP17:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10
278 // CHECK1-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP17]], align 4
279 // CHECK1-NEXT:    [[TMP18:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11
280 // CHECK1-NEXT:    store [3 x i32] [i32 1, i32 0, i32 0], ptr [[TMP18]], align 4
281 // CHECK1-NEXT:    [[TMP19:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12
282 // CHECK1-NEXT:    store i32 0, ptr [[TMP19]], align 4
283 // CHECK1-NEXT:    [[TMP20:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB2]], i64 -1, i32 0, i32 1, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_v_l32.region_id, ptr [[KERNEL_ARGS]])
284 // CHECK1-NEXT:    [[TMP21:%.*]] = icmp ne i32 [[TMP20]], 0
285 // CHECK1-NEXT:    br i1 [[TMP21]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]
286 // CHECK1:       omp_offload.failed:
287 // CHECK1-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_v_l32(i64 [[TMP1]]) #[[ATTR2]]
288 // CHECK1-NEXT:    br label [[OMP_OFFLOAD_CONT]]
289 // CHECK1:       omp_offload.cont:
290 // CHECK1-NEXT:    ret i32 0
291 //
292 //
293 // CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_v_l32
294 // CHECK1-SAME: (i64 noundef [[T_VAR:%.*]]) #[[ATTR1]] {
295 // CHECK1-NEXT:  entry:
296 // CHECK1-NEXT:    [[T_VAR_ADDR:%.*]] = alloca i64, align 8
297 // CHECK1-NEXT:    store i64 [[T_VAR]], ptr [[T_VAR_ADDR]], align 8
298 // CHECK1-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB2]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_v_l32.omp_outlined, ptr [[T_VAR_ADDR]])
299 // CHECK1-NEXT:    ret void
300 //
301 //
302 // CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_v_l32.omp_outlined
303 // CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[T_VAR:%.*]]) #[[ATTR1]] {
304 // CHECK1-NEXT:  entry:
305 // CHECK1-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8
306 // CHECK1-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8
307 // CHECK1-NEXT:    [[T_VAR_ADDR:%.*]] = alloca ptr, align 8
308 // CHECK1-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 4
309 // CHECK1-NEXT:    [[TMP:%.*]] = alloca i32, align 4
310 // CHECK1-NEXT:    [[DOTOMP_LB:%.*]] = alloca i32, align 4
311 // CHECK1-NEXT:    [[DOTOMP_UB:%.*]] = alloca i32, align 4
312 // CHECK1-NEXT:    [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4
313 // CHECK1-NEXT:    [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4
314 // CHECK1-NEXT:    [[T_VAR1:%.*]] = alloca i32, align 4
315 // CHECK1-NEXT:    [[I:%.*]] = alloca i32, align 4
316 // CHECK1-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8
317 // CHECK1-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8
318 // CHECK1-NEXT:    store ptr [[T_VAR]], ptr [[T_VAR_ADDR]], align 8
319 // CHECK1-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[T_VAR_ADDR]], align 8
320 // CHECK1-NEXT:    store i32 0, ptr [[DOTOMP_LB]], align 4
321 // CHECK1-NEXT:    store i32 1, ptr [[DOTOMP_UB]], align 4
322 // CHECK1-NEXT:    store i32 1, ptr [[DOTOMP_STRIDE]], align 4
323 // CHECK1-NEXT:    store i32 0, ptr [[DOTOMP_IS_LAST]], align 4
324 // CHECK1-NEXT:    store i32 0, ptr [[T_VAR1]], align 4
325 // CHECK1-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8
326 // CHECK1-NEXT:    [[TMP2:%.*]] = load i32, ptr [[TMP1]], align 4
327 // CHECK1-NEXT:    call void @__kmpc_for_static_init_4(ptr @[[GLOB1]], i32 [[TMP2]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)
328 // CHECK1-NEXT:    [[TMP3:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4
329 // CHECK1-NEXT:    [[CMP:%.*]] = icmp sgt i32 [[TMP3]], 1
330 // CHECK1-NEXT:    br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]
331 // CHECK1:       cond.true:
332 // CHECK1-NEXT:    br label [[COND_END:%.*]]
333 // CHECK1:       cond.false:
334 // CHECK1-NEXT:    [[TMP4:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4
335 // CHECK1-NEXT:    br label [[COND_END]]
336 // CHECK1:       cond.end:
337 // CHECK1-NEXT:    [[COND:%.*]] = phi i32 [ 1, [[COND_TRUE]] ], [ [[TMP4]], [[COND_FALSE]] ]
338 // CHECK1-NEXT:    store i32 [[COND]], ptr [[DOTOMP_UB]], align 4
339 // CHECK1-NEXT:    [[TMP5:%.*]] = load i32, ptr [[DOTOMP_LB]], align 4
340 // CHECK1-NEXT:    store i32 [[TMP5]], ptr [[DOTOMP_IV]], align 4
341 // CHECK1-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]
342 // CHECK1:       omp.inner.for.cond:
343 // CHECK1-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP11:![0-9]+]]
344 // CHECK1-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4, !llvm.access.group [[ACC_GRP11]]
345 // CHECK1-NEXT:    [[CMP2:%.*]] = icmp sle i32 [[TMP6]], [[TMP7]]
346 // CHECK1-NEXT:    br i1 [[CMP2]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]
347 // CHECK1:       omp.inner.for.body:
348 // CHECK1-NEXT:    [[TMP8:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP11]]
349 // CHECK1-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP8]], 1
350 // CHECK1-NEXT:    [[ADD:%.*]] = add nsw i32 0, [[MUL]]
351 // CHECK1-NEXT:    store i32 [[ADD]], ptr [[I]], align 4, !llvm.access.group [[ACC_GRP11]]
352 // CHECK1-NEXT:    [[TMP9:%.*]] = load i32, ptr [[I]], align 4, !llvm.access.group [[ACC_GRP11]]
353 // CHECK1-NEXT:    [[TMP10:%.*]] = load i32, ptr [[T_VAR1]], align 4, !llvm.access.group [[ACC_GRP11]]
354 // CHECK1-NEXT:    [[ADD3:%.*]] = add nsw i32 [[TMP10]], [[TMP9]]
355 // CHECK1-NEXT:    store i32 [[ADD3]], ptr [[T_VAR1]], align 4, !llvm.access.group [[ACC_GRP11]]
356 // CHECK1-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]
357 // CHECK1:       omp.body.continue:
358 // CHECK1-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]
359 // CHECK1:       omp.inner.for.inc:
360 // CHECK1-NEXT:    [[TMP11:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP11]]
361 // CHECK1-NEXT:    [[ADD4:%.*]] = add nsw i32 [[TMP11]], 1
362 // CHECK1-NEXT:    store i32 [[ADD4]], ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP11]]
363 // CHECK1-NEXT:    br label [[OMP_INNER_FOR_COND]], !llvm.loop [[LOOP12:![0-9]+]]
364 // CHECK1:       omp.inner.for.end:
365 // CHECK1-NEXT:    br label [[OMP_LOOP_EXIT:%.*]]
366 // CHECK1:       omp.loop.exit:
367 // CHECK1-NEXT:    call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP2]])
368 // CHECK1-NEXT:    [[TMP12:%.*]] = load i32, ptr [[DOTOMP_IS_LAST]], align 4
369 // CHECK1-NEXT:    [[TMP13:%.*]] = icmp ne i32 [[TMP12]], 0
370 // CHECK1-NEXT:    br i1 [[TMP13]], label [[DOTOMP_FINAL_THEN:%.*]], label [[DOTOMP_FINAL_DONE:%.*]]
371 // CHECK1:       .omp.final.then:
372 // CHECK1-NEXT:    store i32 2, ptr [[I]], align 4
373 // CHECK1-NEXT:    br label [[DOTOMP_FINAL_DONE]]
374 // CHECK1:       .omp.final.done:
375 // CHECK1-NEXT:    [[TMP14:%.*]] = load i32, ptr [[TMP0]], align 4
376 // CHECK1-NEXT:    [[TMP15:%.*]] = load i32, ptr [[T_VAR1]], align 4
377 // CHECK1-NEXT:    [[ADD5:%.*]] = add nsw i32 [[TMP14]], [[TMP15]]
378 // CHECK1-NEXT:    store i32 [[ADD5]], ptr [[TMP0]], align 4
379 // CHECK1-NEXT:    ret void
380 //
381 //
382 // CHECK3-LABEL: define {{[^@]+}}@main
383 // CHECK3-SAME: () #[[ATTR0:[0-9]+]] {
384 // CHECK3-NEXT:  entry:
385 // CHECK3-NEXT:    [[RETVAL:%.*]] = alloca i32, align 4
386 // CHECK3-NEXT:    [[SIVAR_CASTED:%.*]] = alloca i32, align 4
387 // CHECK3-NEXT:    [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x ptr], align 4
388 // CHECK3-NEXT:    [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x ptr], align 4
389 // CHECK3-NEXT:    [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [1 x ptr], align 4
390 // CHECK3-NEXT:    [[TMP:%.*]] = alloca i32, align 4
391 // CHECK3-NEXT:    [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8
392 // CHECK3-NEXT:    store i32 0, ptr [[RETVAL]], align 4
393 // CHECK3-NEXT:    [[TMP0:%.*]] = load i32, ptr @_ZZ4mainE5sivar, align 4
394 // CHECK3-NEXT:    store i32 [[TMP0]], ptr [[SIVAR_CASTED]], align 4
395 // CHECK3-NEXT:    [[TMP1:%.*]] = load i32, ptr [[SIVAR_CASTED]], align 4
396 // CHECK3-NEXT:    [[TMP2:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
397 // CHECK3-NEXT:    store i32 [[TMP1]], ptr [[TMP2]], align 4
398 // CHECK3-NEXT:    [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0
399 // CHECK3-NEXT:    store i32 [[TMP1]], ptr [[TMP3]], align 4
400 // CHECK3-NEXT:    [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 0
401 // CHECK3-NEXT:    store ptr null, ptr [[TMP4]], align 4
402 // CHECK3-NEXT:    [[TMP5:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
403 // CHECK3-NEXT:    [[TMP6:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0
404 // CHECK3-NEXT:    [[TMP7:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0
405 // CHECK3-NEXT:    store i32 3, ptr [[TMP7]], align 4
406 // CHECK3-NEXT:    [[TMP8:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1
407 // CHECK3-NEXT:    store i32 1, ptr [[TMP8]], align 4
408 // CHECK3-NEXT:    [[TMP9:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2
409 // CHECK3-NEXT:    store ptr [[TMP5]], ptr [[TMP9]], align 4
410 // CHECK3-NEXT:    [[TMP10:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3
411 // CHECK3-NEXT:    store ptr [[TMP6]], ptr [[TMP10]], align 4
412 // CHECK3-NEXT:    [[TMP11:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4
413 // CHECK3-NEXT:    store ptr @.offload_sizes, ptr [[TMP11]], align 4
414 // CHECK3-NEXT:    [[TMP12:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5
415 // CHECK3-NEXT:    store ptr @.offload_maptypes, ptr [[TMP12]], align 4
416 // CHECK3-NEXT:    [[TMP13:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6
417 // CHECK3-NEXT:    store ptr null, ptr [[TMP13]], align 4
418 // CHECK3-NEXT:    [[TMP14:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7
419 // CHECK3-NEXT:    store ptr null, ptr [[TMP14]], align 4
420 // CHECK3-NEXT:    [[TMP15:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8
421 // CHECK3-NEXT:    store i64 2, ptr [[TMP15]], align 8
422 // CHECK3-NEXT:    [[TMP16:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9
423 // CHECK3-NEXT:    store i64 0, ptr [[TMP16]], align 8
424 // CHECK3-NEXT:    [[TMP17:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10
425 // CHECK3-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP17]], align 4
426 // CHECK3-NEXT:    [[TMP18:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11
427 // CHECK3-NEXT:    store [3 x i32] [i32 1, i32 0, i32 0], ptr [[TMP18]], align 4
428 // CHECK3-NEXT:    [[TMP19:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12
429 // CHECK3-NEXT:    store i32 0, ptr [[TMP19]], align 4
430 // CHECK3-NEXT:    [[TMP20:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB2:[0-9]+]], i64 -1, i32 0, i32 1, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l64.region_id, ptr [[KERNEL_ARGS]])
431 // CHECK3-NEXT:    [[TMP21:%.*]] = icmp ne i32 [[TMP20]], 0
432 // CHECK3-NEXT:    br i1 [[TMP21]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]
433 // CHECK3:       omp_offload.failed:
434 // CHECK3-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l64(i32 [[TMP1]]) #[[ATTR2:[0-9]+]]
435 // CHECK3-NEXT:    br label [[OMP_OFFLOAD_CONT]]
436 // CHECK3:       omp_offload.cont:
437 // CHECK3-NEXT:    [[CALL:%.*]] = call noundef i32 @_Z5tmainIiET_v()
438 // CHECK3-NEXT:    ret i32 [[CALL]]
439 //
440 //
441 // CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l64
442 // CHECK3-SAME: (i32 noundef [[SIVAR:%.*]]) #[[ATTR1:[0-9]+]] {
443 // CHECK3-NEXT:  entry:
444 // CHECK3-NEXT:    [[SIVAR_ADDR:%.*]] = alloca i32, align 4
445 // CHECK3-NEXT:    store i32 [[SIVAR]], ptr [[SIVAR_ADDR]], align 4
446 // CHECK3-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB2]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l64.omp_outlined, ptr [[SIVAR_ADDR]])
447 // CHECK3-NEXT:    ret void
448 //
449 //
450 // CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l64.omp_outlined
451 // CHECK3-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[SIVAR:%.*]]) #[[ATTR1]] {
452 // CHECK3-NEXT:  entry:
453 // CHECK3-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 4
454 // CHECK3-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 4
455 // CHECK3-NEXT:    [[SIVAR_ADDR:%.*]] = alloca ptr, align 4
456 // CHECK3-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 4
457 // CHECK3-NEXT:    [[TMP:%.*]] = alloca i32, align 4
458 // CHECK3-NEXT:    [[DOTOMP_LB:%.*]] = alloca i32, align 4
459 // CHECK3-NEXT:    [[DOTOMP_UB:%.*]] = alloca i32, align 4
460 // CHECK3-NEXT:    [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4
461 // CHECK3-NEXT:    [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4
462 // CHECK3-NEXT:    [[SIVAR1:%.*]] = alloca i32, align 4
463 // CHECK3-NEXT:    [[I:%.*]] = alloca i32, align 4
464 // CHECK3-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 4
465 // CHECK3-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 4
466 // CHECK3-NEXT:    store ptr [[SIVAR]], ptr [[SIVAR_ADDR]], align 4
467 // CHECK3-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[SIVAR_ADDR]], align 4
468 // CHECK3-NEXT:    store i32 0, ptr [[DOTOMP_LB]], align 4
469 // CHECK3-NEXT:    store i32 1, ptr [[DOTOMP_UB]], align 4
470 // CHECK3-NEXT:    store i32 1, ptr [[DOTOMP_STRIDE]], align 4
471 // CHECK3-NEXT:    store i32 0, ptr [[DOTOMP_IS_LAST]], align 4
472 // CHECK3-NEXT:    store i32 0, ptr [[SIVAR1]], align 4
473 // CHECK3-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 4
474 // CHECK3-NEXT:    [[TMP2:%.*]] = load i32, ptr [[TMP1]], align 4
475 // CHECK3-NEXT:    call void @__kmpc_for_static_init_4(ptr @[[GLOB1:[0-9]+]], i32 [[TMP2]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)
476 // CHECK3-NEXT:    [[TMP3:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4
477 // CHECK3-NEXT:    [[CMP:%.*]] = icmp sgt i32 [[TMP3]], 1
478 // CHECK3-NEXT:    br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]
479 // CHECK3:       cond.true:
480 // CHECK3-NEXT:    br label [[COND_END:%.*]]
481 // CHECK3:       cond.false:
482 // CHECK3-NEXT:    [[TMP4:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4
483 // CHECK3-NEXT:    br label [[COND_END]]
484 // CHECK3:       cond.end:
485 // CHECK3-NEXT:    [[COND:%.*]] = phi i32 [ 1, [[COND_TRUE]] ], [ [[TMP4]], [[COND_FALSE]] ]
486 // CHECK3-NEXT:    store i32 [[COND]], ptr [[DOTOMP_UB]], align 4
487 // CHECK3-NEXT:    [[TMP5:%.*]] = load i32, ptr [[DOTOMP_LB]], align 4
488 // CHECK3-NEXT:    store i32 [[TMP5]], ptr [[DOTOMP_IV]], align 4
489 // CHECK3-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]
490 // CHECK3:       omp.inner.for.cond:
491 // CHECK3-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP6:![0-9]+]]
492 // CHECK3-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4, !llvm.access.group [[ACC_GRP6]]
493 // CHECK3-NEXT:    [[CMP2:%.*]] = icmp sle i32 [[TMP6]], [[TMP7]]
494 // CHECK3-NEXT:    br i1 [[CMP2]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]
495 // CHECK3:       omp.inner.for.body:
496 // CHECK3-NEXT:    [[TMP8:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP6]]
497 // CHECK3-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP8]], 1
498 // CHECK3-NEXT:    [[ADD:%.*]] = add nsw i32 0, [[MUL]]
499 // CHECK3-NEXT:    store i32 [[ADD]], ptr [[I]], align 4, !llvm.access.group [[ACC_GRP6]]
500 // CHECK3-NEXT:    [[TMP9:%.*]] = load i32, ptr [[I]], align 4, !llvm.access.group [[ACC_GRP6]]
501 // CHECK3-NEXT:    [[TMP10:%.*]] = load i32, ptr [[SIVAR1]], align 4, !llvm.access.group [[ACC_GRP6]]
502 // CHECK3-NEXT:    [[ADD3:%.*]] = add nsw i32 [[TMP10]], [[TMP9]]
503 // CHECK3-NEXT:    store i32 [[ADD3]], ptr [[SIVAR1]], align 4, !llvm.access.group [[ACC_GRP6]]
504 // CHECK3-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]
505 // CHECK3:       omp.body.continue:
506 // CHECK3-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]
507 // CHECK3:       omp.inner.for.inc:
508 // CHECK3-NEXT:    [[TMP11:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP6]]
509 // CHECK3-NEXT:    [[ADD4:%.*]] = add nsw i32 [[TMP11]], 1
510 // CHECK3-NEXT:    store i32 [[ADD4]], ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP6]]
511 // CHECK3-NEXT:    br label [[OMP_INNER_FOR_COND]], !llvm.loop [[LOOP7:![0-9]+]]
512 // CHECK3:       omp.inner.for.end:
513 // CHECK3-NEXT:    br label [[OMP_LOOP_EXIT:%.*]]
514 // CHECK3:       omp.loop.exit:
515 // CHECK3-NEXT:    call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP2]])
516 // CHECK3-NEXT:    [[TMP12:%.*]] = load i32, ptr [[DOTOMP_IS_LAST]], align 4
517 // CHECK3-NEXT:    [[TMP13:%.*]] = icmp ne i32 [[TMP12]], 0
518 // CHECK3-NEXT:    br i1 [[TMP13]], label [[DOTOMP_FINAL_THEN:%.*]], label [[DOTOMP_FINAL_DONE:%.*]]
519 // CHECK3:       .omp.final.then:
520 // CHECK3-NEXT:    store i32 2, ptr [[I]], align 4
521 // CHECK3-NEXT:    br label [[DOTOMP_FINAL_DONE]]
522 // CHECK3:       .omp.final.done:
523 // CHECK3-NEXT:    [[TMP14:%.*]] = load i32, ptr [[TMP0]], align 4
524 // CHECK3-NEXT:    [[TMP15:%.*]] = load i32, ptr [[SIVAR1]], align 4
525 // CHECK3-NEXT:    [[ADD5:%.*]] = add nsw i32 [[TMP14]], [[TMP15]]
526 // CHECK3-NEXT:    store i32 [[ADD5]], ptr [[TMP0]], align 4
527 // CHECK3-NEXT:    ret void
528 //
529 //
530 // CHECK3-LABEL: define {{[^@]+}}@_Z5tmainIiET_v
531 // CHECK3-SAME: () #[[ATTR3:[0-9]+]] comdat {
532 // CHECK3-NEXT:  entry:
533 // CHECK3-NEXT:    [[T_VAR:%.*]] = alloca i32, align 4
534 // CHECK3-NEXT:    [[VEC:%.*]] = alloca [2 x i32], align 4
535 // CHECK3-NEXT:    [[T_VAR_CASTED:%.*]] = alloca i32, align 4
536 // CHECK3-NEXT:    [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x ptr], align 4
537 // CHECK3-NEXT:    [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x ptr], align 4
538 // CHECK3-NEXT:    [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [1 x ptr], align 4
539 // CHECK3-NEXT:    [[TMP:%.*]] = alloca i32, align 4
540 // CHECK3-NEXT:    [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8
541 // CHECK3-NEXT:    store i32 0, ptr [[T_VAR]], align 4
542 // CHECK3-NEXT:    call void @llvm.memcpy.p0.p0.i32(ptr align 4 [[VEC]], ptr align 4 @__const._Z5tmainIiET_v.vec, i32 8, i1 false)
543 // CHECK3-NEXT:    [[TMP0:%.*]] = load i32, ptr [[T_VAR]], align 4
544 // CHECK3-NEXT:    store i32 [[TMP0]], ptr [[T_VAR_CASTED]], align 4
545 // CHECK3-NEXT:    [[TMP1:%.*]] = load i32, ptr [[T_VAR_CASTED]], align 4
546 // CHECK3-NEXT:    [[TMP2:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
547 // CHECK3-NEXT:    store i32 [[TMP1]], ptr [[TMP2]], align 4
548 // CHECK3-NEXT:    [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0
549 // CHECK3-NEXT:    store i32 [[TMP1]], ptr [[TMP3]], align 4
550 // CHECK3-NEXT:    [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 0
551 // CHECK3-NEXT:    store ptr null, ptr [[TMP4]], align 4
552 // CHECK3-NEXT:    [[TMP5:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
553 // CHECK3-NEXT:    [[TMP6:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0
554 // CHECK3-NEXT:    [[TMP7:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0
555 // CHECK3-NEXT:    store i32 3, ptr [[TMP7]], align 4
556 // CHECK3-NEXT:    [[TMP8:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1
557 // CHECK3-NEXT:    store i32 1, ptr [[TMP8]], align 4
558 // CHECK3-NEXT:    [[TMP9:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2
559 // CHECK3-NEXT:    store ptr [[TMP5]], ptr [[TMP9]], align 4
560 // CHECK3-NEXT:    [[TMP10:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3
561 // CHECK3-NEXT:    store ptr [[TMP6]], ptr [[TMP10]], align 4
562 // CHECK3-NEXT:    [[TMP11:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4
563 // CHECK3-NEXT:    store ptr @.offload_sizes.1, ptr [[TMP11]], align 4
564 // CHECK3-NEXT:    [[TMP12:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5
565 // CHECK3-NEXT:    store ptr @.offload_maptypes.2, ptr [[TMP12]], align 4
566 // CHECK3-NEXT:    [[TMP13:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6
567 // CHECK3-NEXT:    store ptr null, ptr [[TMP13]], align 4
568 // CHECK3-NEXT:    [[TMP14:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7
569 // CHECK3-NEXT:    store ptr null, ptr [[TMP14]], align 4
570 // CHECK3-NEXT:    [[TMP15:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8
571 // CHECK3-NEXT:    store i64 2, ptr [[TMP15]], align 8
572 // CHECK3-NEXT:    [[TMP16:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9
573 // CHECK3-NEXT:    store i64 0, ptr [[TMP16]], align 8
574 // CHECK3-NEXT:    [[TMP17:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10
575 // CHECK3-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP17]], align 4
576 // CHECK3-NEXT:    [[TMP18:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11
577 // CHECK3-NEXT:    store [3 x i32] [i32 1, i32 0, i32 0], ptr [[TMP18]], align 4
578 // CHECK3-NEXT:    [[TMP19:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12
579 // CHECK3-NEXT:    store i32 0, ptr [[TMP19]], align 4
580 // CHECK3-NEXT:    [[TMP20:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB2]], i64 -1, i32 0, i32 1, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_v_l32.region_id, ptr [[KERNEL_ARGS]])
581 // CHECK3-NEXT:    [[TMP21:%.*]] = icmp ne i32 [[TMP20]], 0
582 // CHECK3-NEXT:    br i1 [[TMP21]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]
583 // CHECK3:       omp_offload.failed:
584 // CHECK3-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_v_l32(i32 [[TMP1]]) #[[ATTR2]]
585 // CHECK3-NEXT:    br label [[OMP_OFFLOAD_CONT]]
586 // CHECK3:       omp_offload.cont:
587 // CHECK3-NEXT:    ret i32 0
588 //
589 //
590 // CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_v_l32
591 // CHECK3-SAME: (i32 noundef [[T_VAR:%.*]]) #[[ATTR1]] {
592 // CHECK3-NEXT:  entry:
593 // CHECK3-NEXT:    [[T_VAR_ADDR:%.*]] = alloca i32, align 4
594 // CHECK3-NEXT:    store i32 [[T_VAR]], ptr [[T_VAR_ADDR]], align 4
595 // CHECK3-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB2]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_v_l32.omp_outlined, ptr [[T_VAR_ADDR]])
596 // CHECK3-NEXT:    ret void
597 //
598 //
599 // CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_v_l32.omp_outlined
600 // CHECK3-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[T_VAR:%.*]]) #[[ATTR1]] {
601 // CHECK3-NEXT:  entry:
602 // CHECK3-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 4
603 // CHECK3-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 4
604 // CHECK3-NEXT:    [[T_VAR_ADDR:%.*]] = alloca ptr, align 4
605 // CHECK3-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 4
606 // CHECK3-NEXT:    [[TMP:%.*]] = alloca i32, align 4
607 // CHECK3-NEXT:    [[DOTOMP_LB:%.*]] = alloca i32, align 4
608 // CHECK3-NEXT:    [[DOTOMP_UB:%.*]] = alloca i32, align 4
609 // CHECK3-NEXT:    [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4
610 // CHECK3-NEXT:    [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4
611 // CHECK3-NEXT:    [[T_VAR1:%.*]] = alloca i32, align 4
612 // CHECK3-NEXT:    [[I:%.*]] = alloca i32, align 4
613 // CHECK3-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 4
614 // CHECK3-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 4
615 // CHECK3-NEXT:    store ptr [[T_VAR]], ptr [[T_VAR_ADDR]], align 4
616 // CHECK3-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[T_VAR_ADDR]], align 4
617 // CHECK3-NEXT:    store i32 0, ptr [[DOTOMP_LB]], align 4
618 // CHECK3-NEXT:    store i32 1, ptr [[DOTOMP_UB]], align 4
619 // CHECK3-NEXT:    store i32 1, ptr [[DOTOMP_STRIDE]], align 4
620 // CHECK3-NEXT:    store i32 0, ptr [[DOTOMP_IS_LAST]], align 4
621 // CHECK3-NEXT:    store i32 0, ptr [[T_VAR1]], align 4
622 // CHECK3-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 4
623 // CHECK3-NEXT:    [[TMP2:%.*]] = load i32, ptr [[TMP1]], align 4
624 // CHECK3-NEXT:    call void @__kmpc_for_static_init_4(ptr @[[GLOB1]], i32 [[TMP2]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)
625 // CHECK3-NEXT:    [[TMP3:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4
626 // CHECK3-NEXT:    [[CMP:%.*]] = icmp sgt i32 [[TMP3]], 1
627 // CHECK3-NEXT:    br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]
628 // CHECK3:       cond.true:
629 // CHECK3-NEXT:    br label [[COND_END:%.*]]
630 // CHECK3:       cond.false:
631 // CHECK3-NEXT:    [[TMP4:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4
632 // CHECK3-NEXT:    br label [[COND_END]]
633 // CHECK3:       cond.end:
634 // CHECK3-NEXT:    [[COND:%.*]] = phi i32 [ 1, [[COND_TRUE]] ], [ [[TMP4]], [[COND_FALSE]] ]
635 // CHECK3-NEXT:    store i32 [[COND]], ptr [[DOTOMP_UB]], align 4
636 // CHECK3-NEXT:    [[TMP5:%.*]] = load i32, ptr [[DOTOMP_LB]], align 4
637 // CHECK3-NEXT:    store i32 [[TMP5]], ptr [[DOTOMP_IV]], align 4
638 // CHECK3-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]
639 // CHECK3:       omp.inner.for.cond:
640 // CHECK3-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP12:![0-9]+]]
641 // CHECK3-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4, !llvm.access.group [[ACC_GRP12]]
642 // CHECK3-NEXT:    [[CMP2:%.*]] = icmp sle i32 [[TMP6]], [[TMP7]]
643 // CHECK3-NEXT:    br i1 [[CMP2]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]
644 // CHECK3:       omp.inner.for.body:
645 // CHECK3-NEXT:    [[TMP8:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP12]]
646 // CHECK3-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP8]], 1
647 // CHECK3-NEXT:    [[ADD:%.*]] = add nsw i32 0, [[MUL]]
648 // CHECK3-NEXT:    store i32 [[ADD]], ptr [[I]], align 4, !llvm.access.group [[ACC_GRP12]]
649 // CHECK3-NEXT:    [[TMP9:%.*]] = load i32, ptr [[I]], align 4, !llvm.access.group [[ACC_GRP12]]
650 // CHECK3-NEXT:    [[TMP10:%.*]] = load i32, ptr [[T_VAR1]], align 4, !llvm.access.group [[ACC_GRP12]]
651 // CHECK3-NEXT:    [[ADD3:%.*]] = add nsw i32 [[TMP10]], [[TMP9]]
652 // CHECK3-NEXT:    store i32 [[ADD3]], ptr [[T_VAR1]], align 4, !llvm.access.group [[ACC_GRP12]]
653 // CHECK3-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]
654 // CHECK3:       omp.body.continue:
655 // CHECK3-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]
656 // CHECK3:       omp.inner.for.inc:
657 // CHECK3-NEXT:    [[TMP11:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP12]]
658 // CHECK3-NEXT:    [[ADD4:%.*]] = add nsw i32 [[TMP11]], 1
659 // CHECK3-NEXT:    store i32 [[ADD4]], ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP12]]
660 // CHECK3-NEXT:    br label [[OMP_INNER_FOR_COND]], !llvm.loop [[LOOP13:![0-9]+]]
661 // CHECK3:       omp.inner.for.end:
662 // CHECK3-NEXT:    br label [[OMP_LOOP_EXIT:%.*]]
663 // CHECK3:       omp.loop.exit:
664 // CHECK3-NEXT:    call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP2]])
665 // CHECK3-NEXT:    [[TMP12:%.*]] = load i32, ptr [[DOTOMP_IS_LAST]], align 4
666 // CHECK3-NEXT:    [[TMP13:%.*]] = icmp ne i32 [[TMP12]], 0
667 // CHECK3-NEXT:    br i1 [[TMP13]], label [[DOTOMP_FINAL_THEN:%.*]], label [[DOTOMP_FINAL_DONE:%.*]]
668 // CHECK3:       .omp.final.then:
669 // CHECK3-NEXT:    store i32 2, ptr [[I]], align 4
670 // CHECK3-NEXT:    br label [[DOTOMP_FINAL_DONE]]
671 // CHECK3:       .omp.final.done:
672 // CHECK3-NEXT:    [[TMP14:%.*]] = load i32, ptr [[TMP0]], align 4
673 // CHECK3-NEXT:    [[TMP15:%.*]] = load i32, ptr [[T_VAR1]], align 4
674 // CHECK3-NEXT:    [[ADD5:%.*]] = add nsw i32 [[TMP14]], [[TMP15]]
675 // CHECK3-NEXT:    store i32 [[ADD5]], ptr [[TMP0]], align 4
676 // CHECK3-NEXT:    ret void
677 //
678 //
679 // CHECK5-LABEL: define {{[^@]+}}@main
680 // CHECK5-SAME: () #[[ATTR0:[0-9]+]] {
681 // CHECK5-NEXT:  entry:
682 // CHECK5-NEXT:    [[RETVAL:%.*]] = alloca i32, align 4
683 // CHECK5-NEXT:    [[TMP:%.*]] = alloca i32, align 4
684 // CHECK5-NEXT:    [[DOTOMP_LB:%.*]] = alloca i32, align 4
685 // CHECK5-NEXT:    [[DOTOMP_UB:%.*]] = alloca i32, align 4
686 // CHECK5-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 4
687 // CHECK5-NEXT:    [[I:%.*]] = alloca i32, align 4
688 // CHECK5-NEXT:    [[SIVAR:%.*]] = alloca i32, align 4
689 // CHECK5-NEXT:    store i32 0, ptr [[RETVAL]], align 4
690 // CHECK5-NEXT:    store i32 0, ptr [[DOTOMP_LB]], align 4
691 // CHECK5-NEXT:    store i32 1, ptr [[DOTOMP_UB]], align 4
692 // CHECK5-NEXT:    [[TMP0:%.*]] = load i32, ptr [[DOTOMP_LB]], align 4
693 // CHECK5-NEXT:    store i32 [[TMP0]], ptr [[DOTOMP_IV]], align 4
694 // CHECK5-NEXT:    store i32 0, ptr [[SIVAR]], align 4
695 // CHECK5-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]
696 // CHECK5:       omp.inner.for.cond:
697 // CHECK5-NEXT:    [[TMP1:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP2:![0-9]+]]
698 // CHECK5-NEXT:    [[TMP2:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4, !llvm.access.group [[ACC_GRP2]]
699 // CHECK5-NEXT:    [[CMP:%.*]] = icmp sle i32 [[TMP1]], [[TMP2]]
700 // CHECK5-NEXT:    br i1 [[CMP]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]
701 // CHECK5:       omp.inner.for.body:
702 // CHECK5-NEXT:    [[TMP3:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP2]]
703 // CHECK5-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP3]], 1
704 // CHECK5-NEXT:    [[ADD:%.*]] = add nsw i32 0, [[MUL]]
705 // CHECK5-NEXT:    store i32 [[ADD]], ptr [[I]], align 4, !llvm.access.group [[ACC_GRP2]]
706 // CHECK5-NEXT:    [[TMP4:%.*]] = load i32, ptr [[I]], align 4, !llvm.access.group [[ACC_GRP2]]
707 // CHECK5-NEXT:    [[TMP5:%.*]] = load i32, ptr [[SIVAR]], align 4, !llvm.access.group [[ACC_GRP2]]
708 // CHECK5-NEXT:    [[ADD1:%.*]] = add nsw i32 [[TMP5]], [[TMP4]]
709 // CHECK5-NEXT:    store i32 [[ADD1]], ptr [[SIVAR]], align 4, !llvm.access.group [[ACC_GRP2]]
710 // CHECK5-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]
711 // CHECK5:       omp.body.continue:
712 // CHECK5-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]
713 // CHECK5:       omp.inner.for.inc:
714 // CHECK5-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP2]]
715 // CHECK5-NEXT:    [[ADD2:%.*]] = add nsw i32 [[TMP6]], 1
716 // CHECK5-NEXT:    store i32 [[ADD2]], ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP2]]
717 // CHECK5-NEXT:    br label [[OMP_INNER_FOR_COND]], !llvm.loop [[LOOP3:![0-9]+]]
718 // CHECK5:       omp.inner.for.end:
719 // CHECK5-NEXT:    store i32 2, ptr [[I]], align 4
720 // CHECK5-NEXT:    [[TMP7:%.*]] = load i32, ptr @_ZZ4mainE5sivar, align 4
721 // CHECK5-NEXT:    [[TMP8:%.*]] = load i32, ptr [[SIVAR]], align 4
722 // CHECK5-NEXT:    [[ADD3:%.*]] = add nsw i32 [[TMP7]], [[TMP8]]
723 // CHECK5-NEXT:    store i32 [[ADD3]], ptr @_ZZ4mainE5sivar, align 4
724 // CHECK5-NEXT:    [[CALL:%.*]] = call noundef signext i32 @_Z5tmainIiET_v()
725 // CHECK5-NEXT:    ret i32 [[CALL]]
726 //
727 //
728 // CHECK5-LABEL: define {{[^@]+}}@_Z5tmainIiET_v
729 // CHECK5-SAME: () #[[ATTR1:[0-9]+]] comdat {
730 // CHECK5-NEXT:  entry:
731 // CHECK5-NEXT:    [[T_VAR:%.*]] = alloca i32, align 4
732 // CHECK5-NEXT:    [[VEC:%.*]] = alloca [2 x i32], align 4
733 // CHECK5-NEXT:    [[TMP:%.*]] = alloca i32, align 4
734 // CHECK5-NEXT:    [[DOTOMP_LB:%.*]] = alloca i32, align 4
735 // CHECK5-NEXT:    [[DOTOMP_UB:%.*]] = alloca i32, align 4
736 // CHECK5-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 4
737 // CHECK5-NEXT:    [[I:%.*]] = alloca i32, align 4
738 // CHECK5-NEXT:    [[T_VAR1:%.*]] = alloca i32, align 4
739 // CHECK5-NEXT:    store i32 0, ptr [[T_VAR]], align 4
740 // CHECK5-NEXT:    call void @llvm.memcpy.p0.p0.i64(ptr align 4 [[VEC]], ptr align 4 @__const._Z5tmainIiET_v.vec, i64 8, i1 false)
741 // CHECK5-NEXT:    store i32 0, ptr [[DOTOMP_LB]], align 4
742 // CHECK5-NEXT:    store i32 1, ptr [[DOTOMP_UB]], align 4
743 // CHECK5-NEXT:    [[TMP0:%.*]] = load i32, ptr [[DOTOMP_LB]], align 4
744 // CHECK5-NEXT:    store i32 [[TMP0]], ptr [[DOTOMP_IV]], align 4
745 // CHECK5-NEXT:    store i32 0, ptr [[T_VAR1]], align 4
746 // CHECK5-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]
747 // CHECK5:       omp.inner.for.cond:
748 // CHECK5-NEXT:    [[TMP1:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP6:![0-9]+]]
749 // CHECK5-NEXT:    [[TMP2:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4, !llvm.access.group [[ACC_GRP6]]
750 // CHECK5-NEXT:    [[CMP:%.*]] = icmp sle i32 [[TMP1]], [[TMP2]]
751 // CHECK5-NEXT:    br i1 [[CMP]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]
752 // CHECK5:       omp.inner.for.body:
753 // CHECK5-NEXT:    [[TMP3:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP6]]
754 // CHECK5-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP3]], 1
755 // CHECK5-NEXT:    [[ADD:%.*]] = add nsw i32 0, [[MUL]]
756 // CHECK5-NEXT:    store i32 [[ADD]], ptr [[I]], align 4, !llvm.access.group [[ACC_GRP6]]
757 // CHECK5-NEXT:    [[TMP4:%.*]] = load i32, ptr [[I]], align 4, !llvm.access.group [[ACC_GRP6]]
758 // CHECK5-NEXT:    [[TMP5:%.*]] = load i32, ptr [[T_VAR1]], align 4, !llvm.access.group [[ACC_GRP6]]
759 // CHECK5-NEXT:    [[ADD2:%.*]] = add nsw i32 [[TMP5]], [[TMP4]]
760 // CHECK5-NEXT:    store i32 [[ADD2]], ptr [[T_VAR1]], align 4, !llvm.access.group [[ACC_GRP6]]
761 // CHECK5-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]
762 // CHECK5:       omp.body.continue:
763 // CHECK5-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]
764 // CHECK5:       omp.inner.for.inc:
765 // CHECK5-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP6]]
766 // CHECK5-NEXT:    [[ADD3:%.*]] = add nsw i32 [[TMP6]], 1
767 // CHECK5-NEXT:    store i32 [[ADD3]], ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP6]]
768 // CHECK5-NEXT:    br label [[OMP_INNER_FOR_COND]], !llvm.loop [[LOOP7:![0-9]+]]
769 // CHECK5:       omp.inner.for.end:
770 // CHECK5-NEXT:    store i32 2, ptr [[I]], align 4
771 // CHECK5-NEXT:    [[TMP7:%.*]] = load i32, ptr [[T_VAR]], align 4
772 // CHECK5-NEXT:    [[TMP8:%.*]] = load i32, ptr [[T_VAR1]], align 4
773 // CHECK5-NEXT:    [[ADD4:%.*]] = add nsw i32 [[TMP7]], [[TMP8]]
774 // CHECK5-NEXT:    store i32 [[ADD4]], ptr [[T_VAR]], align 4
775 // CHECK5-NEXT:    ret i32 0
776 //
777 //
778 // CHECK7-LABEL: define {{[^@]+}}@main
779 // CHECK7-SAME: () #[[ATTR0:[0-9]+]] {
780 // CHECK7-NEXT:  entry:
781 // CHECK7-NEXT:    [[RETVAL:%.*]] = alloca i32, align 4
782 // CHECK7-NEXT:    [[TMP:%.*]] = alloca i32, align 4
783 // CHECK7-NEXT:    [[DOTOMP_LB:%.*]] = alloca i32, align 4
784 // CHECK7-NEXT:    [[DOTOMP_UB:%.*]] = alloca i32, align 4
785 // CHECK7-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 4
786 // CHECK7-NEXT:    [[I:%.*]] = alloca i32, align 4
787 // CHECK7-NEXT:    [[SIVAR:%.*]] = alloca i32, align 4
788 // CHECK7-NEXT:    store i32 0, ptr [[RETVAL]], align 4
789 // CHECK7-NEXT:    store i32 0, ptr [[DOTOMP_LB]], align 4
790 // CHECK7-NEXT:    store i32 1, ptr [[DOTOMP_UB]], align 4
791 // CHECK7-NEXT:    [[TMP0:%.*]] = load i32, ptr [[DOTOMP_LB]], align 4
792 // CHECK7-NEXT:    store i32 [[TMP0]], ptr [[DOTOMP_IV]], align 4
793 // CHECK7-NEXT:    store i32 0, ptr [[SIVAR]], align 4
794 // CHECK7-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]
795 // CHECK7:       omp.inner.for.cond:
796 // CHECK7-NEXT:    [[TMP1:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP3:![0-9]+]]
797 // CHECK7-NEXT:    [[TMP2:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4, !llvm.access.group [[ACC_GRP3]]
798 // CHECK7-NEXT:    [[CMP:%.*]] = icmp sle i32 [[TMP1]], [[TMP2]]
799 // CHECK7-NEXT:    br i1 [[CMP]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]
800 // CHECK7:       omp.inner.for.body:
801 // CHECK7-NEXT:    [[TMP3:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP3]]
802 // CHECK7-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP3]], 1
803 // CHECK7-NEXT:    [[ADD:%.*]] = add nsw i32 0, [[MUL]]
804 // CHECK7-NEXT:    store i32 [[ADD]], ptr [[I]], align 4, !llvm.access.group [[ACC_GRP3]]
805 // CHECK7-NEXT:    [[TMP4:%.*]] = load i32, ptr [[I]], align 4, !llvm.access.group [[ACC_GRP3]]
806 // CHECK7-NEXT:    [[TMP5:%.*]] = load i32, ptr [[SIVAR]], align 4, !llvm.access.group [[ACC_GRP3]]
807 // CHECK7-NEXT:    [[ADD1:%.*]] = add nsw i32 [[TMP5]], [[TMP4]]
808 // CHECK7-NEXT:    store i32 [[ADD1]], ptr [[SIVAR]], align 4, !llvm.access.group [[ACC_GRP3]]
809 // CHECK7-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]
810 // CHECK7:       omp.body.continue:
811 // CHECK7-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]
812 // CHECK7:       omp.inner.for.inc:
813 // CHECK7-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP3]]
814 // CHECK7-NEXT:    [[ADD2:%.*]] = add nsw i32 [[TMP6]], 1
815 // CHECK7-NEXT:    store i32 [[ADD2]], ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP3]]
816 // CHECK7-NEXT:    br label [[OMP_INNER_FOR_COND]], !llvm.loop [[LOOP4:![0-9]+]]
817 // CHECK7:       omp.inner.for.end:
818 // CHECK7-NEXT:    store i32 2, ptr [[I]], align 4
819 // CHECK7-NEXT:    [[TMP7:%.*]] = load i32, ptr @_ZZ4mainE5sivar, align 4
820 // CHECK7-NEXT:    [[TMP8:%.*]] = load i32, ptr [[SIVAR]], align 4
821 // CHECK7-NEXT:    [[ADD3:%.*]] = add nsw i32 [[TMP7]], [[TMP8]]
822 // CHECK7-NEXT:    store i32 [[ADD3]], ptr @_ZZ4mainE5sivar, align 4
823 // CHECK7-NEXT:    [[CALL:%.*]] = call noundef i32 @_Z5tmainIiET_v()
824 // CHECK7-NEXT:    ret i32 [[CALL]]
825 //
826 //
827 // CHECK7-LABEL: define {{[^@]+}}@_Z5tmainIiET_v
828 // CHECK7-SAME: () #[[ATTR1:[0-9]+]] comdat {
829 // CHECK7-NEXT:  entry:
830 // CHECK7-NEXT:    [[T_VAR:%.*]] = alloca i32, align 4
831 // CHECK7-NEXT:    [[VEC:%.*]] = alloca [2 x i32], align 4
832 // CHECK7-NEXT:    [[TMP:%.*]] = alloca i32, align 4
833 // CHECK7-NEXT:    [[DOTOMP_LB:%.*]] = alloca i32, align 4
834 // CHECK7-NEXT:    [[DOTOMP_UB:%.*]] = alloca i32, align 4
835 // CHECK7-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 4
836 // CHECK7-NEXT:    [[I:%.*]] = alloca i32, align 4
837 // CHECK7-NEXT:    [[T_VAR1:%.*]] = alloca i32, align 4
838 // CHECK7-NEXT:    store i32 0, ptr [[T_VAR]], align 4
839 // CHECK7-NEXT:    call void @llvm.memcpy.p0.p0.i32(ptr align 4 [[VEC]], ptr align 4 @__const._Z5tmainIiET_v.vec, i32 8, i1 false)
840 // CHECK7-NEXT:    store i32 0, ptr [[DOTOMP_LB]], align 4
841 // CHECK7-NEXT:    store i32 1, ptr [[DOTOMP_UB]], align 4
842 // CHECK7-NEXT:    [[TMP0:%.*]] = load i32, ptr [[DOTOMP_LB]], align 4
843 // CHECK7-NEXT:    store i32 [[TMP0]], ptr [[DOTOMP_IV]], align 4
844 // CHECK7-NEXT:    store i32 0, ptr [[T_VAR1]], align 4
845 // CHECK7-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]
846 // CHECK7:       omp.inner.for.cond:
847 // CHECK7-NEXT:    [[TMP1:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP7:![0-9]+]]
848 // CHECK7-NEXT:    [[TMP2:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4, !llvm.access.group [[ACC_GRP7]]
849 // CHECK7-NEXT:    [[CMP:%.*]] = icmp sle i32 [[TMP1]], [[TMP2]]
850 // CHECK7-NEXT:    br i1 [[CMP]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]
851 // CHECK7:       omp.inner.for.body:
852 // CHECK7-NEXT:    [[TMP3:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP7]]
853 // CHECK7-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP3]], 1
854 // CHECK7-NEXT:    [[ADD:%.*]] = add nsw i32 0, [[MUL]]
855 // CHECK7-NEXT:    store i32 [[ADD]], ptr [[I]], align 4, !llvm.access.group [[ACC_GRP7]]
856 // CHECK7-NEXT:    [[TMP4:%.*]] = load i32, ptr [[I]], align 4, !llvm.access.group [[ACC_GRP7]]
857 // CHECK7-NEXT:    [[TMP5:%.*]] = load i32, ptr [[T_VAR1]], align 4, !llvm.access.group [[ACC_GRP7]]
858 // CHECK7-NEXT:    [[ADD2:%.*]] = add nsw i32 [[TMP5]], [[TMP4]]
859 // CHECK7-NEXT:    store i32 [[ADD2]], ptr [[T_VAR1]], align 4, !llvm.access.group [[ACC_GRP7]]
860 // CHECK7-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]
861 // CHECK7:       omp.body.continue:
862 // CHECK7-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]
863 // CHECK7:       omp.inner.for.inc:
864 // CHECK7-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP7]]
865 // CHECK7-NEXT:    [[ADD3:%.*]] = add nsw i32 [[TMP6]], 1
866 // CHECK7-NEXT:    store i32 [[ADD3]], ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP7]]
867 // CHECK7-NEXT:    br label [[OMP_INNER_FOR_COND]], !llvm.loop [[LOOP8:![0-9]+]]
868 // CHECK7:       omp.inner.for.end:
869 // CHECK7-NEXT:    store i32 2, ptr [[I]], align 4
870 // CHECK7-NEXT:    [[TMP7:%.*]] = load i32, ptr [[T_VAR]], align 4
871 // CHECK7-NEXT:    [[TMP8:%.*]] = load i32, ptr [[T_VAR1]], align 4
872 // CHECK7-NEXT:    [[ADD4:%.*]] = add nsw i32 [[TMP7]], [[TMP8]]
873 // CHECK7-NEXT:    store i32 [[ADD4]], ptr [[T_VAR]], align 4
874 // CHECK7-NEXT:    ret i32 0
875 //
876 //
877 // CHECK9-LABEL: define {{[^@]+}}@main
878 // CHECK9-SAME: () #[[ATTR0:[0-9]+]] {
879 // CHECK9-NEXT:  entry:
880 // CHECK9-NEXT:    [[RETVAL:%.*]] = alloca i32, align 4
881 // CHECK9-NEXT:    [[REF_TMP:%.*]] = alloca [[CLASS_ANON:%.*]], align 1
882 // CHECK9-NEXT:    store i32 0, ptr [[RETVAL]], align 4
883 // CHECK9-NEXT:    call void @"_ZZ4mainENK3$_0clEv"(ptr noundef nonnull align 1 dereferenceable(1) [[REF_TMP]])
884 // CHECK9-NEXT:    ret i32 0
885 //
886 //
887 // CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l45
888 // CHECK9-SAME: (i64 noundef [[SIVAR:%.*]]) #[[ATTR2:[0-9]+]] {
889 // CHECK9-NEXT:  entry:
890 // CHECK9-NEXT:    [[SIVAR_ADDR:%.*]] = alloca i64, align 8
891 // CHECK9-NEXT:    store i64 [[SIVAR]], ptr [[SIVAR_ADDR]], align 8
892 // CHECK9-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB2:[0-9]+]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l45.omp_outlined, ptr [[SIVAR_ADDR]])
893 // CHECK9-NEXT:    ret void
894 //
895 //
896 // CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l45.omp_outlined
897 // CHECK9-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[SIVAR:%.*]]) #[[ATTR2]] {
898 // CHECK9-NEXT:  entry:
899 // CHECK9-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8
900 // CHECK9-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8
901 // CHECK9-NEXT:    [[SIVAR_ADDR:%.*]] = alloca ptr, align 8
902 // CHECK9-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 4
903 // CHECK9-NEXT:    [[TMP:%.*]] = alloca i32, align 4
904 // CHECK9-NEXT:    [[DOTOMP_LB:%.*]] = alloca i32, align 4
905 // CHECK9-NEXT:    [[DOTOMP_UB:%.*]] = alloca i32, align 4
906 // CHECK9-NEXT:    [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4
907 // CHECK9-NEXT:    [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4
908 // CHECK9-NEXT:    [[SIVAR1:%.*]] = alloca i32, align 4
909 // CHECK9-NEXT:    [[I:%.*]] = alloca i32, align 4
910 // CHECK9-NEXT:    [[REF_TMP:%.*]] = alloca [[CLASS_ANON_0:%.*]], align 8
911 // CHECK9-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8
912 // CHECK9-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8
913 // CHECK9-NEXT:    store ptr [[SIVAR]], ptr [[SIVAR_ADDR]], align 8
914 // CHECK9-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[SIVAR_ADDR]], align 8
915 // CHECK9-NEXT:    store i32 0, ptr [[DOTOMP_LB]], align 4
916 // CHECK9-NEXT:    store i32 1, ptr [[DOTOMP_UB]], align 4
917 // CHECK9-NEXT:    store i32 1, ptr [[DOTOMP_STRIDE]], align 4
918 // CHECK9-NEXT:    store i32 0, ptr [[DOTOMP_IS_LAST]], align 4
919 // CHECK9-NEXT:    store i32 0, ptr [[SIVAR1]], align 4
920 // CHECK9-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8
921 // CHECK9-NEXT:    [[TMP2:%.*]] = load i32, ptr [[TMP1]], align 4
922 // CHECK9-NEXT:    call void @__kmpc_for_static_init_4(ptr @[[GLOB1:[0-9]+]], i32 [[TMP2]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)
923 // CHECK9-NEXT:    [[TMP3:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4
924 // CHECK9-NEXT:    [[CMP:%.*]] = icmp sgt i32 [[TMP3]], 1
925 // CHECK9-NEXT:    br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]
926 // CHECK9:       cond.true:
927 // CHECK9-NEXT:    br label [[COND_END:%.*]]
928 // CHECK9:       cond.false:
929 // CHECK9-NEXT:    [[TMP4:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4
930 // CHECK9-NEXT:    br label [[COND_END]]
931 // CHECK9:       cond.end:
932 // CHECK9-NEXT:    [[COND:%.*]] = phi i32 [ 1, [[COND_TRUE]] ], [ [[TMP4]], [[COND_FALSE]] ]
933 // CHECK9-NEXT:    store i32 [[COND]], ptr [[DOTOMP_UB]], align 4
934 // CHECK9-NEXT:    [[TMP5:%.*]] = load i32, ptr [[DOTOMP_LB]], align 4
935 // CHECK9-NEXT:    store i32 [[TMP5]], ptr [[DOTOMP_IV]], align 4
936 // CHECK9-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]
937 // CHECK9:       omp.inner.for.cond:
938 // CHECK9-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP4:![0-9]+]]
939 // CHECK9-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4, !llvm.access.group [[ACC_GRP4]]
940 // CHECK9-NEXT:    [[CMP2:%.*]] = icmp sle i32 [[TMP6]], [[TMP7]]
941 // CHECK9-NEXT:    br i1 [[CMP2]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]
942 // CHECK9:       omp.inner.for.body:
943 // CHECK9-NEXT:    [[TMP8:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP4]]
944 // CHECK9-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP8]], 1
945 // CHECK9-NEXT:    [[ADD:%.*]] = add nsw i32 0, [[MUL]]
946 // CHECK9-NEXT:    store i32 [[ADD]], ptr [[I]], align 4, !llvm.access.group [[ACC_GRP4]]
947 // CHECK9-NEXT:    [[TMP9:%.*]] = load i32, ptr [[I]], align 4, !llvm.access.group [[ACC_GRP4]]
948 // CHECK9-NEXT:    [[TMP10:%.*]] = load i32, ptr [[SIVAR1]], align 4, !llvm.access.group [[ACC_GRP4]]
949 // CHECK9-NEXT:    [[ADD3:%.*]] = add nsw i32 [[TMP10]], [[TMP9]]
950 // CHECK9-NEXT:    store i32 [[ADD3]], ptr [[SIVAR1]], align 4, !llvm.access.group [[ACC_GRP4]]
951 // CHECK9-NEXT:    [[TMP11:%.*]] = getelementptr inbounds nuw [[CLASS_ANON_0]], ptr [[REF_TMP]], i32 0, i32 0
952 // CHECK9-NEXT:    store ptr [[SIVAR1]], ptr [[TMP11]], align 8, !llvm.access.group [[ACC_GRP4]]
953 // CHECK9-NEXT:    call void @"_ZZZ4mainENK3$_0clEvENKUlvE_clEv"(ptr noundef nonnull align 8 dereferenceable(8) [[REF_TMP]]), !llvm.access.group [[ACC_GRP4]]
954 // CHECK9-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]
955 // CHECK9:       omp.body.continue:
956 // CHECK9-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]
957 // CHECK9:       omp.inner.for.inc:
958 // CHECK9-NEXT:    [[TMP12:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP4]]
959 // CHECK9-NEXT:    [[ADD4:%.*]] = add nsw i32 [[TMP12]], 1
960 // CHECK9-NEXT:    store i32 [[ADD4]], ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP4]]
961 // CHECK9-NEXT:    br label [[OMP_INNER_FOR_COND]], !llvm.loop [[LOOP5:![0-9]+]]
962 // CHECK9:       omp.inner.for.end:
963 // CHECK9-NEXT:    br label [[OMP_LOOP_EXIT:%.*]]
964 // CHECK9:       omp.loop.exit:
965 // CHECK9-NEXT:    call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP2]])
966 // CHECK9-NEXT:    [[TMP13:%.*]] = load i32, ptr [[DOTOMP_IS_LAST]], align 4
967 // CHECK9-NEXT:    [[TMP14:%.*]] = icmp ne i32 [[TMP13]], 0
968 // CHECK9-NEXT:    br i1 [[TMP14]], label [[DOTOMP_FINAL_THEN:%.*]], label [[DOTOMP_FINAL_DONE:%.*]]
969 // CHECK9:       .omp.final.then:
970 // CHECK9-NEXT:    store i32 2, ptr [[I]], align 4
971 // CHECK9-NEXT:    br label [[DOTOMP_FINAL_DONE]]
972 // CHECK9:       .omp.final.done:
973 // CHECK9-NEXT:    [[TMP15:%.*]] = load i32, ptr [[TMP0]], align 4
974 // CHECK9-NEXT:    [[TMP16:%.*]] = load i32, ptr [[SIVAR1]], align 4
975 // CHECK9-NEXT:    [[ADD5:%.*]] = add nsw i32 [[TMP15]], [[TMP16]]
976 // CHECK9-NEXT:    store i32 [[ADD5]], ptr [[TMP0]], align 4
977 // CHECK9-NEXT:    ret void
978 //
979 //
980 // CHECK11-LABEL: define {{[^@]+}}@main
981 // CHECK11-SAME: () #[[ATTR0:[0-9]+]] {
982 // CHECK11-NEXT:  entry:
983 // CHECK11-NEXT:    [[RETVAL:%.*]] = alloca i32, align 4
984 // CHECK11-NEXT:    [[REF_TMP:%.*]] = alloca [[CLASS_ANON:%.*]], align 1
985 // CHECK11-NEXT:    store i32 0, ptr [[RETVAL]], align 4
986 // CHECK11-NEXT:    call void @"_ZZ4mainENK3$_0clEv"(ptr noundef nonnull align 1 dereferenceable(1) [[REF_TMP]])
987 // CHECK11-NEXT:    ret i32 0
988 //
989