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 tx a = 0; 15 short aa = 0; 16 tx b[10]; 17 18 #pragma omp target teams if(0) 19 { 20 b[2] += 1; 21 } 22 23 #pragma omp target teams if(1) 24 { 25 a = '1'; 26 } 27 28 #pragma omp target teams if(n>40) 29 { 30 aa = 1; 31 } 32 33 #pragma omp target teams 34 { 35 #pragma omp parallel 36 #pragma omp parallel 37 aa = 1; 38 } 39 40 return a; 41 } 42 43 int bar(int n){ 44 int a = 0; 45 46 a += ftemplate<char>(n); 47 48 return a; 49 } 50 51 #endif 52 // CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l23_worker 53 // CHECK1-SAME: () #[[ATTR0:[0-9]+]] { 54 // CHECK1-NEXT: entry: 55 // CHECK1-NEXT: [[WORK_FN:%.*]] = alloca i8*, align 8 56 // CHECK1-NEXT: [[EXEC_STATUS:%.*]] = alloca i8, align 1 57 // CHECK1-NEXT: store i8* null, i8** [[WORK_FN]], align 8 58 // CHECK1-NEXT: store i8 0, i8* [[EXEC_STATUS]], align 1 59 // CHECK1-NEXT: br label [[DOTAWAIT_WORK:%.*]] 60 // CHECK1: .await.work: 61 // CHECK1-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0) 62 // CHECK1-NEXT: [[TMP0:%.*]] = call i1 @__kmpc_kernel_parallel(i8** [[WORK_FN]]) 63 // CHECK1-NEXT: [[TMP1:%.*]] = zext i1 [[TMP0]] to i8 64 // CHECK1-NEXT: store i8 [[TMP1]], i8* [[EXEC_STATUS]], align 1 65 // CHECK1-NEXT: [[TMP2:%.*]] = load i8*, i8** [[WORK_FN]], align 8 66 // CHECK1-NEXT: [[SHOULD_TERMINATE:%.*]] = icmp eq i8* [[TMP2]], null 67 // CHECK1-NEXT: br i1 [[SHOULD_TERMINATE]], label [[DOTEXIT:%.*]], label [[DOTSELECT_WORKERS:%.*]] 68 // CHECK1: .select.workers: 69 // CHECK1-NEXT: [[TMP3:%.*]] = load i8, i8* [[EXEC_STATUS]], align 1 70 // CHECK1-NEXT: [[IS_ACTIVE:%.*]] = icmp ne i8 [[TMP3]], 0 71 // CHECK1-NEXT: br i1 [[IS_ACTIVE]], label [[DOTEXECUTE_PARALLEL:%.*]], label [[DOTBARRIER_PARALLEL:%.*]] 72 // CHECK1: .execute.parallel: 73 // CHECK1-NEXT: [[TMP4:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1:[0-9]+]]) 74 // CHECK1-NEXT: [[TMP5:%.*]] = bitcast i8* [[TMP2]] to void (i16, i32)* 75 // CHECK1-NEXT: call void [[TMP5]](i16 0, i32 [[TMP4]]) 76 // CHECK1-NEXT: br label [[DOTTERMINATE_PARALLEL:%.*]] 77 // CHECK1: .terminate.parallel: 78 // CHECK1-NEXT: call void @__kmpc_kernel_end_parallel() 79 // CHECK1-NEXT: br label [[DOTBARRIER_PARALLEL]] 80 // CHECK1: .barrier.parallel: 81 // CHECK1-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0) 82 // CHECK1-NEXT: br label [[DOTAWAIT_WORK]] 83 // CHECK1: .exit: 84 // CHECK1-NEXT: ret void 85 // 86 // 87 // CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l23 88 // CHECK1-SAME: (i64 [[A:%.*]]) #[[ATTR1:[0-9]+]] { 89 // CHECK1-NEXT: entry: 90 // CHECK1-NEXT: [[A_ADDR:%.*]] = alloca i64, align 8 91 // CHECK1-NEXT: [[A_CASTED:%.*]] = alloca i64, align 8 92 // CHECK1-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4 93 // CHECK1-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 4 94 // CHECK1-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4 95 // CHECK1-NEXT: store i64 [[A]], i64* [[A_ADDR]], align 8 96 // CHECK1-NEXT: [[CONV:%.*]] = bitcast i64* [[A_ADDR]] to i8* 97 // CHECK1-NEXT: [[NVPTX_TID:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x() 98 // CHECK1-NEXT: [[NVPTX_NUM_THREADS:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() 99 // CHECK1-NEXT: [[NVPTX_WARP_SIZE:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize() 100 // CHECK1-NEXT: [[THREAD_LIMIT:%.*]] = sub nuw i32 [[NVPTX_NUM_THREADS]], [[NVPTX_WARP_SIZE]] 101 // CHECK1-NEXT: [[TMP0:%.*]] = icmp ult i32 [[NVPTX_TID]], [[THREAD_LIMIT]] 102 // CHECK1-NEXT: br i1 [[TMP0]], label [[DOTWORKER:%.*]], label [[DOTMASTERCHECK:%.*]] 103 // CHECK1: .worker: 104 // CHECK1-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l23_worker() #[[ATTR3:[0-9]+]] 105 // CHECK1-NEXT: br label [[DOTEXIT:%.*]] 106 // CHECK1: .mastercheck: 107 // CHECK1-NEXT: [[NVPTX_TID1:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x() 108 // CHECK1-NEXT: [[NVPTX_NUM_THREADS2:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() 109 // CHECK1-NEXT: [[NVPTX_WARP_SIZE3:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize() 110 // CHECK1-NEXT: [[TMP1:%.*]] = sub nuw i32 [[NVPTX_WARP_SIZE3]], 1 111 // CHECK1-NEXT: [[TMP2:%.*]] = sub nuw i32 [[NVPTX_NUM_THREADS2]], 1 112 // CHECK1-NEXT: [[TMP3:%.*]] = xor i32 [[TMP1]], -1 113 // CHECK1-NEXT: [[MASTER_TID:%.*]] = and i32 [[TMP2]], [[TMP3]] 114 // CHECK1-NEXT: [[TMP4:%.*]] = icmp eq i32 [[NVPTX_TID1]], [[MASTER_TID]] 115 // CHECK1-NEXT: br i1 [[TMP4]], label [[DOTMASTER:%.*]], label [[DOTEXIT]] 116 // CHECK1: .master: 117 // CHECK1-NEXT: [[NVPTX_NUM_THREADS4:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() 118 // CHECK1-NEXT: [[NVPTX_WARP_SIZE5:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize() 119 // CHECK1-NEXT: [[THREAD_LIMIT6:%.*]] = sub nuw i32 [[NVPTX_NUM_THREADS4]], [[NVPTX_WARP_SIZE5]] 120 // CHECK1-NEXT: call void @__kmpc_kernel_init(i32 [[THREAD_LIMIT6]], i16 1) 121 // CHECK1-NEXT: [[TMP5:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1]]) 122 // CHECK1-NEXT: [[TMP6:%.*]] = load i8, i8* [[CONV]], align 8 123 // CHECK1-NEXT: [[CONV7:%.*]] = bitcast i64* [[A_CASTED]] to i8* 124 // CHECK1-NEXT: store i8 [[TMP6]], i8* [[CONV7]], align 1 125 // CHECK1-NEXT: [[TMP7:%.*]] = load i64, i64* [[A_CASTED]], align 8 126 // CHECK1-NEXT: store i32 [[TMP5]], i32* [[DOTTHREADID_TEMP_]], align 4 127 // CHECK1-NEXT: call void @__omp_outlined__(i32* [[DOTTHREADID_TEMP_]], i32* [[DOTZERO_ADDR]], i64 [[TMP7]]) #[[ATTR3]] 128 // CHECK1-NEXT: br label [[DOTTERMINATION_NOTIFIER:%.*]] 129 // CHECK1: .termination.notifier: 130 // CHECK1-NEXT: call void @__kmpc_kernel_deinit(i16 1) 131 // CHECK1-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0) 132 // CHECK1-NEXT: br label [[DOTEXIT]] 133 // CHECK1: .exit: 134 // CHECK1-NEXT: ret void 135 // 136 // 137 // CHECK1-LABEL: define {{[^@]+}}@__omp_outlined__ 138 // CHECK1-SAME: (i32* noalias [[DOTGLOBAL_TID_:%.*]], i32* noalias [[DOTBOUND_TID_:%.*]], i64 [[A:%.*]]) #[[ATTR1]] { 139 // CHECK1-NEXT: entry: 140 // CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8 141 // CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8 142 // CHECK1-NEXT: [[A_ADDR:%.*]] = alloca i64, align 8 143 // CHECK1-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 8 144 // CHECK1-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 8 145 // CHECK1-NEXT: store i64 [[A]], i64* [[A_ADDR]], align 8 146 // CHECK1-NEXT: [[CONV:%.*]] = bitcast i64* [[A_ADDR]] to i8* 147 // CHECK1-NEXT: store i8 49, i8* [[CONV]], align 8 148 // CHECK1-NEXT: ret void 149 // 150 // 151 // CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l28_worker 152 // CHECK1-SAME: () #[[ATTR0]] { 153 // CHECK1-NEXT: entry: 154 // CHECK1-NEXT: [[WORK_FN:%.*]] = alloca i8*, align 8 155 // CHECK1-NEXT: [[EXEC_STATUS:%.*]] = alloca i8, align 1 156 // CHECK1-NEXT: store i8* null, i8** [[WORK_FN]], align 8 157 // CHECK1-NEXT: store i8 0, i8* [[EXEC_STATUS]], align 1 158 // CHECK1-NEXT: br label [[DOTAWAIT_WORK:%.*]] 159 // CHECK1: .await.work: 160 // CHECK1-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0) 161 // CHECK1-NEXT: [[TMP0:%.*]] = call i1 @__kmpc_kernel_parallel(i8** [[WORK_FN]]) 162 // CHECK1-NEXT: [[TMP1:%.*]] = zext i1 [[TMP0]] to i8 163 // CHECK1-NEXT: store i8 [[TMP1]], i8* [[EXEC_STATUS]], align 1 164 // CHECK1-NEXT: [[TMP2:%.*]] = load i8*, i8** [[WORK_FN]], align 8 165 // CHECK1-NEXT: [[SHOULD_TERMINATE:%.*]] = icmp eq i8* [[TMP2]], null 166 // CHECK1-NEXT: br i1 [[SHOULD_TERMINATE]], label [[DOTEXIT:%.*]], label [[DOTSELECT_WORKERS:%.*]] 167 // CHECK1: .select.workers: 168 // CHECK1-NEXT: [[TMP3:%.*]] = load i8, i8* [[EXEC_STATUS]], align 1 169 // CHECK1-NEXT: [[IS_ACTIVE:%.*]] = icmp ne i8 [[TMP3]], 0 170 // CHECK1-NEXT: br i1 [[IS_ACTIVE]], label [[DOTEXECUTE_PARALLEL:%.*]], label [[DOTBARRIER_PARALLEL:%.*]] 171 // CHECK1: .execute.parallel: 172 // CHECK1-NEXT: [[TMP4:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1]]) 173 // CHECK1-NEXT: [[TMP5:%.*]] = bitcast i8* [[TMP2]] to void (i16, i32)* 174 // CHECK1-NEXT: call void [[TMP5]](i16 0, i32 [[TMP4]]) 175 // CHECK1-NEXT: br label [[DOTTERMINATE_PARALLEL:%.*]] 176 // CHECK1: .terminate.parallel: 177 // CHECK1-NEXT: call void @__kmpc_kernel_end_parallel() 178 // CHECK1-NEXT: br label [[DOTBARRIER_PARALLEL]] 179 // CHECK1: .barrier.parallel: 180 // CHECK1-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0) 181 // CHECK1-NEXT: br label [[DOTAWAIT_WORK]] 182 // CHECK1: .exit: 183 // CHECK1-NEXT: ret void 184 // 185 // 186 // CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l28 187 // CHECK1-SAME: (i64 [[AA:%.*]]) #[[ATTR1]] { 188 // CHECK1-NEXT: entry: 189 // CHECK1-NEXT: [[AA_ADDR:%.*]] = alloca i64, align 8 190 // CHECK1-NEXT: [[AA_CASTED:%.*]] = alloca i64, align 8 191 // CHECK1-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4 192 // CHECK1-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 4 193 // CHECK1-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4 194 // CHECK1-NEXT: store i64 [[AA]], i64* [[AA_ADDR]], align 8 195 // CHECK1-NEXT: [[CONV:%.*]] = bitcast i64* [[AA_ADDR]] to i16* 196 // CHECK1-NEXT: [[NVPTX_TID:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x() 197 // CHECK1-NEXT: [[NVPTX_NUM_THREADS:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() 198 // CHECK1-NEXT: [[NVPTX_WARP_SIZE:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize() 199 // CHECK1-NEXT: [[THREAD_LIMIT:%.*]] = sub nuw i32 [[NVPTX_NUM_THREADS]], [[NVPTX_WARP_SIZE]] 200 // CHECK1-NEXT: [[TMP0:%.*]] = icmp ult i32 [[NVPTX_TID]], [[THREAD_LIMIT]] 201 // CHECK1-NEXT: br i1 [[TMP0]], label [[DOTWORKER:%.*]], label [[DOTMASTERCHECK:%.*]] 202 // CHECK1: .worker: 203 // CHECK1-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l28_worker() #[[ATTR3]] 204 // CHECK1-NEXT: br label [[DOTEXIT:%.*]] 205 // CHECK1: .mastercheck: 206 // CHECK1-NEXT: [[NVPTX_TID1:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x() 207 // CHECK1-NEXT: [[NVPTX_NUM_THREADS2:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() 208 // CHECK1-NEXT: [[NVPTX_WARP_SIZE3:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize() 209 // CHECK1-NEXT: [[TMP1:%.*]] = sub nuw i32 [[NVPTX_WARP_SIZE3]], 1 210 // CHECK1-NEXT: [[TMP2:%.*]] = sub nuw i32 [[NVPTX_NUM_THREADS2]], 1 211 // CHECK1-NEXT: [[TMP3:%.*]] = xor i32 [[TMP1]], -1 212 // CHECK1-NEXT: [[MASTER_TID:%.*]] = and i32 [[TMP2]], [[TMP3]] 213 // CHECK1-NEXT: [[TMP4:%.*]] = icmp eq i32 [[NVPTX_TID1]], [[MASTER_TID]] 214 // CHECK1-NEXT: br i1 [[TMP4]], label [[DOTMASTER:%.*]], label [[DOTEXIT]] 215 // CHECK1: .master: 216 // CHECK1-NEXT: [[NVPTX_NUM_THREADS4:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() 217 // CHECK1-NEXT: [[NVPTX_WARP_SIZE5:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize() 218 // CHECK1-NEXT: [[THREAD_LIMIT6:%.*]] = sub nuw i32 [[NVPTX_NUM_THREADS4]], [[NVPTX_WARP_SIZE5]] 219 // CHECK1-NEXT: call void @__kmpc_kernel_init(i32 [[THREAD_LIMIT6]], i16 1) 220 // CHECK1-NEXT: [[TMP5:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1]]) 221 // CHECK1-NEXT: [[TMP6:%.*]] = load i16, i16* [[CONV]], align 8 222 // CHECK1-NEXT: [[CONV7:%.*]] = bitcast i64* [[AA_CASTED]] to i16* 223 // CHECK1-NEXT: store i16 [[TMP6]], i16* [[CONV7]], align 2 224 // CHECK1-NEXT: [[TMP7:%.*]] = load i64, i64* [[AA_CASTED]], align 8 225 // CHECK1-NEXT: store i32 [[TMP5]], i32* [[DOTTHREADID_TEMP_]], align 4 226 // CHECK1-NEXT: call void @__omp_outlined__1(i32* [[DOTTHREADID_TEMP_]], i32* [[DOTZERO_ADDR]], i64 [[TMP7]]) #[[ATTR3]] 227 // CHECK1-NEXT: br label [[DOTTERMINATION_NOTIFIER:%.*]] 228 // CHECK1: .termination.notifier: 229 // CHECK1-NEXT: call void @__kmpc_kernel_deinit(i16 1) 230 // CHECK1-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0) 231 // CHECK1-NEXT: br label [[DOTEXIT]] 232 // CHECK1: .exit: 233 // CHECK1-NEXT: ret void 234 // 235 // 236 // CHECK1-LABEL: define {{[^@]+}}@__omp_outlined__1 237 // CHECK1-SAME: (i32* noalias [[DOTGLOBAL_TID_:%.*]], i32* noalias [[DOTBOUND_TID_:%.*]], i64 [[AA:%.*]]) #[[ATTR1]] { 238 // CHECK1-NEXT: entry: 239 // CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8 240 // CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8 241 // CHECK1-NEXT: [[AA_ADDR:%.*]] = alloca i64, align 8 242 // CHECK1-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 8 243 // CHECK1-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 8 244 // CHECK1-NEXT: store i64 [[AA]], i64* [[AA_ADDR]], align 8 245 // CHECK1-NEXT: [[CONV:%.*]] = bitcast i64* [[AA_ADDR]] to i16* 246 // CHECK1-NEXT: store i16 1, i16* [[CONV]], align 8 247 // CHECK1-NEXT: ret void 248 // 249 // 250 // CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l33 251 // CHECK1-SAME: (i64 [[AA:%.*]]) #[[ATTR1]] { 252 // CHECK1-NEXT: entry: 253 // CHECK1-NEXT: [[AA_ADDR:%.*]] = alloca i64, align 8 254 // CHECK1-NEXT: [[AA_CASTED:%.*]] = alloca i64, align 8 255 // CHECK1-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4 256 // CHECK1-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 4 257 // CHECK1-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4 258 // CHECK1-NEXT: store i64 [[AA]], i64* [[AA_ADDR]], align 8 259 // CHECK1-NEXT: [[CONV:%.*]] = bitcast i64* [[AA_ADDR]] to i16* 260 // CHECK1-NEXT: [[NVPTX_NUM_THREADS:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() 261 // CHECK1-NEXT: call void @__kmpc_spmd_kernel_init(i32 [[NVPTX_NUM_THREADS]], i16 1) 262 // CHECK1-NEXT: br label [[DOTEXECUTE:%.*]] 263 // CHECK1: .execute: 264 // CHECK1-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB2:[0-9]+]]) 265 // CHECK1-NEXT: [[TMP1:%.*]] = load i16, i16* [[CONV]], align 8 266 // CHECK1-NEXT: [[CONV1:%.*]] = bitcast i64* [[AA_CASTED]] to i16* 267 // CHECK1-NEXT: store i16 [[TMP1]], i16* [[CONV1]], align 2 268 // CHECK1-NEXT: [[TMP2:%.*]] = load i64, i64* [[AA_CASTED]], align 8 269 // CHECK1-NEXT: store i32 [[TMP0]], i32* [[DOTTHREADID_TEMP_]], align 4 270 // CHECK1-NEXT: call void @__omp_outlined__2(i32* [[DOTTHREADID_TEMP_]], i32* [[DOTZERO_ADDR]], i64 [[TMP2]]) #[[ATTR3]] 271 // CHECK1-NEXT: br label [[DOTOMP_DEINIT:%.*]] 272 // CHECK1: .omp.deinit: 273 // CHECK1-NEXT: call void @__kmpc_spmd_kernel_deinit_v2(i16 1) 274 // CHECK1-NEXT: br label [[DOTEXIT:%.*]] 275 // CHECK1: .exit: 276 // CHECK1-NEXT: ret void 277 // 278 // 279 // CHECK1-LABEL: define {{[^@]+}}@__omp_outlined__2 280 // CHECK1-SAME: (i32* noalias [[DOTGLOBAL_TID_:%.*]], i32* noalias [[DOTBOUND_TID_:%.*]], i64 [[AA:%.*]]) #[[ATTR1]] { 281 // CHECK1-NEXT: entry: 282 // CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8 283 // CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8 284 // CHECK1-NEXT: [[AA_ADDR:%.*]] = alloca i64, align 8 285 // CHECK1-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [1 x i8*], align 8 286 // CHECK1-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 8 287 // CHECK1-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 8 288 // CHECK1-NEXT: store i64 [[AA]], i64* [[AA_ADDR]], align 8 289 // CHECK1-NEXT: [[CONV:%.*]] = bitcast i64* [[AA_ADDR]] to i16* 290 // CHECK1-NEXT: [[TMP0:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[CAPTURED_VARS_ADDRS]], i64 0, i64 0 291 // CHECK1-NEXT: [[TMP1:%.*]] = bitcast i16* [[CONV]] to i8* 292 // CHECK1-NEXT: store i8* [[TMP1]], i8** [[TMP0]], align 8 293 // CHECK1-NEXT: [[TMP2:%.*]] = load i32*, i32** [[DOTGLOBAL_TID__ADDR]], align 8 294 // CHECK1-NEXT: [[TMP3:%.*]] = load i32, i32* [[TMP2]], align 4 295 // CHECK1-NEXT: [[TMP4:%.*]] = bitcast [1 x i8*]* [[CAPTURED_VARS_ADDRS]] to i8** 296 // CHECK1-NEXT: call void @__kmpc_parallel_51(%struct.ident_t* @[[GLOB2]], i32 [[TMP3]], i32 1, i32 -1, i32 -1, i8* bitcast (void (i32*, i32*, i16*)* @__omp_outlined__3 to i8*), i8* null, i8** [[TMP4]], i64 1) 297 // CHECK1-NEXT: ret void 298 // 299 // 300 // CHECK1-LABEL: define {{[^@]+}}@__omp_outlined__3 301 // CHECK1-SAME: (i32* noalias [[DOTGLOBAL_TID_:%.*]], i32* noalias [[DOTBOUND_TID_:%.*]], i16* nonnull align 2 dereferenceable(2) [[AA:%.*]]) #[[ATTR1]] { 302 // CHECK1-NEXT: entry: 303 // CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8 304 // CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8 305 // CHECK1-NEXT: [[AA_ADDR:%.*]] = alloca i16*, align 8 306 // CHECK1-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [1 x i8*], align 8 307 // CHECK1-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 8 308 // CHECK1-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 8 309 // CHECK1-NEXT: store i16* [[AA]], i16** [[AA_ADDR]], align 8 310 // CHECK1-NEXT: [[TMP0:%.*]] = load i16*, i16** [[AA_ADDR]], align 8 311 // CHECK1-NEXT: [[TMP1:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[CAPTURED_VARS_ADDRS]], i64 0, i64 0 312 // CHECK1-NEXT: [[TMP2:%.*]] = bitcast i16* [[TMP0]] to i8* 313 // CHECK1-NEXT: store i8* [[TMP2]], i8** [[TMP1]], align 8 314 // CHECK1-NEXT: [[TMP3:%.*]] = load i32*, i32** [[DOTGLOBAL_TID__ADDR]], align 8 315 // CHECK1-NEXT: [[TMP4:%.*]] = load i32, i32* [[TMP3]], align 4 316 // CHECK1-NEXT: [[TMP5:%.*]] = bitcast [1 x i8*]* [[CAPTURED_VARS_ADDRS]] to i8** 317 // CHECK1-NEXT: call void @__kmpc_parallel_51(%struct.ident_t* @[[GLOB2]], i32 [[TMP4]], i32 1, i32 -1, i32 -1, i8* bitcast (void (i32*, i32*, i16*)* @__omp_outlined__4 to i8*), i8* null, i8** [[TMP5]], i64 1) 318 // CHECK1-NEXT: ret void 319 // 320 // 321 // CHECK1-LABEL: define {{[^@]+}}@__omp_outlined__4 322 // CHECK1-SAME: (i32* noalias [[DOTGLOBAL_TID_:%.*]], i32* noalias [[DOTBOUND_TID_:%.*]], i16* nonnull align 2 dereferenceable(2) [[AA:%.*]]) #[[ATTR1]] { 323 // CHECK1-NEXT: entry: 324 // CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8 325 // CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8 326 // CHECK1-NEXT: [[AA_ADDR:%.*]] = alloca i16*, align 8 327 // CHECK1-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 8 328 // CHECK1-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 8 329 // CHECK1-NEXT: store i16* [[AA]], i16** [[AA_ADDR]], align 8 330 // CHECK1-NEXT: [[TMP0:%.*]] = load i16*, i16** [[AA_ADDR]], align 8 331 // CHECK1-NEXT: store i16 1, i16* [[TMP0]], align 2 332 // CHECK1-NEXT: ret void 333 // 334 // 335 // CHECK2-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l23_worker 336 // CHECK2-SAME: () #[[ATTR0:[0-9]+]] { 337 // CHECK2-NEXT: entry: 338 // CHECK2-NEXT: [[WORK_FN:%.*]] = alloca i8*, align 4 339 // CHECK2-NEXT: [[EXEC_STATUS:%.*]] = alloca i8, align 1 340 // CHECK2-NEXT: store i8* null, i8** [[WORK_FN]], align 4 341 // CHECK2-NEXT: store i8 0, i8* [[EXEC_STATUS]], align 1 342 // CHECK2-NEXT: br label [[DOTAWAIT_WORK:%.*]] 343 // CHECK2: .await.work: 344 // CHECK2-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0) 345 // CHECK2-NEXT: [[TMP0:%.*]] = call i1 @__kmpc_kernel_parallel(i8** [[WORK_FN]]) 346 // CHECK2-NEXT: [[TMP1:%.*]] = zext i1 [[TMP0]] to i8 347 // CHECK2-NEXT: store i8 [[TMP1]], i8* [[EXEC_STATUS]], align 1 348 // CHECK2-NEXT: [[TMP2:%.*]] = load i8*, i8** [[WORK_FN]], align 4 349 // CHECK2-NEXT: [[SHOULD_TERMINATE:%.*]] = icmp eq i8* [[TMP2]], null 350 // CHECK2-NEXT: br i1 [[SHOULD_TERMINATE]], label [[DOTEXIT:%.*]], label [[DOTSELECT_WORKERS:%.*]] 351 // CHECK2: .select.workers: 352 // CHECK2-NEXT: [[TMP3:%.*]] = load i8, i8* [[EXEC_STATUS]], align 1 353 // CHECK2-NEXT: [[IS_ACTIVE:%.*]] = icmp ne i8 [[TMP3]], 0 354 // CHECK2-NEXT: br i1 [[IS_ACTIVE]], label [[DOTEXECUTE_PARALLEL:%.*]], label [[DOTBARRIER_PARALLEL:%.*]] 355 // CHECK2: .execute.parallel: 356 // CHECK2-NEXT: [[TMP4:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1:[0-9]+]]) 357 // CHECK2-NEXT: [[TMP5:%.*]] = bitcast i8* [[TMP2]] to void (i16, i32)* 358 // CHECK2-NEXT: call void [[TMP5]](i16 0, i32 [[TMP4]]) 359 // CHECK2-NEXT: br label [[DOTTERMINATE_PARALLEL:%.*]] 360 // CHECK2: .terminate.parallel: 361 // CHECK2-NEXT: call void @__kmpc_kernel_end_parallel() 362 // CHECK2-NEXT: br label [[DOTBARRIER_PARALLEL]] 363 // CHECK2: .barrier.parallel: 364 // CHECK2-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0) 365 // CHECK2-NEXT: br label [[DOTAWAIT_WORK]] 366 // CHECK2: .exit: 367 // CHECK2-NEXT: ret void 368 // 369 // 370 // CHECK2-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l23 371 // CHECK2-SAME: (i32 [[A:%.*]]) #[[ATTR1:[0-9]+]] { 372 // CHECK2-NEXT: entry: 373 // CHECK2-NEXT: [[A_ADDR:%.*]] = alloca i32, align 4 374 // CHECK2-NEXT: [[A_CASTED:%.*]] = alloca i32, align 4 375 // CHECK2-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4 376 // CHECK2-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 4 377 // CHECK2-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4 378 // CHECK2-NEXT: store i32 [[A]], i32* [[A_ADDR]], align 4 379 // CHECK2-NEXT: [[CONV:%.*]] = bitcast i32* [[A_ADDR]] to i8* 380 // CHECK2-NEXT: [[NVPTX_TID:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x() 381 // CHECK2-NEXT: [[NVPTX_NUM_THREADS:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() 382 // CHECK2-NEXT: [[NVPTX_WARP_SIZE:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize() 383 // CHECK2-NEXT: [[THREAD_LIMIT:%.*]] = sub nuw i32 [[NVPTX_NUM_THREADS]], [[NVPTX_WARP_SIZE]] 384 // CHECK2-NEXT: [[TMP0:%.*]] = icmp ult i32 [[NVPTX_TID]], [[THREAD_LIMIT]] 385 // CHECK2-NEXT: br i1 [[TMP0]], label [[DOTWORKER:%.*]], label [[DOTMASTERCHECK:%.*]] 386 // CHECK2: .worker: 387 // CHECK2-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l23_worker() #[[ATTR3:[0-9]+]] 388 // CHECK2-NEXT: br label [[DOTEXIT:%.*]] 389 // CHECK2: .mastercheck: 390 // CHECK2-NEXT: [[NVPTX_TID1:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x() 391 // CHECK2-NEXT: [[NVPTX_NUM_THREADS2:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() 392 // CHECK2-NEXT: [[NVPTX_WARP_SIZE3:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize() 393 // CHECK2-NEXT: [[TMP1:%.*]] = sub nuw i32 [[NVPTX_WARP_SIZE3]], 1 394 // CHECK2-NEXT: [[TMP2:%.*]] = sub nuw i32 [[NVPTX_NUM_THREADS2]], 1 395 // CHECK2-NEXT: [[TMP3:%.*]] = xor i32 [[TMP1]], -1 396 // CHECK2-NEXT: [[MASTER_TID:%.*]] = and i32 [[TMP2]], [[TMP3]] 397 // CHECK2-NEXT: [[TMP4:%.*]] = icmp eq i32 [[NVPTX_TID1]], [[MASTER_TID]] 398 // CHECK2-NEXT: br i1 [[TMP4]], label [[DOTMASTER:%.*]], label [[DOTEXIT]] 399 // CHECK2: .master: 400 // CHECK2-NEXT: [[NVPTX_NUM_THREADS4:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() 401 // CHECK2-NEXT: [[NVPTX_WARP_SIZE5:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize() 402 // CHECK2-NEXT: [[THREAD_LIMIT6:%.*]] = sub nuw i32 [[NVPTX_NUM_THREADS4]], [[NVPTX_WARP_SIZE5]] 403 // CHECK2-NEXT: call void @__kmpc_kernel_init(i32 [[THREAD_LIMIT6]], i16 1) 404 // CHECK2-NEXT: [[TMP5:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1]]) 405 // CHECK2-NEXT: [[TMP6:%.*]] = load i8, i8* [[CONV]], align 4 406 // CHECK2-NEXT: [[CONV7:%.*]] = bitcast i32* [[A_CASTED]] to i8* 407 // CHECK2-NEXT: store i8 [[TMP6]], i8* [[CONV7]], align 1 408 // CHECK2-NEXT: [[TMP7:%.*]] = load i32, i32* [[A_CASTED]], align 4 409 // CHECK2-NEXT: store i32 [[TMP5]], i32* [[DOTTHREADID_TEMP_]], align 4 410 // CHECK2-NEXT: call void @__omp_outlined__(i32* [[DOTTHREADID_TEMP_]], i32* [[DOTZERO_ADDR]], i32 [[TMP7]]) #[[ATTR3]] 411 // CHECK2-NEXT: br label [[DOTTERMINATION_NOTIFIER:%.*]] 412 // CHECK2: .termination.notifier: 413 // CHECK2-NEXT: call void @__kmpc_kernel_deinit(i16 1) 414 // CHECK2-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0) 415 // CHECK2-NEXT: br label [[DOTEXIT]] 416 // CHECK2: .exit: 417 // CHECK2-NEXT: ret void 418 // 419 // 420 // CHECK2-LABEL: define {{[^@]+}}@__omp_outlined__ 421 // CHECK2-SAME: (i32* noalias [[DOTGLOBAL_TID_:%.*]], i32* noalias [[DOTBOUND_TID_:%.*]], i32 [[A:%.*]]) #[[ATTR1]] { 422 // CHECK2-NEXT: entry: 423 // CHECK2-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 4 424 // CHECK2-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 4 425 // CHECK2-NEXT: [[A_ADDR:%.*]] = alloca i32, align 4 426 // CHECK2-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 4 427 // CHECK2-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 4 428 // CHECK2-NEXT: store i32 [[A]], i32* [[A_ADDR]], align 4 429 // CHECK2-NEXT: [[CONV:%.*]] = bitcast i32* [[A_ADDR]] to i8* 430 // CHECK2-NEXT: store i8 49, i8* [[CONV]], align 4 431 // CHECK2-NEXT: ret void 432 // 433 // 434 // CHECK2-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l28_worker 435 // CHECK2-SAME: () #[[ATTR0]] { 436 // CHECK2-NEXT: entry: 437 // CHECK2-NEXT: [[WORK_FN:%.*]] = alloca i8*, align 4 438 // CHECK2-NEXT: [[EXEC_STATUS:%.*]] = alloca i8, align 1 439 // CHECK2-NEXT: store i8* null, i8** [[WORK_FN]], align 4 440 // CHECK2-NEXT: store i8 0, i8* [[EXEC_STATUS]], align 1 441 // CHECK2-NEXT: br label [[DOTAWAIT_WORK:%.*]] 442 // CHECK2: .await.work: 443 // CHECK2-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0) 444 // CHECK2-NEXT: [[TMP0:%.*]] = call i1 @__kmpc_kernel_parallel(i8** [[WORK_FN]]) 445 // CHECK2-NEXT: [[TMP1:%.*]] = zext i1 [[TMP0]] to i8 446 // CHECK2-NEXT: store i8 [[TMP1]], i8* [[EXEC_STATUS]], align 1 447 // CHECK2-NEXT: [[TMP2:%.*]] = load i8*, i8** [[WORK_FN]], align 4 448 // CHECK2-NEXT: [[SHOULD_TERMINATE:%.*]] = icmp eq i8* [[TMP2]], null 449 // CHECK2-NEXT: br i1 [[SHOULD_TERMINATE]], label [[DOTEXIT:%.*]], label [[DOTSELECT_WORKERS:%.*]] 450 // CHECK2: .select.workers: 451 // CHECK2-NEXT: [[TMP3:%.*]] = load i8, i8* [[EXEC_STATUS]], align 1 452 // CHECK2-NEXT: [[IS_ACTIVE:%.*]] = icmp ne i8 [[TMP3]], 0 453 // CHECK2-NEXT: br i1 [[IS_ACTIVE]], label [[DOTEXECUTE_PARALLEL:%.*]], label [[DOTBARRIER_PARALLEL:%.*]] 454 // CHECK2: .execute.parallel: 455 // CHECK2-NEXT: [[TMP4:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1]]) 456 // CHECK2-NEXT: [[TMP5:%.*]] = bitcast i8* [[TMP2]] to void (i16, i32)* 457 // CHECK2-NEXT: call void [[TMP5]](i16 0, i32 [[TMP4]]) 458 // CHECK2-NEXT: br label [[DOTTERMINATE_PARALLEL:%.*]] 459 // CHECK2: .terminate.parallel: 460 // CHECK2-NEXT: call void @__kmpc_kernel_end_parallel() 461 // CHECK2-NEXT: br label [[DOTBARRIER_PARALLEL]] 462 // CHECK2: .barrier.parallel: 463 // CHECK2-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0) 464 // CHECK2-NEXT: br label [[DOTAWAIT_WORK]] 465 // CHECK2: .exit: 466 // CHECK2-NEXT: ret void 467 // 468 // 469 // CHECK2-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l28 470 // CHECK2-SAME: (i32 [[AA:%.*]]) #[[ATTR1]] { 471 // CHECK2-NEXT: entry: 472 // CHECK2-NEXT: [[AA_ADDR:%.*]] = alloca i32, align 4 473 // CHECK2-NEXT: [[AA_CASTED:%.*]] = alloca i32, align 4 474 // CHECK2-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4 475 // CHECK2-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 4 476 // CHECK2-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4 477 // CHECK2-NEXT: store i32 [[AA]], i32* [[AA_ADDR]], align 4 478 // CHECK2-NEXT: [[CONV:%.*]] = bitcast i32* [[AA_ADDR]] to i16* 479 // CHECK2-NEXT: [[NVPTX_TID:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x() 480 // CHECK2-NEXT: [[NVPTX_NUM_THREADS:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() 481 // CHECK2-NEXT: [[NVPTX_WARP_SIZE:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize() 482 // CHECK2-NEXT: [[THREAD_LIMIT:%.*]] = sub nuw i32 [[NVPTX_NUM_THREADS]], [[NVPTX_WARP_SIZE]] 483 // CHECK2-NEXT: [[TMP0:%.*]] = icmp ult i32 [[NVPTX_TID]], [[THREAD_LIMIT]] 484 // CHECK2-NEXT: br i1 [[TMP0]], label [[DOTWORKER:%.*]], label [[DOTMASTERCHECK:%.*]] 485 // CHECK2: .worker: 486 // CHECK2-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l28_worker() #[[ATTR3]] 487 // CHECK2-NEXT: br label [[DOTEXIT:%.*]] 488 // CHECK2: .mastercheck: 489 // CHECK2-NEXT: [[NVPTX_TID1:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x() 490 // CHECK2-NEXT: [[NVPTX_NUM_THREADS2:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() 491 // CHECK2-NEXT: [[NVPTX_WARP_SIZE3:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize() 492 // CHECK2-NEXT: [[TMP1:%.*]] = sub nuw i32 [[NVPTX_WARP_SIZE3]], 1 493 // CHECK2-NEXT: [[TMP2:%.*]] = sub nuw i32 [[NVPTX_NUM_THREADS2]], 1 494 // CHECK2-NEXT: [[TMP3:%.*]] = xor i32 [[TMP1]], -1 495 // CHECK2-NEXT: [[MASTER_TID:%.*]] = and i32 [[TMP2]], [[TMP3]] 496 // CHECK2-NEXT: [[TMP4:%.*]] = icmp eq i32 [[NVPTX_TID1]], [[MASTER_TID]] 497 // CHECK2-NEXT: br i1 [[TMP4]], label [[DOTMASTER:%.*]], label [[DOTEXIT]] 498 // CHECK2: .master: 499 // CHECK2-NEXT: [[NVPTX_NUM_THREADS4:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() 500 // CHECK2-NEXT: [[NVPTX_WARP_SIZE5:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize() 501 // CHECK2-NEXT: [[THREAD_LIMIT6:%.*]] = sub nuw i32 [[NVPTX_NUM_THREADS4]], [[NVPTX_WARP_SIZE5]] 502 // CHECK2-NEXT: call void @__kmpc_kernel_init(i32 [[THREAD_LIMIT6]], i16 1) 503 // CHECK2-NEXT: [[TMP5:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1]]) 504 // CHECK2-NEXT: [[TMP6:%.*]] = load i16, i16* [[CONV]], align 4 505 // CHECK2-NEXT: [[CONV7:%.*]] = bitcast i32* [[AA_CASTED]] to i16* 506 // CHECK2-NEXT: store i16 [[TMP6]], i16* [[CONV7]], align 2 507 // CHECK2-NEXT: [[TMP7:%.*]] = load i32, i32* [[AA_CASTED]], align 4 508 // CHECK2-NEXT: store i32 [[TMP5]], i32* [[DOTTHREADID_TEMP_]], align 4 509 // CHECK2-NEXT: call void @__omp_outlined__1(i32* [[DOTTHREADID_TEMP_]], i32* [[DOTZERO_ADDR]], i32 [[TMP7]]) #[[ATTR3]] 510 // CHECK2-NEXT: br label [[DOTTERMINATION_NOTIFIER:%.*]] 511 // CHECK2: .termination.notifier: 512 // CHECK2-NEXT: call void @__kmpc_kernel_deinit(i16 1) 513 // CHECK2-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0) 514 // CHECK2-NEXT: br label [[DOTEXIT]] 515 // CHECK2: .exit: 516 // CHECK2-NEXT: ret void 517 // 518 // 519 // CHECK2-LABEL: define {{[^@]+}}@__omp_outlined__1 520 // CHECK2-SAME: (i32* noalias [[DOTGLOBAL_TID_:%.*]], i32* noalias [[DOTBOUND_TID_:%.*]], i32 [[AA:%.*]]) #[[ATTR1]] { 521 // CHECK2-NEXT: entry: 522 // CHECK2-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 4 523 // CHECK2-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 4 524 // CHECK2-NEXT: [[AA_ADDR:%.*]] = alloca i32, align 4 525 // CHECK2-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 4 526 // CHECK2-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 4 527 // CHECK2-NEXT: store i32 [[AA]], i32* [[AA_ADDR]], align 4 528 // CHECK2-NEXT: [[CONV:%.*]] = bitcast i32* [[AA_ADDR]] to i16* 529 // CHECK2-NEXT: store i16 1, i16* [[CONV]], align 4 530 // CHECK2-NEXT: ret void 531 // 532 // 533 // CHECK2-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l33 534 // CHECK2-SAME: (i32 [[AA:%.*]]) #[[ATTR1]] { 535 // CHECK2-NEXT: entry: 536 // CHECK2-NEXT: [[AA_ADDR:%.*]] = alloca i32, align 4 537 // CHECK2-NEXT: [[AA_CASTED:%.*]] = alloca i32, align 4 538 // CHECK2-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4 539 // CHECK2-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 4 540 // CHECK2-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4 541 // CHECK2-NEXT: store i32 [[AA]], i32* [[AA_ADDR]], align 4 542 // CHECK2-NEXT: [[CONV:%.*]] = bitcast i32* [[AA_ADDR]] to i16* 543 // CHECK2-NEXT: [[NVPTX_NUM_THREADS:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() 544 // CHECK2-NEXT: call void @__kmpc_spmd_kernel_init(i32 [[NVPTX_NUM_THREADS]], i16 1) 545 // CHECK2-NEXT: br label [[DOTEXECUTE:%.*]] 546 // CHECK2: .execute: 547 // CHECK2-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB2:[0-9]+]]) 548 // CHECK2-NEXT: [[TMP1:%.*]] = load i16, i16* [[CONV]], align 4 549 // CHECK2-NEXT: [[CONV1:%.*]] = bitcast i32* [[AA_CASTED]] to i16* 550 // CHECK2-NEXT: store i16 [[TMP1]], i16* [[CONV1]], align 2 551 // CHECK2-NEXT: [[TMP2:%.*]] = load i32, i32* [[AA_CASTED]], align 4 552 // CHECK2-NEXT: store i32 [[TMP0]], i32* [[DOTTHREADID_TEMP_]], align 4 553 // CHECK2-NEXT: call void @__omp_outlined__2(i32* [[DOTTHREADID_TEMP_]], i32* [[DOTZERO_ADDR]], i32 [[TMP2]]) #[[ATTR3]] 554 // CHECK2-NEXT: br label [[DOTOMP_DEINIT:%.*]] 555 // CHECK2: .omp.deinit: 556 // CHECK2-NEXT: call void @__kmpc_spmd_kernel_deinit_v2(i16 1) 557 // CHECK2-NEXT: br label [[DOTEXIT:%.*]] 558 // CHECK2: .exit: 559 // CHECK2-NEXT: ret void 560 // 561 // 562 // CHECK2-LABEL: define {{[^@]+}}@__omp_outlined__2 563 // CHECK2-SAME: (i32* noalias [[DOTGLOBAL_TID_:%.*]], i32* noalias [[DOTBOUND_TID_:%.*]], i32 [[AA:%.*]]) #[[ATTR1]] { 564 // CHECK2-NEXT: entry: 565 // CHECK2-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 4 566 // CHECK2-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 4 567 // CHECK2-NEXT: [[AA_ADDR:%.*]] = alloca i32, align 4 568 // CHECK2-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [1 x i8*], align 4 569 // CHECK2-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 4 570 // CHECK2-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 4 571 // CHECK2-NEXT: store i32 [[AA]], i32* [[AA_ADDR]], align 4 572 // CHECK2-NEXT: [[CONV:%.*]] = bitcast i32* [[AA_ADDR]] to i16* 573 // CHECK2-NEXT: [[TMP0:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[CAPTURED_VARS_ADDRS]], i32 0, i32 0 574 // CHECK2-NEXT: [[TMP1:%.*]] = bitcast i16* [[CONV]] to i8* 575 // CHECK2-NEXT: store i8* [[TMP1]], i8** [[TMP0]], align 4 576 // CHECK2-NEXT: [[TMP2:%.*]] = load i32*, i32** [[DOTGLOBAL_TID__ADDR]], align 4 577 // CHECK2-NEXT: [[TMP3:%.*]] = load i32, i32* [[TMP2]], align 4 578 // CHECK2-NEXT: [[TMP4:%.*]] = bitcast [1 x i8*]* [[CAPTURED_VARS_ADDRS]] to i8** 579 // CHECK2-NEXT: call void @__kmpc_parallel_51(%struct.ident_t* @[[GLOB2]], i32 [[TMP3]], i32 1, i32 -1, i32 -1, i8* bitcast (void (i32*, i32*, i16*)* @__omp_outlined__3 to i8*), i8* null, i8** [[TMP4]], i32 1) 580 // CHECK2-NEXT: ret void 581 // 582 // 583 // CHECK2-LABEL: define {{[^@]+}}@__omp_outlined__3 584 // CHECK2-SAME: (i32* noalias [[DOTGLOBAL_TID_:%.*]], i32* noalias [[DOTBOUND_TID_:%.*]], i16* nonnull align 2 dereferenceable(2) [[AA:%.*]]) #[[ATTR1]] { 585 // CHECK2-NEXT: entry: 586 // CHECK2-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 4 587 // CHECK2-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 4 588 // CHECK2-NEXT: [[AA_ADDR:%.*]] = alloca i16*, align 4 589 // CHECK2-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [1 x i8*], align 4 590 // CHECK2-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 4 591 // CHECK2-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 4 592 // CHECK2-NEXT: store i16* [[AA]], i16** [[AA_ADDR]], align 4 593 // CHECK2-NEXT: [[TMP0:%.*]] = load i16*, i16** [[AA_ADDR]], align 4 594 // CHECK2-NEXT: [[TMP1:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[CAPTURED_VARS_ADDRS]], i32 0, i32 0 595 // CHECK2-NEXT: [[TMP2:%.*]] = bitcast i16* [[TMP0]] to i8* 596 // CHECK2-NEXT: store i8* [[TMP2]], i8** [[TMP1]], align 4 597 // CHECK2-NEXT: [[TMP3:%.*]] = load i32*, i32** [[DOTGLOBAL_TID__ADDR]], align 4 598 // CHECK2-NEXT: [[TMP4:%.*]] = load i32, i32* [[TMP3]], align 4 599 // CHECK2-NEXT: [[TMP5:%.*]] = bitcast [1 x i8*]* [[CAPTURED_VARS_ADDRS]] to i8** 600 // CHECK2-NEXT: call void @__kmpc_parallel_51(%struct.ident_t* @[[GLOB2]], i32 [[TMP4]], i32 1, i32 -1, i32 -1, i8* bitcast (void (i32*, i32*, i16*)* @__omp_outlined__4 to i8*), i8* null, i8** [[TMP5]], i32 1) 601 // CHECK2-NEXT: ret void 602 // 603 // 604 // CHECK2-LABEL: define {{[^@]+}}@__omp_outlined__4 605 // CHECK2-SAME: (i32* noalias [[DOTGLOBAL_TID_:%.*]], i32* noalias [[DOTBOUND_TID_:%.*]], i16* nonnull align 2 dereferenceable(2) [[AA:%.*]]) #[[ATTR1]] { 606 // CHECK2-NEXT: entry: 607 // CHECK2-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 4 608 // CHECK2-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 4 609 // CHECK2-NEXT: [[AA_ADDR:%.*]] = alloca i16*, align 4 610 // CHECK2-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 4 611 // CHECK2-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 4 612 // CHECK2-NEXT: store i16* [[AA]], i16** [[AA_ADDR]], align 4 613 // CHECK2-NEXT: [[TMP0:%.*]] = load i16*, i16** [[AA_ADDR]], align 4 614 // CHECK2-NEXT: store i16 1, i16* [[TMP0]], align 2 615 // CHECK2-NEXT: ret void 616 // 617 // 618 // CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l23_worker 619 // CHECK3-SAME: () #[[ATTR0:[0-9]+]] { 620 // CHECK3-NEXT: entry: 621 // CHECK3-NEXT: [[WORK_FN:%.*]] = alloca i8*, align 4 622 // CHECK3-NEXT: [[EXEC_STATUS:%.*]] = alloca i8, align 1 623 // CHECK3-NEXT: store i8* null, i8** [[WORK_FN]], align 4 624 // CHECK3-NEXT: store i8 0, i8* [[EXEC_STATUS]], align 1 625 // CHECK3-NEXT: br label [[DOTAWAIT_WORK:%.*]] 626 // CHECK3: .await.work: 627 // CHECK3-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0) 628 // CHECK3-NEXT: [[TMP0:%.*]] = call i1 @__kmpc_kernel_parallel(i8** [[WORK_FN]]) 629 // CHECK3-NEXT: [[TMP1:%.*]] = zext i1 [[TMP0]] to i8 630 // CHECK3-NEXT: store i8 [[TMP1]], i8* [[EXEC_STATUS]], align 1 631 // CHECK3-NEXT: [[TMP2:%.*]] = load i8*, i8** [[WORK_FN]], align 4 632 // CHECK3-NEXT: [[SHOULD_TERMINATE:%.*]] = icmp eq i8* [[TMP2]], null 633 // CHECK3-NEXT: br i1 [[SHOULD_TERMINATE]], label [[DOTEXIT:%.*]], label [[DOTSELECT_WORKERS:%.*]] 634 // CHECK3: .select.workers: 635 // CHECK3-NEXT: [[TMP3:%.*]] = load i8, i8* [[EXEC_STATUS]], align 1 636 // CHECK3-NEXT: [[IS_ACTIVE:%.*]] = icmp ne i8 [[TMP3]], 0 637 // CHECK3-NEXT: br i1 [[IS_ACTIVE]], label [[DOTEXECUTE_PARALLEL:%.*]], label [[DOTBARRIER_PARALLEL:%.*]] 638 // CHECK3: .execute.parallel: 639 // CHECK3-NEXT: [[TMP4:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1:[0-9]+]]) 640 // CHECK3-NEXT: [[TMP5:%.*]] = bitcast i8* [[TMP2]] to void (i16, i32)* 641 // CHECK3-NEXT: call void [[TMP5]](i16 0, i32 [[TMP4]]) 642 // CHECK3-NEXT: br label [[DOTTERMINATE_PARALLEL:%.*]] 643 // CHECK3: .terminate.parallel: 644 // CHECK3-NEXT: call void @__kmpc_kernel_end_parallel() 645 // CHECK3-NEXT: br label [[DOTBARRIER_PARALLEL]] 646 // CHECK3: .barrier.parallel: 647 // CHECK3-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0) 648 // CHECK3-NEXT: br label [[DOTAWAIT_WORK]] 649 // CHECK3: .exit: 650 // CHECK3-NEXT: ret void 651 // 652 // 653 // CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l23 654 // CHECK3-SAME: (i32 [[A:%.*]]) #[[ATTR1:[0-9]+]] { 655 // CHECK3-NEXT: entry: 656 // CHECK3-NEXT: [[A_ADDR:%.*]] = alloca i32, align 4 657 // CHECK3-NEXT: [[A_CASTED:%.*]] = alloca i32, align 4 658 // CHECK3-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4 659 // CHECK3-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 4 660 // CHECK3-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4 661 // CHECK3-NEXT: store i32 [[A]], i32* [[A_ADDR]], align 4 662 // CHECK3-NEXT: [[CONV:%.*]] = bitcast i32* [[A_ADDR]] to i8* 663 // CHECK3-NEXT: [[NVPTX_TID:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x() 664 // CHECK3-NEXT: [[NVPTX_NUM_THREADS:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() 665 // CHECK3-NEXT: [[NVPTX_WARP_SIZE:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize() 666 // CHECK3-NEXT: [[THREAD_LIMIT:%.*]] = sub nuw i32 [[NVPTX_NUM_THREADS]], [[NVPTX_WARP_SIZE]] 667 // CHECK3-NEXT: [[TMP0:%.*]] = icmp ult i32 [[NVPTX_TID]], [[THREAD_LIMIT]] 668 // CHECK3-NEXT: br i1 [[TMP0]], label [[DOTWORKER:%.*]], label [[DOTMASTERCHECK:%.*]] 669 // CHECK3: .worker: 670 // CHECK3-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l23_worker() #[[ATTR3:[0-9]+]] 671 // CHECK3-NEXT: br label [[DOTEXIT:%.*]] 672 // CHECK3: .mastercheck: 673 // CHECK3-NEXT: [[NVPTX_TID1:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x() 674 // CHECK3-NEXT: [[NVPTX_NUM_THREADS2:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() 675 // CHECK3-NEXT: [[NVPTX_WARP_SIZE3:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize() 676 // CHECK3-NEXT: [[TMP1:%.*]] = sub nuw i32 [[NVPTX_WARP_SIZE3]], 1 677 // CHECK3-NEXT: [[TMP2:%.*]] = sub nuw i32 [[NVPTX_NUM_THREADS2]], 1 678 // CHECK3-NEXT: [[TMP3:%.*]] = xor i32 [[TMP1]], -1 679 // CHECK3-NEXT: [[MASTER_TID:%.*]] = and i32 [[TMP2]], [[TMP3]] 680 // CHECK3-NEXT: [[TMP4:%.*]] = icmp eq i32 [[NVPTX_TID1]], [[MASTER_TID]] 681 // CHECK3-NEXT: br i1 [[TMP4]], label [[DOTMASTER:%.*]], label [[DOTEXIT]] 682 // CHECK3: .master: 683 // CHECK3-NEXT: [[NVPTX_NUM_THREADS4:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() 684 // CHECK3-NEXT: [[NVPTX_WARP_SIZE5:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize() 685 // CHECK3-NEXT: [[THREAD_LIMIT6:%.*]] = sub nuw i32 [[NVPTX_NUM_THREADS4]], [[NVPTX_WARP_SIZE5]] 686 // CHECK3-NEXT: call void @__kmpc_kernel_init(i32 [[THREAD_LIMIT6]], i16 1) 687 // CHECK3-NEXT: [[TMP5:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1]]) 688 // CHECK3-NEXT: [[TMP6:%.*]] = load i8, i8* [[CONV]], align 4 689 // CHECK3-NEXT: [[CONV7:%.*]] = bitcast i32* [[A_CASTED]] to i8* 690 // CHECK3-NEXT: store i8 [[TMP6]], i8* [[CONV7]], align 1 691 // CHECK3-NEXT: [[TMP7:%.*]] = load i32, i32* [[A_CASTED]], align 4 692 // CHECK3-NEXT: store i32 [[TMP5]], i32* [[DOTTHREADID_TEMP_]], align 4 693 // CHECK3-NEXT: call void @__omp_outlined__(i32* [[DOTTHREADID_TEMP_]], i32* [[DOTZERO_ADDR]], i32 [[TMP7]]) #[[ATTR3]] 694 // CHECK3-NEXT: br label [[DOTTERMINATION_NOTIFIER:%.*]] 695 // CHECK3: .termination.notifier: 696 // CHECK3-NEXT: call void @__kmpc_kernel_deinit(i16 1) 697 // CHECK3-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0) 698 // CHECK3-NEXT: br label [[DOTEXIT]] 699 // CHECK3: .exit: 700 // CHECK3-NEXT: ret void 701 // 702 // 703 // CHECK3-LABEL: define {{[^@]+}}@__omp_outlined__ 704 // CHECK3-SAME: (i32* noalias [[DOTGLOBAL_TID_:%.*]], i32* noalias [[DOTBOUND_TID_:%.*]], i32 [[A:%.*]]) #[[ATTR1]] { 705 // CHECK3-NEXT: entry: 706 // CHECK3-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 4 707 // CHECK3-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 4 708 // CHECK3-NEXT: [[A_ADDR:%.*]] = alloca i32, align 4 709 // CHECK3-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 4 710 // CHECK3-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 4 711 // CHECK3-NEXT: store i32 [[A]], i32* [[A_ADDR]], align 4 712 // CHECK3-NEXT: [[CONV:%.*]] = bitcast i32* [[A_ADDR]] to i8* 713 // CHECK3-NEXT: store i8 49, i8* [[CONV]], align 4 714 // CHECK3-NEXT: ret void 715 // 716 // 717 // CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l28_worker 718 // CHECK3-SAME: () #[[ATTR0]] { 719 // CHECK3-NEXT: entry: 720 // CHECK3-NEXT: [[WORK_FN:%.*]] = alloca i8*, align 4 721 // CHECK3-NEXT: [[EXEC_STATUS:%.*]] = alloca i8, align 1 722 // CHECK3-NEXT: store i8* null, i8** [[WORK_FN]], align 4 723 // CHECK3-NEXT: store i8 0, i8* [[EXEC_STATUS]], align 1 724 // CHECK3-NEXT: br label [[DOTAWAIT_WORK:%.*]] 725 // CHECK3: .await.work: 726 // CHECK3-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0) 727 // CHECK3-NEXT: [[TMP0:%.*]] = call i1 @__kmpc_kernel_parallel(i8** [[WORK_FN]]) 728 // CHECK3-NEXT: [[TMP1:%.*]] = zext i1 [[TMP0]] to i8 729 // CHECK3-NEXT: store i8 [[TMP1]], i8* [[EXEC_STATUS]], align 1 730 // CHECK3-NEXT: [[TMP2:%.*]] = load i8*, i8** [[WORK_FN]], align 4 731 // CHECK3-NEXT: [[SHOULD_TERMINATE:%.*]] = icmp eq i8* [[TMP2]], null 732 // CHECK3-NEXT: br i1 [[SHOULD_TERMINATE]], label [[DOTEXIT:%.*]], label [[DOTSELECT_WORKERS:%.*]] 733 // CHECK3: .select.workers: 734 // CHECK3-NEXT: [[TMP3:%.*]] = load i8, i8* [[EXEC_STATUS]], align 1 735 // CHECK3-NEXT: [[IS_ACTIVE:%.*]] = icmp ne i8 [[TMP3]], 0 736 // CHECK3-NEXT: br i1 [[IS_ACTIVE]], label [[DOTEXECUTE_PARALLEL:%.*]], label [[DOTBARRIER_PARALLEL:%.*]] 737 // CHECK3: .execute.parallel: 738 // CHECK3-NEXT: [[TMP4:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1]]) 739 // CHECK3-NEXT: [[TMP5:%.*]] = bitcast i8* [[TMP2]] to void (i16, i32)* 740 // CHECK3-NEXT: call void [[TMP5]](i16 0, i32 [[TMP4]]) 741 // CHECK3-NEXT: br label [[DOTTERMINATE_PARALLEL:%.*]] 742 // CHECK3: .terminate.parallel: 743 // CHECK3-NEXT: call void @__kmpc_kernel_end_parallel() 744 // CHECK3-NEXT: br label [[DOTBARRIER_PARALLEL]] 745 // CHECK3: .barrier.parallel: 746 // CHECK3-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0) 747 // CHECK3-NEXT: br label [[DOTAWAIT_WORK]] 748 // CHECK3: .exit: 749 // CHECK3-NEXT: ret void 750 // 751 // 752 // CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l28 753 // CHECK3-SAME: (i32 [[AA:%.*]]) #[[ATTR1]] { 754 // CHECK3-NEXT: entry: 755 // CHECK3-NEXT: [[AA_ADDR:%.*]] = alloca i32, align 4 756 // CHECK3-NEXT: [[AA_CASTED:%.*]] = alloca i32, align 4 757 // CHECK3-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4 758 // CHECK3-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 4 759 // CHECK3-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4 760 // CHECK3-NEXT: store i32 [[AA]], i32* [[AA_ADDR]], align 4 761 // CHECK3-NEXT: [[CONV:%.*]] = bitcast i32* [[AA_ADDR]] to i16* 762 // CHECK3-NEXT: [[NVPTX_TID:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x() 763 // CHECK3-NEXT: [[NVPTX_NUM_THREADS:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() 764 // CHECK3-NEXT: [[NVPTX_WARP_SIZE:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize() 765 // CHECK3-NEXT: [[THREAD_LIMIT:%.*]] = sub nuw i32 [[NVPTX_NUM_THREADS]], [[NVPTX_WARP_SIZE]] 766 // CHECK3-NEXT: [[TMP0:%.*]] = icmp ult i32 [[NVPTX_TID]], [[THREAD_LIMIT]] 767 // CHECK3-NEXT: br i1 [[TMP0]], label [[DOTWORKER:%.*]], label [[DOTMASTERCHECK:%.*]] 768 // CHECK3: .worker: 769 // CHECK3-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l28_worker() #[[ATTR3]] 770 // CHECK3-NEXT: br label [[DOTEXIT:%.*]] 771 // CHECK3: .mastercheck: 772 // CHECK3-NEXT: [[NVPTX_TID1:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x() 773 // CHECK3-NEXT: [[NVPTX_NUM_THREADS2:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() 774 // CHECK3-NEXT: [[NVPTX_WARP_SIZE3:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize() 775 // CHECK3-NEXT: [[TMP1:%.*]] = sub nuw i32 [[NVPTX_WARP_SIZE3]], 1 776 // CHECK3-NEXT: [[TMP2:%.*]] = sub nuw i32 [[NVPTX_NUM_THREADS2]], 1 777 // CHECK3-NEXT: [[TMP3:%.*]] = xor i32 [[TMP1]], -1 778 // CHECK3-NEXT: [[MASTER_TID:%.*]] = and i32 [[TMP2]], [[TMP3]] 779 // CHECK3-NEXT: [[TMP4:%.*]] = icmp eq i32 [[NVPTX_TID1]], [[MASTER_TID]] 780 // CHECK3-NEXT: br i1 [[TMP4]], label [[DOTMASTER:%.*]], label [[DOTEXIT]] 781 // CHECK3: .master: 782 // CHECK3-NEXT: [[NVPTX_NUM_THREADS4:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() 783 // CHECK3-NEXT: [[NVPTX_WARP_SIZE5:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize() 784 // CHECK3-NEXT: [[THREAD_LIMIT6:%.*]] = sub nuw i32 [[NVPTX_NUM_THREADS4]], [[NVPTX_WARP_SIZE5]] 785 // CHECK3-NEXT: call void @__kmpc_kernel_init(i32 [[THREAD_LIMIT6]], i16 1) 786 // CHECK3-NEXT: [[TMP5:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1]]) 787 // CHECK3-NEXT: [[TMP6:%.*]] = load i16, i16* [[CONV]], align 4 788 // CHECK3-NEXT: [[CONV7:%.*]] = bitcast i32* [[AA_CASTED]] to i16* 789 // CHECK3-NEXT: store i16 [[TMP6]], i16* [[CONV7]], align 2 790 // CHECK3-NEXT: [[TMP7:%.*]] = load i32, i32* [[AA_CASTED]], align 4 791 // CHECK3-NEXT: store i32 [[TMP5]], i32* [[DOTTHREADID_TEMP_]], align 4 792 // CHECK3-NEXT: call void @__omp_outlined__1(i32* [[DOTTHREADID_TEMP_]], i32* [[DOTZERO_ADDR]], i32 [[TMP7]]) #[[ATTR3]] 793 // CHECK3-NEXT: br label [[DOTTERMINATION_NOTIFIER:%.*]] 794 // CHECK3: .termination.notifier: 795 // CHECK3-NEXT: call void @__kmpc_kernel_deinit(i16 1) 796 // CHECK3-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0) 797 // CHECK3-NEXT: br label [[DOTEXIT]] 798 // CHECK3: .exit: 799 // CHECK3-NEXT: ret void 800 // 801 // 802 // CHECK3-LABEL: define {{[^@]+}}@__omp_outlined__1 803 // CHECK3-SAME: (i32* noalias [[DOTGLOBAL_TID_:%.*]], i32* noalias [[DOTBOUND_TID_:%.*]], i32 [[AA:%.*]]) #[[ATTR1]] { 804 // CHECK3-NEXT: entry: 805 // CHECK3-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 4 806 // CHECK3-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 4 807 // CHECK3-NEXT: [[AA_ADDR:%.*]] = alloca i32, align 4 808 // CHECK3-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 4 809 // CHECK3-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 4 810 // CHECK3-NEXT: store i32 [[AA]], i32* [[AA_ADDR]], align 4 811 // CHECK3-NEXT: [[CONV:%.*]] = bitcast i32* [[AA_ADDR]] to i16* 812 // CHECK3-NEXT: store i16 1, i16* [[CONV]], align 4 813 // CHECK3-NEXT: ret void 814 // 815 // 816 // CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l33 817 // CHECK3-SAME: (i32 [[AA:%.*]]) #[[ATTR1]] { 818 // CHECK3-NEXT: entry: 819 // CHECK3-NEXT: [[AA_ADDR:%.*]] = alloca i32, align 4 820 // CHECK3-NEXT: [[AA_CASTED:%.*]] = alloca i32, align 4 821 // CHECK3-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4 822 // CHECK3-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 4 823 // CHECK3-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4 824 // CHECK3-NEXT: store i32 [[AA]], i32* [[AA_ADDR]], align 4 825 // CHECK3-NEXT: [[CONV:%.*]] = bitcast i32* [[AA_ADDR]] to i16* 826 // CHECK3-NEXT: [[NVPTX_NUM_THREADS:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() 827 // CHECK3-NEXT: call void @__kmpc_spmd_kernel_init(i32 [[NVPTX_NUM_THREADS]], i16 1) 828 // CHECK3-NEXT: br label [[DOTEXECUTE:%.*]] 829 // CHECK3: .execute: 830 // CHECK3-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB2:[0-9]+]]) 831 // CHECK3-NEXT: [[TMP1:%.*]] = load i16, i16* [[CONV]], align 4 832 // CHECK3-NEXT: [[CONV1:%.*]] = bitcast i32* [[AA_CASTED]] to i16* 833 // CHECK3-NEXT: store i16 [[TMP1]], i16* [[CONV1]], align 2 834 // CHECK3-NEXT: [[TMP2:%.*]] = load i32, i32* [[AA_CASTED]], align 4 835 // CHECK3-NEXT: store i32 [[TMP0]], i32* [[DOTTHREADID_TEMP_]], align 4 836 // CHECK3-NEXT: call void @__omp_outlined__2(i32* [[DOTTHREADID_TEMP_]], i32* [[DOTZERO_ADDR]], i32 [[TMP2]]) #[[ATTR3]] 837 // CHECK3-NEXT: br label [[DOTOMP_DEINIT:%.*]] 838 // CHECK3: .omp.deinit: 839 // CHECK3-NEXT: call void @__kmpc_spmd_kernel_deinit_v2(i16 1) 840 // CHECK3-NEXT: br label [[DOTEXIT:%.*]] 841 // CHECK3: .exit: 842 // CHECK3-NEXT: ret void 843 // 844 // 845 // CHECK3-LABEL: define {{[^@]+}}@__omp_outlined__2 846 // CHECK3-SAME: (i32* noalias [[DOTGLOBAL_TID_:%.*]], i32* noalias [[DOTBOUND_TID_:%.*]], i32 [[AA:%.*]]) #[[ATTR1]] { 847 // CHECK3-NEXT: entry: 848 // CHECK3-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 4 849 // CHECK3-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 4 850 // CHECK3-NEXT: [[AA_ADDR:%.*]] = 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: store i32 [[AA]], i32* [[AA_ADDR]], align 4 855 // CHECK3-NEXT: [[CONV:%.*]] = bitcast i32* [[AA_ADDR]] to i16* 856 // CHECK3-NEXT: [[TMP0:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[CAPTURED_VARS_ADDRS]], i32 0, i32 0 857 // CHECK3-NEXT: [[TMP1:%.*]] = bitcast i16* [[CONV]] to i8* 858 // CHECK3-NEXT: store i8* [[TMP1]], i8** [[TMP0]], align 4 859 // CHECK3-NEXT: [[TMP2:%.*]] = load i32*, i32** [[DOTGLOBAL_TID__ADDR]], align 4 860 // CHECK3-NEXT: [[TMP3:%.*]] = load i32, i32* [[TMP2]], align 4 861 // CHECK3-NEXT: [[TMP4:%.*]] = bitcast [1 x i8*]* [[CAPTURED_VARS_ADDRS]] to i8** 862 // CHECK3-NEXT: call void @__kmpc_parallel_51(%struct.ident_t* @[[GLOB2]], i32 [[TMP3]], i32 1, i32 -1, i32 -1, i8* bitcast (void (i32*, i32*, i16*)* @__omp_outlined__3 to i8*), i8* null, i8** [[TMP4]], i32 1) 863 // CHECK3-NEXT: ret void 864 // 865 // 866 // CHECK3-LABEL: define {{[^@]+}}@__omp_outlined__3 867 // CHECK3-SAME: (i32* noalias [[DOTGLOBAL_TID_:%.*]], i32* noalias [[DOTBOUND_TID_:%.*]], i16* nonnull align 2 dereferenceable(2) [[AA:%.*]]) #[[ATTR1]] { 868 // CHECK3-NEXT: entry: 869 // CHECK3-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 4 870 // CHECK3-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 4 871 // CHECK3-NEXT: [[AA_ADDR:%.*]] = alloca i16*, align 4 872 // CHECK3-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [1 x i8*], align 4 873 // CHECK3-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 4 874 // CHECK3-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 4 875 // CHECK3-NEXT: store i16* [[AA]], i16** [[AA_ADDR]], align 4 876 // CHECK3-NEXT: [[TMP0:%.*]] = load i16*, i16** [[AA_ADDR]], align 4 877 // CHECK3-NEXT: [[TMP1:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[CAPTURED_VARS_ADDRS]], i32 0, i32 0 878 // CHECK3-NEXT: [[TMP2:%.*]] = bitcast i16* [[TMP0]] to i8* 879 // CHECK3-NEXT: store i8* [[TMP2]], i8** [[TMP1]], align 4 880 // CHECK3-NEXT: [[TMP3:%.*]] = load i32*, i32** [[DOTGLOBAL_TID__ADDR]], align 4 881 // CHECK3-NEXT: [[TMP4:%.*]] = load i32, i32* [[TMP3]], align 4 882 // CHECK3-NEXT: [[TMP5:%.*]] = bitcast [1 x i8*]* [[CAPTURED_VARS_ADDRS]] to i8** 883 // CHECK3-NEXT: call void @__kmpc_parallel_51(%struct.ident_t* @[[GLOB2]], i32 [[TMP4]], i32 1, i32 -1, i32 -1, i8* bitcast (void (i32*, i32*, i16*)* @__omp_outlined__4 to i8*), i8* null, i8** [[TMP5]], i32 1) 884 // CHECK3-NEXT: ret void 885 // 886 // 887 // CHECK3-LABEL: define {{[^@]+}}@__omp_outlined__4 888 // CHECK3-SAME: (i32* noalias [[DOTGLOBAL_TID_:%.*]], i32* noalias [[DOTBOUND_TID_:%.*]], i16* nonnull align 2 dereferenceable(2) [[AA:%.*]]) #[[ATTR1]] { 889 // CHECK3-NEXT: entry: 890 // CHECK3-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 4 891 // CHECK3-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 4 892 // CHECK3-NEXT: [[AA_ADDR:%.*]] = alloca i16*, align 4 893 // CHECK3-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 4 894 // CHECK3-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 4 895 // CHECK3-NEXT: store i16* [[AA]], i16** [[AA_ADDR]], align 4 896 // CHECK3-NEXT: [[TMP0:%.*]] = load i16*, i16** [[AA_ADDR]], align 4 897 // CHECK3-NEXT: store i16 1, i16* [[TMP0]], align 2 898 // CHECK3-NEXT: ret void 899 // 900