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