1 ; NOTE: Assertions have been autogenerated by utils/update_test_checks.py UTC_ARGS: --check-globals
2 ; RUN: opt -S -passes=openmp-opt < %s | FileCheck %s
10 ; #pragma omp parallel
23 target triple = "nvptx64"
25 %struct.ident_t = type { i32, i32, i32, i32, ptr }
26 %struct.KernelEnvironmentTy = type { %struct.ConfigurationEnvironmentTy, ptr, ptr }
27 %struct.ConfigurationEnvironmentTy = type { i8, i8, i8, i32, i32, i32, i32 }
29 @0 = private unnamed_addr constant [23 x i8] c";unknown;unknown;0;0;;\00", align 1
30 @1 = private unnamed_addr constant %struct.ident_t { i32 0, i32 2, i32 0, i32 22, ptr @0 }, align 8
31 @i_shared = internal addrspace(3) global [4 x i8] undef, align 16
32 @i.i_shared = internal addrspace(3) global [4 x i8] undef, align 16
34 @__omp_offloading_10302_bd7e0_main_l13_kernel_environment = local_unnamed_addr constant %struct.KernelEnvironmentTy { %struct.ConfigurationEnvironmentTy { i8 0, i8 0, i8 2, i32 0, i32 0, i32 0, i32 0 }, ptr @1, ptr null }
35 @__omp_offloading_10302_bd7e0_main_l16_kernel_environment = local_unnamed_addr constant %struct.KernelEnvironmentTy { %struct.ConfigurationEnvironmentTy { i8 1, i8 0, i8 1, i32 0, i32 0, i32 0, i32 0 }, ptr @1, ptr null }
39 ; CHECK: @[[GLOB0:[0-9]+]] = private unnamed_addr constant [23 x i8] c"
40 ; CHECK: @[[GLOB1:[0-9]+]] = private unnamed_addr constant [[STRUCT_IDENT_T:%.*]] { i32 0, i32 2, i32 0, i32 22, ptr @[[GLOB0]] }, align 8
41 ; CHECK: @[[I_SHARED:[a-zA-Z0-9_$"\\.-]+]] = internal addrspace(3) global [4 x i8] undef, align 16
42 ; CHECK: @[[I_I_SHARED:[a-zA-Z0-9_$"\\.-]+]] = internal addrspace(3) global [4 x i8] undef, align 16
43 ; CHECK: @[[__OMP_OFFLOADING_10302_BD7E0_MAIN_L13_KERNEL_ENVIRONMENT:[a-zA-Z0-9_$"\\.-]+]] = local_unnamed_addr constant [[STRUCT_KERNELENVIRONMENTTY:%.*]] { [[STRUCT_CONFIGURATIONENVIRONMENTTY:%.*]] { i8 0, i8 0, i8 2, i32 0, i32 0, i32 0, i32 0 }, ptr @[[GLOB1]], ptr null }
44 ; CHECK: @[[__OMP_OFFLOADING_10302_BD7E0_MAIN_L16_KERNEL_ENVIRONMENT:[a-zA-Z0-9_$"\\.-]+]] = local_unnamed_addr constant [[STRUCT_KERNELENVIRONMENTTY:%.*]] { [[STRUCT_CONFIGURATIONENVIRONMENTTY:%.*]] { i8 1, i8 0, i8 1, i32 0, i32 0, i32 0, i32 0 }, ptr @[[GLOB1]], ptr null }
46 define weak_odr protected void @__omp_offloading_10302_bd7e0_main_l13(ptr %dyn, i64 noundef %i) local_unnamed_addr "kernel" {
47 ; CHECK-LABEL: @__omp_offloading_10302_bd7e0_main_l13(
49 ; CHECK-NEXT: [[CAPTURED_VARS_ADDRS_I:%.*]] = alloca [1 x ptr], align 8
50 ; CHECK-NEXT: [[TMP0:%.*]] = tail call i32 @__kmpc_target_init(ptr @__omp_offloading_10302_bd7e0_main_l13_kernel_environment, ptr [[DYN:%.*]])
51 ; CHECK-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP0]], -1
52 ; CHECK-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[COMMON_RET:%.*]]
54 ; CHECK-NEXT: ret void
55 ; CHECK: user_code.entry:
56 ; CHECK-NEXT: call void @llvm.lifetime.start.p0(i64 8, ptr nonnull [[CAPTURED_VARS_ADDRS_I]])
57 ; CHECK-NEXT: [[TMP1:%.*]] = tail call i32 @__kmpc_global_thread_num(ptr nonnull @[[GLOB1]]) #[[ATTR2:[0-9]+]]
58 ; CHECK-NEXT: [[TMP2:%.*]] = tail call i32 @__kmpc_get_hardware_thread_id_in_block() #[[ATTR2]]
59 ; CHECK-NEXT: [[TMP3:%.*]] = icmp eq i32 [[TMP2]], 0
60 ; CHECK-NEXT: br i1 [[TMP3]], label [[REGION_GUARDED_I:%.*]], label [[_Z3FOOI_INTERNALIZED_EXIT:%.*]]
61 ; CHECK: region.guarded.i:
62 ; CHECK-NEXT: [[I_ADDR_SROA_0_0_EXTRACT_TRUNC:%.*]] = trunc i64 [[I:%.*]] to i32
63 ; CHECK-NEXT: store i32 [[I_ADDR_SROA_0_0_EXTRACT_TRUNC]], ptr addrspacecast (ptr addrspace(3) @i_shared to ptr), align 16
64 ; CHECK-NEXT: br label [[_Z3FOOI_INTERNALIZED_EXIT]]
65 ; CHECK: _Z3fooi.internalized.exit:
66 ; CHECK-NEXT: tail call void @__kmpc_barrier_simple_spmd(ptr nonnull @[[GLOB1]], i32 [[TMP2]]) #[[ATTR2]]
67 ; CHECK-NEXT: store ptr addrspacecast (ptr addrspace(3) @i_shared to ptr), ptr [[CAPTURED_VARS_ADDRS_I]], align 8
68 ; CHECK-NEXT: call void @__kmpc_parallel_51(ptr nonnull @[[GLOB1]], i32 [[TMP1]], i32 1, i32 -1, i32 -1, ptr nonnull @__omp_outlined__, ptr nonnull @__omp_outlined___wrapper, ptr nonnull [[CAPTURED_VARS_ADDRS_I]], i64 1)
69 ; CHECK-NEXT: call void @llvm.lifetime.end.p0(i64 8, ptr nonnull [[CAPTURED_VARS_ADDRS_I]])
70 ; CHECK-NEXT: call void @__kmpc_target_deinit()
71 ; CHECK-NEXT: br label [[COMMON_RET]]
74 %captured_vars_addrs.i = alloca [1 x ptr], align 8
75 %0 = tail call i32 @__kmpc_target_init(ptr @__omp_offloading_10302_bd7e0_main_l13_kernel_environment, ptr %dyn) #6
76 %exec_user_code = icmp eq i32 %0, -1
77 br i1 %exec_user_code, label %user_code.entry, label %common.ret
79 common.ret: ; preds = %entry, %_Z3fooi.internalized.exit
82 user_code.entry: ; preds = %entry
83 call void @llvm.lifetime.start.p0(i64 8, ptr nonnull %captured_vars_addrs.i)
84 %1 = tail call i32 @__kmpc_global_thread_num(ptr nonnull @1) #6
85 %2 = tail call i32 @__kmpc_get_hardware_thread_id_in_block() #6
86 %3 = icmp eq i32 %2, 0
87 br i1 %3, label %region.guarded.i, label %_Z3fooi.internalized.exit
89 region.guarded.i: ; preds = %user_code.entry
90 %i.addr.sroa.0.0.extract.trunc = trunc i64 %i to i32
91 store i32 %i.addr.sroa.0.0.extract.trunc, ptr addrspacecast (ptr addrspace(3) @i_shared to ptr), align 16
92 br label %_Z3fooi.internalized.exit
94 _Z3fooi.internalized.exit: ; preds = %user_code.entry, %region.guarded.i
95 tail call void @__kmpc_barrier_simple_spmd(ptr nonnull @1, i32 %2)
96 store ptr addrspacecast (ptr addrspace(3) @i_shared to ptr), ptr %captured_vars_addrs.i, align 8
97 call void @__kmpc_parallel_51(ptr nonnull @1, i32 %1, i32 1, i32 -1, i32 -1, ptr nonnull @__omp_outlined__, ptr nonnull @__omp_outlined___wrapper, ptr nonnull %captured_vars_addrs.i, i64 1) #6
98 call void @llvm.lifetime.end.p0(i64 8, ptr nonnull %captured_vars_addrs.i)
99 call void @__kmpc_target_deinit() #6
103 declare i32 @__kmpc_target_init(ptr, ptr) local_unnamed_addr
105 define hidden void @_Z3fooi(i32 noundef %i1) local_unnamed_addr #1 {
106 ; CHECK-LABEL: @_Z3fooi(
108 ; CHECK-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [1 x ptr], align 8
109 ; CHECK-NEXT: [[TMP0:%.*]] = tail call i32 @__kmpc_global_thread_num(ptr nonnull @[[GLOB1]]) #[[ATTR2]]
110 ; CHECK-NEXT: [[I:%.*]] = tail call align 16 dereferenceable_or_null(4) ptr @__kmpc_alloc_shared(i64 4) #[[ATTR2]]
111 ; CHECK-NEXT: store i32 [[I1:%.*]], ptr [[I]], align 16
112 ; CHECK-NEXT: store ptr [[I]], ptr [[CAPTURED_VARS_ADDRS]], align 8
113 ; CHECK-NEXT: call void @__kmpc_parallel_51(ptr nonnull @[[GLOB1]], i32 [[TMP0]], i32 1, i32 -1, i32 -1, ptr nonnull @__omp_outlined__, ptr nonnull @__omp_outlined___wrapper, ptr nonnull [[CAPTURED_VARS_ADDRS]], i64 1)
114 ; CHECK-NEXT: call void @__kmpc_free_shared(ptr [[I]], i64 4) #[[ATTR2]]
115 ; CHECK-NEXT: ret void
118 %captured_vars_addrs = alloca [1 x ptr], align 8
119 %0 = tail call i32 @__kmpc_global_thread_num(ptr nonnull @1) #6
120 %i = tail call align 16 dereferenceable_or_null(4) ptr @__kmpc_alloc_shared(i64 4)
121 store i32 %i1, ptr %i, align 16
122 store ptr %i, ptr %captured_vars_addrs, align 8
123 call void @__kmpc_parallel_51(ptr nonnull @1, i32 %0, i32 1, i32 -1, i32 -1, ptr nonnull @__omp_outlined__, ptr nonnull @__omp_outlined___wrapper, ptr nonnull %captured_vars_addrs, i64 1) #6
124 call void @__kmpc_free_shared(ptr %i, i64 4)
128 declare void @__kmpc_target_deinit(ptr, i8) local_unnamed_addr
130 define weak_odr protected void @__omp_offloading_10302_bd7e0_main_l16(ptr %dyn, i64 noundef %i) local_unnamed_addr "kernel" {
131 ; CHECK-LABEL: @__omp_offloading_10302_bd7e0_main_l16(
133 ; CHECK-NEXT: [[CAPTURED_VARS_ADDRS_I:%.*]] = alloca [1 x ptr], align 8
134 ; CHECK-NEXT: [[TMP0:%.*]] = tail call i32 @__kmpc_target_init(ptr @__omp_offloading_10302_bd7e0_main_l16_kernel_environment, ptr [[DYN:%.*]])
135 ; CHECK-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP0]], -1
136 ; CHECK-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[COMMON_RET:%.*]]
138 ; CHECK-NEXT: ret void
139 ; CHECK: user_code.entry:
140 ; CHECK-NEXT: [[I_ADDR_SROA_0_0_EXTRACT_TRUNC:%.*]] = trunc i64 [[I:%.*]] to i32
141 ; CHECK-NEXT: call void @llvm.lifetime.start.p0(i64 8, ptr nonnull [[CAPTURED_VARS_ADDRS_I]])
142 ; CHECK-NEXT: [[TMP1:%.*]] = tail call i32 @__kmpc_global_thread_num(ptr nonnull @[[GLOB1]]) #[[ATTR2]]
143 ; CHECK-NEXT: store i32 [[I_ADDR_SROA_0_0_EXTRACT_TRUNC]], ptr addrspacecast (ptr addrspace(3) @i.i_shared to ptr), align 16
144 ; CHECK-NEXT: store ptr addrspacecast (ptr addrspace(3) @i.i_shared to ptr), ptr [[CAPTURED_VARS_ADDRS_I]], align 8
145 ; CHECK-NEXT: call void @__kmpc_parallel_51(ptr nonnull @[[GLOB1]], i32 [[TMP1]], i32 1, i32 -1, i32 -1, ptr nonnull @__omp_outlined__1, ptr nonnull @__omp_outlined__1_wrapper, ptr nonnull [[CAPTURED_VARS_ADDRS_I]], i64 1)
146 ; CHECK-NEXT: call void @llvm.lifetime.end.p0(i64 8, ptr nonnull [[CAPTURED_VARS_ADDRS_I]])
147 ; CHECK-NEXT: call void @__kmpc_target_deinit()
148 ; CHECK-NEXT: br label [[COMMON_RET]]
151 %captured_vars_addrs.i = alloca [1 x ptr], align 8
152 %0 = tail call i32 @__kmpc_target_init(ptr @__omp_offloading_10302_bd7e0_main_l16_kernel_environment, ptr %dyn) #6
153 %exec_user_code = icmp eq i32 %0, -1
154 br i1 %exec_user_code, label %user_code.entry, label %common.ret
156 common.ret: ; preds = %entry, %user_code.entry
159 user_code.entry: ; preds = %entry
160 %i.addr.sroa.0.0.extract.trunc = trunc i64 %i to i32
161 call void @llvm.lifetime.start.p0(i64 8, ptr nonnull %captured_vars_addrs.i)
162 %1 = tail call i32 @__kmpc_global_thread_num(ptr nonnull @1) #6
163 store i32 %i.addr.sroa.0.0.extract.trunc, ptr addrspacecast (ptr addrspace(3) @i.i_shared to ptr), align 16
164 store ptr addrspacecast (ptr addrspace(3) @i.i_shared to ptr), ptr %captured_vars_addrs.i, align 8
165 call void @__kmpc_parallel_51(ptr nonnull @1, i32 %1, i32 1, i32 -1, i32 -1, ptr nonnull @__omp_outlined__1, ptr nonnull @__omp_outlined__1_wrapper, ptr nonnull %captured_vars_addrs.i, i64 1) #6
166 call void @llvm.lifetime.end.p0(i64 8, ptr nonnull %captured_vars_addrs.i)
167 call void @__kmpc_target_deinit() #6
171 define hidden void @_Z4foo1i(i32 noundef %i1) local_unnamed_addr #1 {
172 ; CHECK-LABEL: @_Z4foo1i(
174 ; CHECK-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [1 x ptr], align 8
175 ; CHECK-NEXT: [[TMP0:%.*]] = tail call i32 @__kmpc_global_thread_num(ptr nonnull @[[GLOB1]]) #[[ATTR2]]
176 ; CHECK-NEXT: [[I:%.*]] = tail call align 16 dereferenceable_or_null(4) ptr @__kmpc_alloc_shared(i64 4) #[[ATTR2]]
177 ; CHECK-NEXT: store i32 [[I1:%.*]], ptr [[I]], align 16
178 ; CHECK-NEXT: store ptr [[I]], ptr [[CAPTURED_VARS_ADDRS]], align 8
179 ; CHECK-NEXT: call void @__kmpc_parallel_51(ptr nonnull @[[GLOB1]], i32 [[TMP0]], i32 1, i32 -1, i32 -1, ptr nonnull @__omp_outlined__1, ptr nonnull @__omp_outlined__1_wrapper, ptr nonnull [[CAPTURED_VARS_ADDRS]], i64 1)
180 ; CHECK-NEXT: call void @__kmpc_free_shared(ptr [[I]], i64 4) #[[ATTR2]]
181 ; CHECK-NEXT: ret void
184 %captured_vars_addrs = alloca [1 x ptr], align 8
185 %0 = tail call i32 @__kmpc_global_thread_num(ptr nonnull @1) #6
186 %i = tail call align 16 dereferenceable_or_null(4) ptr @__kmpc_alloc_shared(i64 4)
187 store i32 %i1, ptr %i, align 16
188 store ptr %i, ptr %captured_vars_addrs, align 8
189 call void @__kmpc_parallel_51(ptr nonnull @1, i32 %0, i32 1, i32 -1, i32 -1, ptr nonnull @__omp_outlined__1, ptr nonnull @__omp_outlined__1_wrapper, ptr nonnull %captured_vars_addrs, i64 1) #6
190 call void @__kmpc_free_shared(ptr %i, i64 4)
194 declare ptr @__kmpc_alloc_shared(i64) local_unnamed_addr #3
196 define internal void @__omp_outlined__(ptr noalias nocapture readnone %.global_tid., ptr noalias nocapture readnone %.bound_tid., ptr nocapture noundef nonnull readonly align 4 dereferenceable(4) %i) #4 {
197 ; CHECK-LABEL: @__omp_outlined__(
199 ; CHECK-NEXT: [[CAPTURED_VARS_ADDRS_I:%.*]] = alloca [1 x ptr], align 8
200 ; CHECK-NEXT: [[TMP0:%.*]] = load i32, ptr [[I:%.*]], align 4
201 ; CHECK-NEXT: call void @llvm.lifetime.start.p0(i64 8, ptr nonnull [[CAPTURED_VARS_ADDRS_I]])
202 ; CHECK-NEXT: [[TMP1:%.*]] = tail call i32 @__kmpc_global_thread_num(ptr nonnull @[[GLOB1]]) #[[ATTR2]]
203 ; CHECK-NEXT: [[I_I:%.*]] = tail call align 16 dereferenceable_or_null(4) ptr @__kmpc_alloc_shared(i64 4) #[[ATTR2]]
204 ; CHECK-NEXT: store i32 [[TMP0]], ptr [[I_I]], align 16
205 ; CHECK-NEXT: store ptr [[I_I]], ptr [[CAPTURED_VARS_ADDRS_I]], align 8
206 ; CHECK-NEXT: call void @__kmpc_parallel_51(ptr nonnull @[[GLOB1]], i32 [[TMP1]], i32 1, i32 -1, i32 -1, ptr nonnull @__omp_outlined__1, ptr nonnull @__omp_outlined__1_wrapper, ptr nonnull [[CAPTURED_VARS_ADDRS_I]], i64 1)
207 ; CHECK-NEXT: call void @__kmpc_free_shared(ptr [[I_I]], i64 4) #[[ATTR2]]
208 ; CHECK-NEXT: call void @llvm.lifetime.end.p0(i64 8, ptr nonnull [[CAPTURED_VARS_ADDRS_I]])
209 ; CHECK-NEXT: ret void
212 %captured_vars_addrs.i = alloca [1 x ptr], align 8
213 %0 = load i32, ptr %i, align 4
214 call void @llvm.lifetime.start.p0(i64 8, ptr nonnull %captured_vars_addrs.i)
215 %1 = tail call i32 @__kmpc_global_thread_num(ptr nonnull @1) #6
216 %i.i = tail call align 16 dereferenceable_or_null(4) ptr @__kmpc_alloc_shared(i64 4) #6
217 store i32 %0, ptr %i.i, align 16
218 store ptr %i.i, ptr %captured_vars_addrs.i, align 8
219 call void @__kmpc_parallel_51(ptr nonnull @1, i32 %1, i32 1, i32 -1, i32 -1, ptr nonnull @__omp_outlined__1, ptr nonnull @__omp_outlined__1_wrapper, ptr nonnull %captured_vars_addrs.i, i64 1) #6
220 call void @__kmpc_free_shared(ptr %i.i, i64 4) #6
221 call void @llvm.lifetime.end.p0(i64 8, ptr nonnull %captured_vars_addrs.i)
225 define internal void @__omp_outlined___wrapper(i16 zeroext %0, i32 %1) #5 {
226 ; CHECK-LABEL: @__omp_outlined___wrapper(
228 ; CHECK-NEXT: [[CAPTURED_VARS_ADDRS_I_I:%.*]] = alloca [1 x ptr], align 8
229 ; CHECK-NEXT: [[GLOBAL_ARGS:%.*]] = alloca ptr, align 8
230 ; CHECK-NEXT: call void @__kmpc_get_shared_variables(ptr nonnull [[GLOBAL_ARGS]])
231 ; CHECK-NEXT: [[TMP2:%.*]] = load ptr, ptr [[GLOBAL_ARGS]], align 8
232 ; CHECK-NEXT: [[TMP3:%.*]] = load ptr, ptr [[TMP2]], align 8
233 ; CHECK-NEXT: [[TMP4:%.*]] = load i32, ptr [[TMP3]], align 4
234 ; CHECK-NEXT: call void @llvm.lifetime.start.p0(i64 8, ptr nonnull [[CAPTURED_VARS_ADDRS_I_I]])
235 ; CHECK-NEXT: [[TMP5:%.*]] = call i32 @__kmpc_global_thread_num(ptr nonnull @[[GLOB1]]) #[[ATTR2]]
236 ; CHECK-NEXT: [[I_I_I:%.*]] = call align 16 dereferenceable_or_null(4) ptr @__kmpc_alloc_shared(i64 4) #[[ATTR2]]
237 ; CHECK-NEXT: store i32 [[TMP4]], ptr [[I_I_I]], align 16
238 ; CHECK-NEXT: store ptr [[I_I_I]], ptr [[CAPTURED_VARS_ADDRS_I_I]], align 8
239 ; CHECK-NEXT: call void @__kmpc_parallel_51(ptr nonnull @[[GLOB1]], i32 [[TMP5]], i32 1, i32 -1, i32 -1, ptr nonnull @__omp_outlined__1, ptr nonnull @__omp_outlined__1_wrapper, ptr nonnull [[CAPTURED_VARS_ADDRS_I_I]], i64 1)
240 ; CHECK-NEXT: call void @__kmpc_free_shared(ptr [[I_I_I]], i64 4) #[[ATTR2]]
241 ; CHECK-NEXT: call void @llvm.lifetime.end.p0(i64 8, ptr nonnull [[CAPTURED_VARS_ADDRS_I_I]])
242 ; CHECK-NEXT: ret void
245 %captured_vars_addrs.i.i = alloca [1 x ptr], align 8
246 %global_args = alloca ptr, align 8
247 call void @__kmpc_get_shared_variables(ptr nonnull %global_args) #6
248 %2 = load ptr, ptr %global_args, align 8
249 %3 = load ptr, ptr %2, align 8
250 %4 = load i32, ptr %3, align 4
251 call void @llvm.lifetime.start.p0(i64 8, ptr nonnull %captured_vars_addrs.i.i)
252 %5 = call i32 @__kmpc_global_thread_num(ptr nonnull @1) #6
253 %i.i.i = call align 16 dereferenceable_or_null(4) ptr @__kmpc_alloc_shared(i64 4) #6
254 store i32 %4, ptr %i.i.i, align 16
255 store ptr %i.i.i, ptr %captured_vars_addrs.i.i, align 8
256 call void @__kmpc_parallel_51(ptr nonnull @1, i32 %5, i32 1, i32 -1, i32 -1, ptr nonnull @__omp_outlined__1, ptr nonnull @__omp_outlined__1_wrapper, ptr nonnull %captured_vars_addrs.i.i, i64 1) #6
257 call void @__kmpc_free_shared(ptr %i.i.i, i64 4) #6
258 call void @llvm.lifetime.end.p0(i64 8, ptr nonnull %captured_vars_addrs.i.i)
262 declare void @__kmpc_get_shared_variables(ptr) local_unnamed_addr
264 declare i32 @__kmpc_global_thread_num(ptr) local_unnamed_addr #6
266 declare void @__kmpc_parallel_51(ptr, i32, i32, i32, i32, ptr, ptr, ptr, i64) local_unnamed_addr #7
268 declare void @__kmpc_free_shared(ptr allocptr nocapture, i64) local_unnamed_addr #8
270 define internal void @__omp_outlined__1(ptr noalias nocapture readnone %.global_tid., ptr noalias nocapture readnone %.bound_tid., ptr nocapture noundef nonnull align 4 dereferenceable(4) %i) #9 {
271 ; CHECK-LABEL: @__omp_outlined__1(
273 ; CHECK-NEXT: [[TMP0:%.*]] = load i32, ptr [[I:%.*]], align 4
274 ; CHECK-NEXT: [[INC:%.*]] = add nsw i32 [[TMP0]], 1
275 ; CHECK-NEXT: store i32 [[INC]], ptr [[I]], align 4
276 ; CHECK-NEXT: ret void
279 %0 = load i32, ptr %i, align 4
280 %inc = add nsw i32 %0, 1
281 store i32 %inc, ptr %i, align 4
285 define internal void @__omp_outlined__1_wrapper(i16 zeroext %0, i32 %1) #5 {
286 ; CHECK-LABEL: @__omp_outlined__1_wrapper(
288 ; CHECK-NEXT: [[GLOBAL_ARGS:%.*]] = alloca ptr, align 8
289 ; CHECK-NEXT: call void @__kmpc_get_shared_variables(ptr nonnull [[GLOBAL_ARGS]])
290 ; CHECK-NEXT: [[TMP2:%.*]] = load ptr, ptr [[GLOBAL_ARGS]], align 8
291 ; CHECK-NEXT: [[TMP3:%.*]] = load ptr, ptr [[TMP2]], align 8
292 ; CHECK-NEXT: [[TMP4:%.*]] = load i32, ptr [[TMP3]], align 4
293 ; CHECK-NEXT: [[INC_I:%.*]] = add nsw i32 [[TMP4]], 1
294 ; CHECK-NEXT: store i32 [[INC_I]], ptr [[TMP3]], align 4
295 ; CHECK-NEXT: ret void
298 %global_args = alloca ptr, align 8
299 call void @__kmpc_get_shared_variables(ptr nonnull %global_args) #6
300 %2 = load ptr, ptr %global_args, align 8
301 %3 = load ptr, ptr %2, align 8
302 %4 = load i32, ptr %3, align 4
303 %inc.i = add nsw i32 %4, 1
304 store i32 %inc.i, ptr %3, align 4
308 declare i32 @__kmpc_get_hardware_thread_id_in_block() local_unnamed_addr
310 declare void @__kmpc_barrier_simple_spmd(ptr, i32) local_unnamed_addr #10
312 declare void @llvm.lifetime.start.p0(i64 immarg, ptr nocapture) #11
314 declare void @llvm.lifetime.end.p0(i64 immarg, ptr nocapture) #11
317 !omp_offload.info = !{!0, !1}
318 !nvvm.annotations = !{!2, !3}
319 !llvm.module.flags = !{!4, !5}
321 !0 = !{i32 0, i32 66306, i32 776160, !"main", i32 13, i32 0, i32 0}
322 !1 = !{i32 0, i32 66306, i32 776160, !"main", i32 16, i32 0, i32 1}
323 !2 = !{ptr @__omp_offloading_10302_bd7e0_main_l13, !"kernel", i32 1}
324 !3 = !{ptr @__omp_offloading_10302_bd7e0_main_l16, !"kernel", i32 1}
326 !4 = !{i32 7, !"openmp", i32 50}
327 !5 = !{i32 7, !"openmp-device", i32 50}
329 ; CHECK: attributes #[[ATTR0:[0-9]+]] = { "kernel" }
330 ; CHECK: attributes #[[ATTR1:[0-9]+]] = { nosync nounwind allocsize(0) }
331 ; CHECK: attributes #[[ATTR2]] = { nounwind }
332 ; CHECK: attributes #[[ATTR3:[0-9]+]] = { alwaysinline }
333 ; CHECK: attributes #[[ATTR4:[0-9]+]] = { nosync nounwind }
334 ; CHECK: attributes #[[ATTR5:[0-9]+]] = { convergent nounwind }
335 ; CHECK: attributes #[[ATTR6:[0-9]+]] = { nocallback nofree nosync nounwind willreturn memory(argmem: readwrite) }
337 ; CHECK: [[META0:![0-9]+]] = !{i32 0, i32 66306, i32 776160, !"main", i32 13, i32 0, i32 0}
338 ; CHECK: [[META1:![0-9]+]] = !{i32 0, i32 66306, i32 776160, !"main", i32 16, i32 0, i32 1}
339 ; CHECK: [[META2:![0-9]+]] = !{ptr @__omp_offloading_10302_bd7e0_main_l13, !"kernel", i32 1}
340 ; CHECK: [[META3:![0-9]+]] = !{ptr @__omp_offloading_10302_bd7e0_main_l16, !"kernel", i32 1}
341 ; CHECK: [[META4:![0-9]+]] = !{i32 7, !"openmp", i32 50}
342 ; CHECK: [[META5:![0-9]+]] = !{i32 7, !"openmp-device", i32 50}