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 -fopenmp-cuda-mode -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 -fopenmp-cuda-mode -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 - -disable-llvm-optzns | FileCheck %s --check-prefix=CHECK1 5 // RUN: %clang_cc1 -verify -fopenmp -fopenmp-cuda-mode -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 -fopenmp-cuda-mode -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 - -disable-llvm-optzns | FileCheck %s --check-prefix=CHECK2 7 // RUN: %clang_cc1 -verify -fopenmp -fopenmp-cuda-mode -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 - -disable-llvm-optzns | FileCheck %s --check-prefix=CHECK3 8 9 // expected-no-diagnostics 10 #ifndef HEADER 11 #define HEADER 12 13 void work(int *C) { 14 #pragma omp atomic 15 ++(*C); 16 } 17 18 void use(int *C) { 19 #pragma omp parallel num_threads(2) 20 work(C); 21 } 22 23 int main() { 24 int C = 0; 25 #pragma omp target map(C) 26 { 27 use(&C); 28 #pragma omp parallel num_threads(2) 29 use(&C); 30 } 31 32 return C; 33 } 34 35 #endif 36 // CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l25_worker 37 // CHECK1-SAME: () #[[ATTR0:[0-9]+]] { 38 // CHECK1-NEXT: entry: 39 // CHECK1-NEXT: [[WORK_FN:%.*]] = alloca i8*, align 8 40 // CHECK1-NEXT: [[EXEC_STATUS:%.*]] = alloca i8, align 1 41 // CHECK1-NEXT: store i8* null, i8** [[WORK_FN]], align 8 42 // CHECK1-NEXT: store i8 0, i8* [[EXEC_STATUS]], align 1 43 // CHECK1-NEXT: br label [[DOTAWAIT_WORK:%.*]] 44 // CHECK1: .await.work: 45 // CHECK1-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0) 46 // CHECK1-NEXT: [[TMP0:%.*]] = call i1 @__kmpc_kernel_parallel(i8** [[WORK_FN]]) 47 // CHECK1-NEXT: [[TMP1:%.*]] = zext i1 [[TMP0]] to i8 48 // CHECK1-NEXT: store i8 [[TMP1]], i8* [[EXEC_STATUS]], align 1 49 // CHECK1-NEXT: [[TMP2:%.*]] = load i8*, i8** [[WORK_FN]], align 8 50 // CHECK1-NEXT: [[SHOULD_TERMINATE:%.*]] = icmp eq i8* [[TMP2]], null 51 // CHECK1-NEXT: br i1 [[SHOULD_TERMINATE]], label [[DOTEXIT:%.*]], label [[DOTSELECT_WORKERS:%.*]] 52 // CHECK1: .select.workers: 53 // CHECK1-NEXT: [[TMP3:%.*]] = load i8, i8* [[EXEC_STATUS]], align 1 54 // CHECK1-NEXT: [[IS_ACTIVE:%.*]] = icmp ne i8 [[TMP3]], 0 55 // CHECK1-NEXT: br i1 [[IS_ACTIVE]], label [[DOTEXECUTE_PARALLEL:%.*]], label [[DOTBARRIER_PARALLEL:%.*]] 56 // CHECK1: .execute.parallel: 57 // CHECK1-NEXT: [[TMP4:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1:[0-9]+]]) 58 // CHECK1-NEXT: [[TMP5:%.*]] = load i8*, i8** [[WORK_FN]], align 8 59 // CHECK1-NEXT: [[WORK_MATCH:%.*]] = icmp eq i8* [[TMP5]], bitcast (void (i16, i32)* @__omp_outlined___wrapper to i8*) 60 // CHECK1-NEXT: br i1 [[WORK_MATCH]], label [[DOTEXECUTE_FN:%.*]], label [[DOTCHECK_NEXT:%.*]] 61 // CHECK1: .execute.fn: 62 // CHECK1-NEXT: call void @__omp_outlined___wrapper(i16 0, i32 [[TMP4]]) #[[ATTR4:[0-9]+]] 63 // CHECK1-NEXT: br label [[DOTTERMINATE_PARALLEL:%.*]] 64 // CHECK1: .check.next: 65 // CHECK1-NEXT: [[TMP6:%.*]] = bitcast i8* [[TMP2]] to void (i16, i32)* 66 // CHECK1-NEXT: call void [[TMP6]](i16 0, i32 [[TMP4]]) 67 // CHECK1-NEXT: br label [[DOTTERMINATE_PARALLEL]] 68 // CHECK1: .terminate.parallel: 69 // CHECK1-NEXT: call void @__kmpc_kernel_end_parallel() 70 // CHECK1-NEXT: br label [[DOTBARRIER_PARALLEL]] 71 // CHECK1: .barrier.parallel: 72 // CHECK1-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0) 73 // CHECK1-NEXT: br label [[DOTAWAIT_WORK]] 74 // CHECK1: .exit: 75 // CHECK1-NEXT: ret void 76 // 77 // 78 // CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l25 79 // CHECK1-SAME: (i32* nonnull align 4 dereferenceable(4) [[C:%.*]]) #[[ATTR1:[0-9]+]] { 80 // CHECK1-NEXT: entry: 81 // CHECK1-NEXT: [[C_ADDR:%.*]] = alloca i32*, align 8 82 // CHECK1-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [1 x i8*], align 8 83 // CHECK1-NEXT: store i32* [[C]], i32** [[C_ADDR]], align 8 84 // CHECK1-NEXT: [[TMP0:%.*]] = load i32*, i32** [[C_ADDR]], align 8 85 // CHECK1-NEXT: [[NVPTX_TID:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x() 86 // CHECK1-NEXT: [[NVPTX_NUM_THREADS:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() 87 // CHECK1-NEXT: [[NVPTX_WARP_SIZE:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize() 88 // CHECK1-NEXT: [[THREAD_LIMIT:%.*]] = sub nuw i32 [[NVPTX_NUM_THREADS]], [[NVPTX_WARP_SIZE]] 89 // CHECK1-NEXT: [[TMP1:%.*]] = icmp ult i32 [[NVPTX_TID]], [[THREAD_LIMIT]] 90 // CHECK1-NEXT: br i1 [[TMP1]], label [[DOTWORKER:%.*]], label [[DOTMASTERCHECK:%.*]] 91 // CHECK1: .worker: 92 // CHECK1-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l25_worker() #[[ATTR4]] 93 // CHECK1-NEXT: br label [[DOTEXIT:%.*]] 94 // CHECK1: .mastercheck: 95 // CHECK1-NEXT: [[NVPTX_TID1:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x() 96 // CHECK1-NEXT: [[NVPTX_NUM_THREADS2:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() 97 // CHECK1-NEXT: [[NVPTX_WARP_SIZE3:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize() 98 // CHECK1-NEXT: [[TMP2:%.*]] = sub nuw i32 [[NVPTX_WARP_SIZE3]], 1 99 // CHECK1-NEXT: [[TMP3:%.*]] = sub nuw i32 [[NVPTX_NUM_THREADS2]], 1 100 // CHECK1-NEXT: [[TMP4:%.*]] = xor i32 [[TMP2]], -1 101 // CHECK1-NEXT: [[MASTER_TID:%.*]] = and i32 [[TMP3]], [[TMP4]] 102 // CHECK1-NEXT: [[TMP5:%.*]] = icmp eq i32 [[NVPTX_TID1]], [[MASTER_TID]] 103 // CHECK1-NEXT: br i1 [[TMP5]], label [[DOTMASTER:%.*]], label [[DOTEXIT]] 104 // CHECK1: .master: 105 // CHECK1-NEXT: [[NVPTX_NUM_THREADS4:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() 106 // CHECK1-NEXT: [[NVPTX_WARP_SIZE5:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize() 107 // CHECK1-NEXT: [[THREAD_LIMIT6:%.*]] = sub nuw i32 [[NVPTX_NUM_THREADS4]], [[NVPTX_WARP_SIZE5]] 108 // CHECK1-NEXT: call void @__kmpc_kernel_init(i32 [[THREAD_LIMIT6]], i16 1) 109 // CHECK1-NEXT: call void @__kmpc_data_sharing_init_stack() 110 // CHECK1-NEXT: [[TMP6:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1]]) 111 // CHECK1-NEXT: call void @_Z3usePi(i32* [[TMP0]]) #[[ATTR7:[0-9]+]] 112 // CHECK1-NEXT: call void @__kmpc_push_num_threads(%struct.ident_t* @[[GLOB1]], i32 [[TMP6]], i32 2) 113 // CHECK1-NEXT: [[TMP7:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[CAPTURED_VARS_ADDRS]], i64 0, i64 0 114 // CHECK1-NEXT: [[TMP8:%.*]] = bitcast i32* [[TMP0]] to i8* 115 // CHECK1-NEXT: store i8* [[TMP8]], i8** [[TMP7]], align 8 116 // CHECK1-NEXT: [[TMP9:%.*]] = bitcast [1 x i8*]* [[CAPTURED_VARS_ADDRS]] to i8** 117 // CHECK1-NEXT: call void @__kmpc_parallel_51(%struct.ident_t* @[[GLOB1]], i32 [[TMP6]], i32 1, i32 -1, i32 -1, i8* bitcast (void (i32*, i32*, i32*)* @__omp_outlined__ to i8*), i8* bitcast (void (i16, i32)* @__omp_outlined___wrapper to i8*), i8** [[TMP9]], i64 1) 118 // CHECK1-NEXT: br label [[DOTTERMINATION_NOTIFIER:%.*]] 119 // CHECK1: .termination.notifier: 120 // CHECK1-NEXT: call void @__kmpc_kernel_deinit(i16 1) 121 // CHECK1-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0) 122 // CHECK1-NEXT: br label [[DOTEXIT]] 123 // CHECK1: .exit: 124 // CHECK1-NEXT: ret void 125 // 126 // 127 // CHECK1-LABEL: define {{[^@]+}}@_Z3usePi 128 // CHECK1-SAME: (i32* [[C:%.*]]) #[[ATTR3:[0-9]+]] { 129 // CHECK1-NEXT: entry: 130 // CHECK1-NEXT: [[C_ADDR:%.*]] = alloca i32*, align 8 131 // CHECK1-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [1 x i8*], align 8 132 // CHECK1-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB2:[0-9]+]]) 133 // CHECK1-NEXT: store i32* [[C]], i32** [[C_ADDR]], align 8 134 // CHECK1-NEXT: call void @__kmpc_push_num_threads(%struct.ident_t* @[[GLOB2]], i32 [[TMP0]], i32 2) 135 // CHECK1-NEXT: [[TMP1:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[CAPTURED_VARS_ADDRS]], i64 0, i64 0 136 // CHECK1-NEXT: [[TMP2:%.*]] = bitcast i32** [[C_ADDR]] to i8* 137 // CHECK1-NEXT: store i8* [[TMP2]], i8** [[TMP1]], align 8 138 // CHECK1-NEXT: [[TMP3:%.*]] = bitcast [1 x i8*]* [[CAPTURED_VARS_ADDRS]] to i8** 139 // CHECK1-NEXT: call void @__kmpc_parallel_51(%struct.ident_t* @[[GLOB2]], i32 [[TMP0]], i32 1, i32 -1, i32 -1, i8* bitcast (void (i32*, i32*, i32**)* @__omp_outlined__1 to i8*), i8* bitcast (void (i16, i32)* @__omp_outlined__1_wrapper to i8*), i8** [[TMP3]], i64 1) 140 // CHECK1-NEXT: ret void 141 // 142 // 143 // CHECK1-LABEL: define {{[^@]+}}@__omp_outlined__ 144 // CHECK1-SAME: (i32* noalias [[DOTGLOBAL_TID_:%.*]], i32* noalias [[DOTBOUND_TID_:%.*]], i32* nonnull align 4 dereferenceable(4) [[C:%.*]]) #[[ATTR1]] { 145 // CHECK1-NEXT: entry: 146 // CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8 147 // CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8 148 // CHECK1-NEXT: [[C_ADDR:%.*]] = alloca i32*, align 8 149 // CHECK1-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 8 150 // CHECK1-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 8 151 // CHECK1-NEXT: store i32* [[C]], i32** [[C_ADDR]], align 8 152 // CHECK1-NEXT: [[TMP0:%.*]] = load i32*, i32** [[C_ADDR]], align 8 153 // CHECK1-NEXT: call void @_Z3usePi(i32* [[TMP0]]) #[[ATTR7]] 154 // CHECK1-NEXT: ret void 155 // 156 // 157 // CHECK1-LABEL: define {{[^@]+}}@__omp_outlined___wrapper 158 // CHECK1-SAME: (i16 zeroext [[TMP0:%.*]], i32 [[TMP1:%.*]]) #[[ATTR0]] { 159 // CHECK1-NEXT: entry: 160 // CHECK1-NEXT: [[DOTADDR:%.*]] = alloca i16, align 2 161 // CHECK1-NEXT: [[DOTADDR1:%.*]] = alloca i32, align 4 162 // CHECK1-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4 163 // CHECK1-NEXT: [[GLOBAL_ARGS:%.*]] = alloca i8**, align 8 164 // CHECK1-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4 165 // CHECK1-NEXT: store i16 [[TMP0]], i16* [[DOTADDR]], align 2 166 // CHECK1-NEXT: store i32 [[TMP1]], i32* [[DOTADDR1]], align 4 167 // CHECK1-NEXT: call void @__kmpc_get_shared_variables(i8*** [[GLOBAL_ARGS]]) 168 // CHECK1-NEXT: [[TMP2:%.*]] = load i8**, i8*** [[GLOBAL_ARGS]], align 8 169 // CHECK1-NEXT: [[TMP3:%.*]] = getelementptr inbounds i8*, i8** [[TMP2]], i64 0 170 // CHECK1-NEXT: [[TMP4:%.*]] = bitcast i8** [[TMP3]] to i32** 171 // CHECK1-NEXT: [[TMP5:%.*]] = load i32*, i32** [[TMP4]], align 8 172 // CHECK1-NEXT: call void @__omp_outlined__(i32* [[DOTADDR1]], i32* [[DOTZERO_ADDR]], i32* [[TMP5]]) #[[ATTR4]] 173 // CHECK1-NEXT: ret void 174 // 175 // 176 // CHECK1-LABEL: define {{[^@]+}}@__omp_outlined__1 177 // CHECK1-SAME: (i32* noalias [[DOTGLOBAL_TID_:%.*]], i32* noalias [[DOTBOUND_TID_:%.*]], i32** nonnull align 8 dereferenceable(8) [[C:%.*]]) #[[ATTR1]] { 178 // CHECK1-NEXT: entry: 179 // CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8 180 // CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8 181 // CHECK1-NEXT: [[C_ADDR:%.*]] = alloca i32**, align 8 182 // CHECK1-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 8 183 // CHECK1-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 8 184 // CHECK1-NEXT: store i32** [[C]], i32*** [[C_ADDR]], align 8 185 // CHECK1-NEXT: [[TMP0:%.*]] = load i32**, i32*** [[C_ADDR]], align 8 186 // CHECK1-NEXT: [[TMP1:%.*]] = load i32*, i32** [[TMP0]], align 8 187 // CHECK1-NEXT: call void @_Z4workPi(i32* [[TMP1]]) #[[ATTR7]] 188 // CHECK1-NEXT: ret void 189 // 190 // 191 // CHECK1-LABEL: define {{[^@]+}}@_Z4workPi 192 // CHECK1-SAME: (i32* [[C:%.*]]) #[[ATTR3]] { 193 // CHECK1-NEXT: entry: 194 // CHECK1-NEXT: [[C_ADDR:%.*]] = alloca i32*, align 8 195 // CHECK1-NEXT: [[ATOMIC_TEMP:%.*]] = alloca i32, align 4 196 // CHECK1-NEXT: [[ATOMIC_TEMP1:%.*]] = alloca i32, align 4 197 // CHECK1-NEXT: store i32* [[C]], i32** [[C_ADDR]], align 8 198 // CHECK1-NEXT: [[TMP0:%.*]] = load i32*, i32** [[C_ADDR]], align 8 199 // CHECK1-NEXT: [[TMP1:%.*]] = bitcast i32* [[TMP0]] to i8* 200 // CHECK1-NEXT: [[TMP2:%.*]] = bitcast i32* [[ATOMIC_TEMP]] to i8* 201 // CHECK1-NEXT: call void @__atomic_load(i64 4, i8* [[TMP1]], i8* [[TMP2]], i32 0) #[[ATTR7]] 202 // CHECK1-NEXT: br label [[ATOMIC_CONT:%.*]] 203 // CHECK1: atomic_cont: 204 // CHECK1-NEXT: [[TMP3:%.*]] = load i32, i32* [[ATOMIC_TEMP]], align 4 205 // CHECK1-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP3]], 1 206 // CHECK1-NEXT: store i32 [[ADD]], i32* [[ATOMIC_TEMP1]], align 4 207 // CHECK1-NEXT: [[TMP4:%.*]] = bitcast i32* [[TMP0]] to i8* 208 // CHECK1-NEXT: [[TMP5:%.*]] = bitcast i32* [[ATOMIC_TEMP]] to i8* 209 // CHECK1-NEXT: [[TMP6:%.*]] = bitcast i32* [[ATOMIC_TEMP1]] to i8* 210 // CHECK1-NEXT: [[CALL:%.*]] = call zeroext i1 @__atomic_compare_exchange(i64 4, i8* [[TMP4]], i8* [[TMP5]], i8* [[TMP6]], i32 0, i32 0) #[[ATTR7]] 211 // CHECK1-NEXT: br i1 [[CALL]], label [[ATOMIC_EXIT:%.*]], label [[ATOMIC_CONT]] 212 // CHECK1: atomic_exit: 213 // CHECK1-NEXT: ret void 214 // 215 // 216 // CHECK1-LABEL: define {{[^@]+}}@__omp_outlined__1_wrapper 217 // CHECK1-SAME: (i16 zeroext [[TMP0:%.*]], i32 [[TMP1:%.*]]) #[[ATTR0]] { 218 // CHECK1-NEXT: entry: 219 // CHECK1-NEXT: [[DOTADDR:%.*]] = alloca i16, align 2 220 // CHECK1-NEXT: [[DOTADDR1:%.*]] = alloca i32, align 4 221 // CHECK1-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4 222 // CHECK1-NEXT: [[GLOBAL_ARGS:%.*]] = alloca i8**, align 8 223 // CHECK1-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4 224 // CHECK1-NEXT: store i16 [[TMP0]], i16* [[DOTADDR]], align 2 225 // CHECK1-NEXT: store i32 [[TMP1]], i32* [[DOTADDR1]], align 4 226 // CHECK1-NEXT: call void @__kmpc_get_shared_variables(i8*** [[GLOBAL_ARGS]]) 227 // CHECK1-NEXT: [[TMP2:%.*]] = load i8**, i8*** [[GLOBAL_ARGS]], align 8 228 // CHECK1-NEXT: [[TMP3:%.*]] = getelementptr inbounds i8*, i8** [[TMP2]], i64 0 229 // CHECK1-NEXT: [[TMP4:%.*]] = bitcast i8** [[TMP3]] to i32*** 230 // CHECK1-NEXT: [[TMP5:%.*]] = load i32**, i32*** [[TMP4]], align 8 231 // CHECK1-NEXT: call void @__omp_outlined__1(i32* [[DOTADDR1]], i32* [[DOTZERO_ADDR]], i32** [[TMP5]]) #[[ATTR4]] 232 // CHECK1-NEXT: ret void 233 // 234 // 235 // CHECK2-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l25_worker 236 // CHECK2-SAME: () #[[ATTR0:[0-9]+]] { 237 // CHECK2-NEXT: entry: 238 // CHECK2-NEXT: [[WORK_FN:%.*]] = alloca i8*, align 4 239 // CHECK2-NEXT: [[EXEC_STATUS:%.*]] = alloca i8, align 1 240 // CHECK2-NEXT: store i8* null, i8** [[WORK_FN]], align 4 241 // CHECK2-NEXT: store i8 0, i8* [[EXEC_STATUS]], align 1 242 // CHECK2-NEXT: br label [[DOTAWAIT_WORK:%.*]] 243 // CHECK2: .await.work: 244 // CHECK2-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0) 245 // CHECK2-NEXT: [[TMP0:%.*]] = call i1 @__kmpc_kernel_parallel(i8** [[WORK_FN]]) 246 // CHECK2-NEXT: [[TMP1:%.*]] = zext i1 [[TMP0]] to i8 247 // CHECK2-NEXT: store i8 [[TMP1]], i8* [[EXEC_STATUS]], align 1 248 // CHECK2-NEXT: [[TMP2:%.*]] = load i8*, i8** [[WORK_FN]], align 4 249 // CHECK2-NEXT: [[SHOULD_TERMINATE:%.*]] = icmp eq i8* [[TMP2]], null 250 // CHECK2-NEXT: br i1 [[SHOULD_TERMINATE]], label [[DOTEXIT:%.*]], label [[DOTSELECT_WORKERS:%.*]] 251 // CHECK2: .select.workers: 252 // CHECK2-NEXT: [[TMP3:%.*]] = load i8, i8* [[EXEC_STATUS]], align 1 253 // CHECK2-NEXT: [[IS_ACTIVE:%.*]] = icmp ne i8 [[TMP3]], 0 254 // CHECK2-NEXT: br i1 [[IS_ACTIVE]], label [[DOTEXECUTE_PARALLEL:%.*]], label [[DOTBARRIER_PARALLEL:%.*]] 255 // CHECK2: .execute.parallel: 256 // CHECK2-NEXT: [[TMP4:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1:[0-9]+]]) 257 // CHECK2-NEXT: [[TMP5:%.*]] = load i8*, i8** [[WORK_FN]], align 4 258 // CHECK2-NEXT: [[WORK_MATCH:%.*]] = icmp eq i8* [[TMP5]], bitcast (void (i16, i32)* @__omp_outlined___wrapper to i8*) 259 // CHECK2-NEXT: br i1 [[WORK_MATCH]], label [[DOTEXECUTE_FN:%.*]], label [[DOTCHECK_NEXT:%.*]] 260 // CHECK2: .execute.fn: 261 // CHECK2-NEXT: call void @__omp_outlined___wrapper(i16 0, i32 [[TMP4]]) #[[ATTR4:[0-9]+]] 262 // CHECK2-NEXT: br label [[DOTTERMINATE_PARALLEL:%.*]] 263 // CHECK2: .check.next: 264 // CHECK2-NEXT: [[TMP6:%.*]] = bitcast i8* [[TMP2]] to void (i16, i32)* 265 // CHECK2-NEXT: call void [[TMP6]](i16 0, i32 [[TMP4]]) 266 // CHECK2-NEXT: br label [[DOTTERMINATE_PARALLEL]] 267 // CHECK2: .terminate.parallel: 268 // CHECK2-NEXT: call void @__kmpc_kernel_end_parallel() 269 // CHECK2-NEXT: br label [[DOTBARRIER_PARALLEL]] 270 // CHECK2: .barrier.parallel: 271 // CHECK2-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0) 272 // CHECK2-NEXT: br label [[DOTAWAIT_WORK]] 273 // CHECK2: .exit: 274 // CHECK2-NEXT: ret void 275 // 276 // 277 // CHECK2-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l25 278 // CHECK2-SAME: (i32* nonnull align 4 dereferenceable(4) [[C:%.*]]) #[[ATTR1:[0-9]+]] { 279 // CHECK2-NEXT: entry: 280 // CHECK2-NEXT: [[C_ADDR:%.*]] = alloca i32*, align 4 281 // CHECK2-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [1 x i8*], align 4 282 // CHECK2-NEXT: store i32* [[C]], i32** [[C_ADDR]], align 4 283 // CHECK2-NEXT: [[TMP0:%.*]] = load i32*, i32** [[C_ADDR]], align 4 284 // CHECK2-NEXT: [[NVPTX_TID:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x() 285 // CHECK2-NEXT: [[NVPTX_NUM_THREADS:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() 286 // CHECK2-NEXT: [[NVPTX_WARP_SIZE:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize() 287 // CHECK2-NEXT: [[THREAD_LIMIT:%.*]] = sub nuw i32 [[NVPTX_NUM_THREADS]], [[NVPTX_WARP_SIZE]] 288 // CHECK2-NEXT: [[TMP1:%.*]] = icmp ult i32 [[NVPTX_TID]], [[THREAD_LIMIT]] 289 // CHECK2-NEXT: br i1 [[TMP1]], label [[DOTWORKER:%.*]], label [[DOTMASTERCHECK:%.*]] 290 // CHECK2: .worker: 291 // CHECK2-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l25_worker() #[[ATTR4]] 292 // CHECK2-NEXT: br label [[DOTEXIT:%.*]] 293 // CHECK2: .mastercheck: 294 // CHECK2-NEXT: [[NVPTX_TID1:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x() 295 // CHECK2-NEXT: [[NVPTX_NUM_THREADS2:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() 296 // CHECK2-NEXT: [[NVPTX_WARP_SIZE3:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize() 297 // CHECK2-NEXT: [[TMP2:%.*]] = sub nuw i32 [[NVPTX_WARP_SIZE3]], 1 298 // CHECK2-NEXT: [[TMP3:%.*]] = sub nuw i32 [[NVPTX_NUM_THREADS2]], 1 299 // CHECK2-NEXT: [[TMP4:%.*]] = xor i32 [[TMP2]], -1 300 // CHECK2-NEXT: [[MASTER_TID:%.*]] = and i32 [[TMP3]], [[TMP4]] 301 // CHECK2-NEXT: [[TMP5:%.*]] = icmp eq i32 [[NVPTX_TID1]], [[MASTER_TID]] 302 // CHECK2-NEXT: br i1 [[TMP5]], label [[DOTMASTER:%.*]], label [[DOTEXIT]] 303 // CHECK2: .master: 304 // CHECK2-NEXT: [[NVPTX_NUM_THREADS4:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() 305 // CHECK2-NEXT: [[NVPTX_WARP_SIZE5:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize() 306 // CHECK2-NEXT: [[THREAD_LIMIT6:%.*]] = sub nuw i32 [[NVPTX_NUM_THREADS4]], [[NVPTX_WARP_SIZE5]] 307 // CHECK2-NEXT: call void @__kmpc_kernel_init(i32 [[THREAD_LIMIT6]], i16 1) 308 // CHECK2-NEXT: call void @__kmpc_data_sharing_init_stack() 309 // CHECK2-NEXT: [[TMP6:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1]]) 310 // CHECK2-NEXT: call void @_Z3usePi(i32* [[TMP0]]) #[[ATTR7:[0-9]+]] 311 // CHECK2-NEXT: call void @__kmpc_push_num_threads(%struct.ident_t* @[[GLOB1]], i32 [[TMP6]], i32 2) 312 // CHECK2-NEXT: [[TMP7:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[CAPTURED_VARS_ADDRS]], i32 0, i32 0 313 // CHECK2-NEXT: [[TMP8:%.*]] = bitcast i32* [[TMP0]] to i8* 314 // CHECK2-NEXT: store i8* [[TMP8]], i8** [[TMP7]], align 4 315 // CHECK2-NEXT: [[TMP9:%.*]] = bitcast [1 x i8*]* [[CAPTURED_VARS_ADDRS]] to i8** 316 // CHECK2-NEXT: call void @__kmpc_parallel_51(%struct.ident_t* @[[GLOB1]], i32 [[TMP6]], i32 1, i32 -1, i32 -1, i8* bitcast (void (i32*, i32*, i32*)* @__omp_outlined__ to i8*), i8* bitcast (void (i16, i32)* @__omp_outlined___wrapper to i8*), i8** [[TMP9]], i32 1) 317 // CHECK2-NEXT: br label [[DOTTERMINATION_NOTIFIER:%.*]] 318 // CHECK2: .termination.notifier: 319 // CHECK2-NEXT: call void @__kmpc_kernel_deinit(i16 1) 320 // CHECK2-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0) 321 // CHECK2-NEXT: br label [[DOTEXIT]] 322 // CHECK2: .exit: 323 // CHECK2-NEXT: ret void 324 // 325 // 326 // CHECK2-LABEL: define {{[^@]+}}@_Z3usePi 327 // CHECK2-SAME: (i32* [[C:%.*]]) #[[ATTR3:[0-9]+]] { 328 // CHECK2-NEXT: entry: 329 // CHECK2-NEXT: [[C_ADDR:%.*]] = alloca i32*, align 4 330 // CHECK2-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [1 x i8*], align 4 331 // CHECK2-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB2:[0-9]+]]) 332 // CHECK2-NEXT: store i32* [[C]], i32** [[C_ADDR]], align 4 333 // CHECK2-NEXT: call void @__kmpc_push_num_threads(%struct.ident_t* @[[GLOB2]], i32 [[TMP0]], i32 2) 334 // CHECK2-NEXT: [[TMP1:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[CAPTURED_VARS_ADDRS]], i32 0, i32 0 335 // CHECK2-NEXT: [[TMP2:%.*]] = bitcast i32** [[C_ADDR]] to i8* 336 // CHECK2-NEXT: store i8* [[TMP2]], i8** [[TMP1]], align 4 337 // CHECK2-NEXT: [[TMP3:%.*]] = bitcast [1 x i8*]* [[CAPTURED_VARS_ADDRS]] to i8** 338 // CHECK2-NEXT: call void @__kmpc_parallel_51(%struct.ident_t* @[[GLOB2]], i32 [[TMP0]], i32 1, i32 -1, i32 -1, i8* bitcast (void (i32*, i32*, i32**)* @__omp_outlined__1 to i8*), i8* bitcast (void (i16, i32)* @__omp_outlined__1_wrapper to i8*), i8** [[TMP3]], i32 1) 339 // CHECK2-NEXT: ret void 340 // 341 // 342 // CHECK2-LABEL: define {{[^@]+}}@__omp_outlined__ 343 // CHECK2-SAME: (i32* noalias [[DOTGLOBAL_TID_:%.*]], i32* noalias [[DOTBOUND_TID_:%.*]], i32* nonnull align 4 dereferenceable(4) [[C:%.*]]) #[[ATTR1]] { 344 // CHECK2-NEXT: entry: 345 // CHECK2-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 4 346 // CHECK2-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 4 347 // CHECK2-NEXT: [[C_ADDR:%.*]] = alloca i32*, align 4 348 // CHECK2-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 4 349 // CHECK2-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 4 350 // CHECK2-NEXT: store i32* [[C]], i32** [[C_ADDR]], align 4 351 // CHECK2-NEXT: [[TMP0:%.*]] = load i32*, i32** [[C_ADDR]], align 4 352 // CHECK2-NEXT: call void @_Z3usePi(i32* [[TMP0]]) #[[ATTR7]] 353 // CHECK2-NEXT: ret void 354 // 355 // 356 // CHECK2-LABEL: define {{[^@]+}}@__omp_outlined___wrapper 357 // CHECK2-SAME: (i16 zeroext [[TMP0:%.*]], i32 [[TMP1:%.*]]) #[[ATTR0]] { 358 // CHECK2-NEXT: entry: 359 // CHECK2-NEXT: [[DOTADDR:%.*]] = alloca i16, align 2 360 // CHECK2-NEXT: [[DOTADDR1:%.*]] = alloca i32, align 4 361 // CHECK2-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4 362 // CHECK2-NEXT: [[GLOBAL_ARGS:%.*]] = alloca i8**, align 4 363 // CHECK2-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4 364 // CHECK2-NEXT: store i16 [[TMP0]], i16* [[DOTADDR]], align 2 365 // CHECK2-NEXT: store i32 [[TMP1]], i32* [[DOTADDR1]], align 4 366 // CHECK2-NEXT: call void @__kmpc_get_shared_variables(i8*** [[GLOBAL_ARGS]]) 367 // CHECK2-NEXT: [[TMP2:%.*]] = load i8**, i8*** [[GLOBAL_ARGS]], align 4 368 // CHECK2-NEXT: [[TMP3:%.*]] = getelementptr inbounds i8*, i8** [[TMP2]], i32 0 369 // CHECK2-NEXT: [[TMP4:%.*]] = bitcast i8** [[TMP3]] to i32** 370 // CHECK2-NEXT: [[TMP5:%.*]] = load i32*, i32** [[TMP4]], align 4 371 // CHECK2-NEXT: call void @__omp_outlined__(i32* [[DOTADDR1]], i32* [[DOTZERO_ADDR]], i32* [[TMP5]]) #[[ATTR4]] 372 // CHECK2-NEXT: ret void 373 // 374 // 375 // CHECK2-LABEL: define {{[^@]+}}@__omp_outlined__1 376 // CHECK2-SAME: (i32* noalias [[DOTGLOBAL_TID_:%.*]], i32* noalias [[DOTBOUND_TID_:%.*]], i32** nonnull align 4 dereferenceable(4) [[C:%.*]]) #[[ATTR1]] { 377 // CHECK2-NEXT: entry: 378 // CHECK2-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 4 379 // CHECK2-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 4 380 // CHECK2-NEXT: [[C_ADDR:%.*]] = alloca i32**, align 4 381 // CHECK2-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 4 382 // CHECK2-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 4 383 // CHECK2-NEXT: store i32** [[C]], i32*** [[C_ADDR]], align 4 384 // CHECK2-NEXT: [[TMP0:%.*]] = load i32**, i32*** [[C_ADDR]], align 4 385 // CHECK2-NEXT: [[TMP1:%.*]] = load i32*, i32** [[TMP0]], align 4 386 // CHECK2-NEXT: call void @_Z4workPi(i32* [[TMP1]]) #[[ATTR7]] 387 // CHECK2-NEXT: ret void 388 // 389 // 390 // CHECK2-LABEL: define {{[^@]+}}@_Z4workPi 391 // CHECK2-SAME: (i32* [[C:%.*]]) #[[ATTR3]] { 392 // CHECK2-NEXT: entry: 393 // CHECK2-NEXT: [[C_ADDR:%.*]] = alloca i32*, align 4 394 // CHECK2-NEXT: [[ATOMIC_TEMP:%.*]] = alloca i32, align 4 395 // CHECK2-NEXT: [[ATOMIC_TEMP1:%.*]] = alloca i32, align 4 396 // CHECK2-NEXT: store i32* [[C]], i32** [[C_ADDR]], align 4 397 // CHECK2-NEXT: [[TMP0:%.*]] = load i32*, i32** [[C_ADDR]], align 4 398 // CHECK2-NEXT: [[TMP1:%.*]] = bitcast i32* [[TMP0]] to i8* 399 // CHECK2-NEXT: [[TMP2:%.*]] = bitcast i32* [[ATOMIC_TEMP]] to i8* 400 // CHECK2-NEXT: call void @__atomic_load(i32 4, i8* [[TMP1]], i8* [[TMP2]], i32 0) #[[ATTR7]] 401 // CHECK2-NEXT: br label [[ATOMIC_CONT:%.*]] 402 // CHECK2: atomic_cont: 403 // CHECK2-NEXT: [[TMP3:%.*]] = load i32, i32* [[ATOMIC_TEMP]], align 4 404 // CHECK2-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP3]], 1 405 // CHECK2-NEXT: store i32 [[ADD]], i32* [[ATOMIC_TEMP1]], align 4 406 // CHECK2-NEXT: [[TMP4:%.*]] = bitcast i32* [[TMP0]] to i8* 407 // CHECK2-NEXT: [[TMP5:%.*]] = bitcast i32* [[ATOMIC_TEMP]] to i8* 408 // CHECK2-NEXT: [[TMP6:%.*]] = bitcast i32* [[ATOMIC_TEMP1]] to i8* 409 // CHECK2-NEXT: [[CALL:%.*]] = call zeroext i1 @__atomic_compare_exchange(i32 4, i8* [[TMP4]], i8* [[TMP5]], i8* [[TMP6]], i32 0, i32 0) #[[ATTR7]] 410 // CHECK2-NEXT: br i1 [[CALL]], label [[ATOMIC_EXIT:%.*]], label [[ATOMIC_CONT]] 411 // CHECK2: atomic_exit: 412 // CHECK2-NEXT: ret void 413 // 414 // 415 // CHECK2-LABEL: define {{[^@]+}}@__omp_outlined__1_wrapper 416 // CHECK2-SAME: (i16 zeroext [[TMP0:%.*]], i32 [[TMP1:%.*]]) #[[ATTR0]] { 417 // CHECK2-NEXT: entry: 418 // CHECK2-NEXT: [[DOTADDR:%.*]] = alloca i16, align 2 419 // CHECK2-NEXT: [[DOTADDR1:%.*]] = alloca i32, align 4 420 // CHECK2-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4 421 // CHECK2-NEXT: [[GLOBAL_ARGS:%.*]] = alloca i8**, align 4 422 // CHECK2-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4 423 // CHECK2-NEXT: store i16 [[TMP0]], i16* [[DOTADDR]], align 2 424 // CHECK2-NEXT: store i32 [[TMP1]], i32* [[DOTADDR1]], align 4 425 // CHECK2-NEXT: call void @__kmpc_get_shared_variables(i8*** [[GLOBAL_ARGS]]) 426 // CHECK2-NEXT: [[TMP2:%.*]] = load i8**, i8*** [[GLOBAL_ARGS]], align 4 427 // CHECK2-NEXT: [[TMP3:%.*]] = getelementptr inbounds i8*, i8** [[TMP2]], i32 0 428 // CHECK2-NEXT: [[TMP4:%.*]] = bitcast i8** [[TMP3]] to i32*** 429 // CHECK2-NEXT: [[TMP5:%.*]] = load i32**, i32*** [[TMP4]], align 4 430 // CHECK2-NEXT: call void @__omp_outlined__1(i32* [[DOTADDR1]], i32* [[DOTZERO_ADDR]], i32** [[TMP5]]) #[[ATTR4]] 431 // CHECK2-NEXT: ret void 432 // 433 // 434 // CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l25_worker 435 // CHECK3-SAME: () #[[ATTR0:[0-9]+]] { 436 // CHECK3-NEXT: entry: 437 // CHECK3-NEXT: [[WORK_FN:%.*]] = alloca i8*, align 4 438 // CHECK3-NEXT: [[EXEC_STATUS:%.*]] = alloca i8, align 1 439 // CHECK3-NEXT: store i8* null, i8** [[WORK_FN]], align 4 440 // CHECK3-NEXT: store i8 0, i8* [[EXEC_STATUS]], align 1 441 // CHECK3-NEXT: br label [[DOTAWAIT_WORK:%.*]] 442 // CHECK3: .await.work: 443 // CHECK3-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0) 444 // CHECK3-NEXT: [[TMP0:%.*]] = call i1 @__kmpc_kernel_parallel(i8** [[WORK_FN]]) 445 // CHECK3-NEXT: [[TMP1:%.*]] = zext i1 [[TMP0]] to i8 446 // CHECK3-NEXT: store i8 [[TMP1]], i8* [[EXEC_STATUS]], align 1 447 // CHECK3-NEXT: [[TMP2:%.*]] = load i8*, i8** [[WORK_FN]], align 4 448 // CHECK3-NEXT: [[SHOULD_TERMINATE:%.*]] = icmp eq i8* [[TMP2]], null 449 // CHECK3-NEXT: br i1 [[SHOULD_TERMINATE]], label [[DOTEXIT:%.*]], label [[DOTSELECT_WORKERS:%.*]] 450 // CHECK3: .select.workers: 451 // CHECK3-NEXT: [[TMP3:%.*]] = load i8, i8* [[EXEC_STATUS]], align 1 452 // CHECK3-NEXT: [[IS_ACTIVE:%.*]] = icmp ne i8 [[TMP3]], 0 453 // CHECK3-NEXT: br i1 [[IS_ACTIVE]], label [[DOTEXECUTE_PARALLEL:%.*]], label [[DOTBARRIER_PARALLEL:%.*]] 454 // CHECK3: .execute.parallel: 455 // CHECK3-NEXT: [[TMP4:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1:[0-9]+]]) 456 // CHECK3-NEXT: [[TMP5:%.*]] = load i8*, i8** [[WORK_FN]], align 4 457 // CHECK3-NEXT: [[WORK_MATCH:%.*]] = icmp eq i8* [[TMP5]], bitcast (void (i16, i32)* @__omp_outlined___wrapper to i8*) 458 // CHECK3-NEXT: br i1 [[WORK_MATCH]], label [[DOTEXECUTE_FN:%.*]], label [[DOTCHECK_NEXT:%.*]] 459 // CHECK3: .execute.fn: 460 // CHECK3-NEXT: call void @__omp_outlined___wrapper(i16 0, i32 [[TMP4]]) #[[ATTR4:[0-9]+]] 461 // CHECK3-NEXT: br label [[DOTTERMINATE_PARALLEL:%.*]] 462 // CHECK3: .check.next: 463 // CHECK3-NEXT: [[TMP6:%.*]] = bitcast i8* [[TMP2]] to void (i16, i32)* 464 // CHECK3-NEXT: call void [[TMP6]](i16 0, i32 [[TMP4]]) 465 // CHECK3-NEXT: br label [[DOTTERMINATE_PARALLEL]] 466 // CHECK3: .terminate.parallel: 467 // CHECK3-NEXT: call void @__kmpc_kernel_end_parallel() 468 // CHECK3-NEXT: br label [[DOTBARRIER_PARALLEL]] 469 // CHECK3: .barrier.parallel: 470 // CHECK3-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0) 471 // CHECK3-NEXT: br label [[DOTAWAIT_WORK]] 472 // CHECK3: .exit: 473 // CHECK3-NEXT: ret void 474 // 475 // 476 // CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l25 477 // CHECK3-SAME: (i32* nonnull align 4 dereferenceable(4) [[C:%.*]]) #[[ATTR1:[0-9]+]] { 478 // CHECK3-NEXT: entry: 479 // CHECK3-NEXT: [[C_ADDR:%.*]] = alloca i32*, align 4 480 // CHECK3-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [1 x i8*], align 4 481 // CHECK3-NEXT: store i32* [[C]], i32** [[C_ADDR]], align 4 482 // CHECK3-NEXT: [[TMP0:%.*]] = load i32*, i32** [[C_ADDR]], align 4 483 // CHECK3-NEXT: [[NVPTX_TID:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x() 484 // CHECK3-NEXT: [[NVPTX_NUM_THREADS:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() 485 // CHECK3-NEXT: [[NVPTX_WARP_SIZE:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize() 486 // CHECK3-NEXT: [[THREAD_LIMIT:%.*]] = sub nuw i32 [[NVPTX_NUM_THREADS]], [[NVPTX_WARP_SIZE]] 487 // CHECK3-NEXT: [[TMP1:%.*]] = icmp ult i32 [[NVPTX_TID]], [[THREAD_LIMIT]] 488 // CHECK3-NEXT: br i1 [[TMP1]], label [[DOTWORKER:%.*]], label [[DOTMASTERCHECK:%.*]] 489 // CHECK3: .worker: 490 // CHECK3-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l25_worker() #[[ATTR4]] 491 // CHECK3-NEXT: br label [[DOTEXIT:%.*]] 492 // CHECK3: .mastercheck: 493 // CHECK3-NEXT: [[NVPTX_TID1:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x() 494 // CHECK3-NEXT: [[NVPTX_NUM_THREADS2:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() 495 // CHECK3-NEXT: [[NVPTX_WARP_SIZE3:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize() 496 // CHECK3-NEXT: [[TMP2:%.*]] = sub nuw i32 [[NVPTX_WARP_SIZE3]], 1 497 // CHECK3-NEXT: [[TMP3:%.*]] = sub nuw i32 [[NVPTX_NUM_THREADS2]], 1 498 // CHECK3-NEXT: [[TMP4:%.*]] = xor i32 [[TMP2]], -1 499 // CHECK3-NEXT: [[MASTER_TID:%.*]] = and i32 [[TMP3]], [[TMP4]] 500 // CHECK3-NEXT: [[TMP5:%.*]] = icmp eq i32 [[NVPTX_TID1]], [[MASTER_TID]] 501 // CHECK3-NEXT: br i1 [[TMP5]], label [[DOTMASTER:%.*]], label [[DOTEXIT]] 502 // CHECK3: .master: 503 // CHECK3-NEXT: [[NVPTX_NUM_THREADS4:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() 504 // CHECK3-NEXT: [[NVPTX_WARP_SIZE5:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize() 505 // CHECK3-NEXT: [[THREAD_LIMIT6:%.*]] = sub nuw i32 [[NVPTX_NUM_THREADS4]], [[NVPTX_WARP_SIZE5]] 506 // CHECK3-NEXT: call void @__kmpc_kernel_init(i32 [[THREAD_LIMIT6]], i16 1) 507 // CHECK3-NEXT: call void @__kmpc_data_sharing_init_stack() 508 // CHECK3-NEXT: [[TMP6:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1]]) 509 // CHECK3-NEXT: call void @_Z3usePi(i32* [[TMP0]]) #[[ATTR7:[0-9]+]] 510 // CHECK3-NEXT: call void @__kmpc_push_num_threads(%struct.ident_t* @[[GLOB1]], i32 [[TMP6]], i32 2) 511 // CHECK3-NEXT: [[TMP7:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[CAPTURED_VARS_ADDRS]], i32 0, i32 0 512 // CHECK3-NEXT: [[TMP8:%.*]] = bitcast i32* [[TMP0]] to i8* 513 // CHECK3-NEXT: store i8* [[TMP8]], i8** [[TMP7]], align 4 514 // CHECK3-NEXT: [[TMP9:%.*]] = bitcast [1 x i8*]* [[CAPTURED_VARS_ADDRS]] to i8** 515 // CHECK3-NEXT: call void @__kmpc_parallel_51(%struct.ident_t* @[[GLOB1]], i32 [[TMP6]], i32 1, i32 -1, i32 -1, i8* bitcast (void (i32*, i32*, i32*)* @__omp_outlined__ to i8*), i8* bitcast (void (i16, i32)* @__omp_outlined___wrapper to i8*), i8** [[TMP9]], i32 1) 516 // CHECK3-NEXT: br label [[DOTTERMINATION_NOTIFIER:%.*]] 517 // CHECK3: .termination.notifier: 518 // CHECK3-NEXT: call void @__kmpc_kernel_deinit(i16 1) 519 // CHECK3-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0) 520 // CHECK3-NEXT: br label [[DOTEXIT]] 521 // CHECK3: .exit: 522 // CHECK3-NEXT: ret void 523 // 524 // 525 // CHECK3-LABEL: define {{[^@]+}}@_Z3usePi 526 // CHECK3-SAME: (i32* [[C:%.*]]) #[[ATTR3:[0-9]+]] { 527 // CHECK3-NEXT: entry: 528 // CHECK3-NEXT: [[C_ADDR:%.*]] = alloca i32*, align 4 529 // CHECK3-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [1 x i8*], align 4 530 // CHECK3-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB2:[0-9]+]]) 531 // CHECK3-NEXT: store i32* [[C]], i32** [[C_ADDR]], align 4 532 // CHECK3-NEXT: call void @__kmpc_push_num_threads(%struct.ident_t* @[[GLOB2]], i32 [[TMP0]], i32 2) 533 // CHECK3-NEXT: [[TMP1:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[CAPTURED_VARS_ADDRS]], i32 0, i32 0 534 // CHECK3-NEXT: [[TMP2:%.*]] = bitcast i32** [[C_ADDR]] to i8* 535 // CHECK3-NEXT: store i8* [[TMP2]], i8** [[TMP1]], align 4 536 // CHECK3-NEXT: [[TMP3:%.*]] = bitcast [1 x i8*]* [[CAPTURED_VARS_ADDRS]] to i8** 537 // CHECK3-NEXT: call void @__kmpc_parallel_51(%struct.ident_t* @[[GLOB2]], i32 [[TMP0]], i32 1, i32 -1, i32 -1, i8* bitcast (void (i32*, i32*, i32**)* @__omp_outlined__1 to i8*), i8* bitcast (void (i16, i32)* @__omp_outlined__1_wrapper to i8*), i8** [[TMP3]], i32 1) 538 // CHECK3-NEXT: ret void 539 // 540 // 541 // CHECK3-LABEL: define {{[^@]+}}@__omp_outlined__ 542 // CHECK3-SAME: (i32* noalias [[DOTGLOBAL_TID_:%.*]], i32* noalias [[DOTBOUND_TID_:%.*]], i32* nonnull align 4 dereferenceable(4) [[C:%.*]]) #[[ATTR1]] { 543 // CHECK3-NEXT: entry: 544 // CHECK3-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 4 545 // CHECK3-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 4 546 // CHECK3-NEXT: [[C_ADDR:%.*]] = alloca i32*, align 4 547 // CHECK3-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 4 548 // CHECK3-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 4 549 // CHECK3-NEXT: store i32* [[C]], i32** [[C_ADDR]], align 4 550 // CHECK3-NEXT: [[TMP0:%.*]] = load i32*, i32** [[C_ADDR]], align 4 551 // CHECK3-NEXT: call void @_Z3usePi(i32* [[TMP0]]) #[[ATTR7]] 552 // CHECK3-NEXT: ret void 553 // 554 // 555 // CHECK3-LABEL: define {{[^@]+}}@__omp_outlined___wrapper 556 // CHECK3-SAME: (i16 zeroext [[TMP0:%.*]], i32 [[TMP1:%.*]]) #[[ATTR0]] { 557 // CHECK3-NEXT: entry: 558 // CHECK3-NEXT: [[DOTADDR:%.*]] = alloca i16, align 2 559 // CHECK3-NEXT: [[DOTADDR1:%.*]] = alloca i32, align 4 560 // CHECK3-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4 561 // CHECK3-NEXT: [[GLOBAL_ARGS:%.*]] = alloca i8**, align 4 562 // CHECK3-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4 563 // CHECK3-NEXT: store i16 [[TMP0]], i16* [[DOTADDR]], align 2 564 // CHECK3-NEXT: store i32 [[TMP1]], i32* [[DOTADDR1]], align 4 565 // CHECK3-NEXT: call void @__kmpc_get_shared_variables(i8*** [[GLOBAL_ARGS]]) 566 // CHECK3-NEXT: [[TMP2:%.*]] = load i8**, i8*** [[GLOBAL_ARGS]], align 4 567 // CHECK3-NEXT: [[TMP3:%.*]] = getelementptr inbounds i8*, i8** [[TMP2]], i32 0 568 // CHECK3-NEXT: [[TMP4:%.*]] = bitcast i8** [[TMP3]] to i32** 569 // CHECK3-NEXT: [[TMP5:%.*]] = load i32*, i32** [[TMP4]], align 4 570 // CHECK3-NEXT: call void @__omp_outlined__(i32* [[DOTADDR1]], i32* [[DOTZERO_ADDR]], i32* [[TMP5]]) #[[ATTR4]] 571 // CHECK3-NEXT: ret void 572 // 573 // 574 // CHECK3-LABEL: define {{[^@]+}}@__omp_outlined__1 575 // CHECK3-SAME: (i32* noalias [[DOTGLOBAL_TID_:%.*]], i32* noalias [[DOTBOUND_TID_:%.*]], i32** nonnull align 4 dereferenceable(4) [[C:%.*]]) #[[ATTR1]] { 576 // CHECK3-NEXT: entry: 577 // CHECK3-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 4 578 // CHECK3-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 4 579 // CHECK3-NEXT: [[C_ADDR:%.*]] = alloca i32**, align 4 580 // CHECK3-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 4 581 // CHECK3-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 4 582 // CHECK3-NEXT: store i32** [[C]], i32*** [[C_ADDR]], align 4 583 // CHECK3-NEXT: [[TMP0:%.*]] = load i32**, i32*** [[C_ADDR]], align 4 584 // CHECK3-NEXT: [[TMP1:%.*]] = load i32*, i32** [[TMP0]], align 4 585 // CHECK3-NEXT: call void @_Z4workPi(i32* [[TMP1]]) #[[ATTR7]] 586 // CHECK3-NEXT: ret void 587 // 588 // 589 // CHECK3-LABEL: define {{[^@]+}}@_Z4workPi 590 // CHECK3-SAME: (i32* [[C:%.*]]) #[[ATTR3]] { 591 // CHECK3-NEXT: entry: 592 // CHECK3-NEXT: [[C_ADDR:%.*]] = alloca i32*, align 4 593 // CHECK3-NEXT: [[ATOMIC_TEMP:%.*]] = alloca i32, align 4 594 // CHECK3-NEXT: [[ATOMIC_TEMP1:%.*]] = alloca i32, align 4 595 // CHECK3-NEXT: store i32* [[C]], i32** [[C_ADDR]], align 4 596 // CHECK3-NEXT: [[TMP0:%.*]] = load i32*, i32** [[C_ADDR]], align 4 597 // CHECK3-NEXT: [[TMP1:%.*]] = bitcast i32* [[TMP0]] to i8* 598 // CHECK3-NEXT: [[TMP2:%.*]] = bitcast i32* [[ATOMIC_TEMP]] to i8* 599 // CHECK3-NEXT: call void @__atomic_load(i32 4, i8* [[TMP1]], i8* [[TMP2]], i32 0) #[[ATTR7]] 600 // CHECK3-NEXT: br label [[ATOMIC_CONT:%.*]] 601 // CHECK3: atomic_cont: 602 // CHECK3-NEXT: [[TMP3:%.*]] = load i32, i32* [[ATOMIC_TEMP]], align 4 603 // CHECK3-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP3]], 1 604 // CHECK3-NEXT: store i32 [[ADD]], i32* [[ATOMIC_TEMP1]], align 4 605 // CHECK3-NEXT: [[TMP4:%.*]] = bitcast i32* [[TMP0]] to i8* 606 // CHECK3-NEXT: [[TMP5:%.*]] = bitcast i32* [[ATOMIC_TEMP]] to i8* 607 // CHECK3-NEXT: [[TMP6:%.*]] = bitcast i32* [[ATOMIC_TEMP1]] to i8* 608 // CHECK3-NEXT: [[CALL:%.*]] = call zeroext i1 @__atomic_compare_exchange(i32 4, i8* [[TMP4]], i8* [[TMP5]], i8* [[TMP6]], i32 0, i32 0) #[[ATTR7]] 609 // CHECK3-NEXT: br i1 [[CALL]], label [[ATOMIC_EXIT:%.*]], label [[ATOMIC_CONT]] 610 // CHECK3: atomic_exit: 611 // CHECK3-NEXT: ret void 612 // 613 // 614 // CHECK3-LABEL: define {{[^@]+}}@__omp_outlined__1_wrapper 615 // CHECK3-SAME: (i16 zeroext [[TMP0:%.*]], i32 [[TMP1:%.*]]) #[[ATTR0]] { 616 // CHECK3-NEXT: entry: 617 // CHECK3-NEXT: [[DOTADDR:%.*]] = alloca i16, align 2 618 // CHECK3-NEXT: [[DOTADDR1:%.*]] = alloca i32, align 4 619 // CHECK3-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4 620 // CHECK3-NEXT: [[GLOBAL_ARGS:%.*]] = alloca i8**, align 4 621 // CHECK3-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4 622 // CHECK3-NEXT: store i16 [[TMP0]], i16* [[DOTADDR]], align 2 623 // CHECK3-NEXT: store i32 [[TMP1]], i32* [[DOTADDR1]], align 4 624 // CHECK3-NEXT: call void @__kmpc_get_shared_variables(i8*** [[GLOBAL_ARGS]]) 625 // CHECK3-NEXT: [[TMP2:%.*]] = load i8**, i8*** [[GLOBAL_ARGS]], align 4 626 // CHECK3-NEXT: [[TMP3:%.*]] = getelementptr inbounds i8*, i8** [[TMP2]], i32 0 627 // CHECK3-NEXT: [[TMP4:%.*]] = bitcast i8** [[TMP3]] to i32*** 628 // CHECK3-NEXT: [[TMP5:%.*]] = load i32**, i32*** [[TMP4]], align 4 629 // CHECK3-NEXT: call void @__omp_outlined__1(i32* [[DOTADDR1]], i32* [[DOTZERO_ADDR]], i32** [[TMP5]]) #[[ATTR4]] 630 // CHECK3-NEXT: ret void 631 // 632