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