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 -DCK1 -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 -DCK1 -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 -DCK1 -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 -DCK1 -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 // expected-no-diagnostics 8 #ifndef HEADER 9 #define HEADER 10 11 #ifdef CK1 12 13 template <typename T> 14 int tmain(T argc) { 15 #pragma omp target 16 #pragma omp teams 17 argc = 0; 18 return 0; 19 } 20 21 22 int main (int argc, char **argv) { 23 #pragma omp target 24 #pragma omp teams 25 { 26 argc = 0; 27 } 28 return tmain(argv); 29 } 30 31 32 // only nvptx side: do not outline teams region and do not call fork_teams 33 34 35 // target region in template 36 37 38 39 #endif // CK1 40 41 // Test target codegen - host bc file has to be created first. 42 // RUN: %clang_cc1 -DCK2 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc 43 // RUN: %clang_cc1 -DCK2 -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=CHECK3 44 // RUN: %clang_cc1 -DCK2 -verify -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm-bc %s -o %t-x86-host.bc 45 // RUN: %clang_cc1 -DCK2 -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=CHECK4 46 // expected-no-diagnostics 47 #ifdef CK2 48 49 template <typename T> 50 int tmain(T argc) { 51 int a = 10; 52 int b = 5; 53 #pragma omp target 54 #pragma omp teams num_teams(a) thread_limit(b) 55 { 56 argc = 0; 57 } 58 return 0; 59 } 60 61 int main (int argc, char **argv) { 62 int a = 20; 63 int b = 5; 64 #pragma omp target 65 #pragma omp teams num_teams(a) thread_limit(b) 66 { 67 argc = 0; 68 } 69 return tmain(argv); 70 } 71 72 73 74 75 76 77 #endif // CK2 78 #endif 79 // CHECK5-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l68_worker 80 // CHECK5-SAME: () #[[ATTR0:[0-9]+]] { 81 // CHECK5-NEXT: entry: 82 // CHECK5-NEXT: [[WORK_FN:%.*]] = alloca i8*, align 8 83 // CHECK5-NEXT: [[EXEC_STATUS:%.*]] = alloca i8, align 1 84 // CHECK5-NEXT: store i8* null, i8** [[WORK_FN]], align 8 85 // CHECK5-NEXT: store i8 0, i8* [[EXEC_STATUS]], align 1 86 // CHECK5-NEXT: br label [[DOTAWAIT_WORK:%.*]] 87 // CHECK5: .await.work: 88 // CHECK5-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0) 89 // CHECK5-NEXT: [[TMP0:%.*]] = call i1 @__kmpc_kernel_parallel(i8** [[WORK_FN]]) 90 // CHECK5-NEXT: [[TMP1:%.*]] = zext i1 [[TMP0]] to i8 91 // CHECK5-NEXT: store i8 [[TMP1]], i8* [[EXEC_STATUS]], align 1 92 // CHECK5-NEXT: [[TMP2:%.*]] = load i8*, i8** [[WORK_FN]], align 8 93 // CHECK5-NEXT: [[SHOULD_TERMINATE:%.*]] = icmp eq i8* [[TMP2]], null 94 // CHECK5-NEXT: br i1 [[SHOULD_TERMINATE]], label [[DOTEXIT:%.*]], label [[DOTSELECT_WORKERS:%.*]] 95 // CHECK5: .select.workers: 96 // CHECK5-NEXT: [[TMP3:%.*]] = load i8, i8* [[EXEC_STATUS]], align 1 97 // CHECK5-NEXT: [[IS_ACTIVE:%.*]] = icmp ne i8 [[TMP3]], 0 98 // CHECK5-NEXT: br i1 [[IS_ACTIVE]], label [[DOTEXECUTE_PARALLEL:%.*]], label [[DOTBARRIER_PARALLEL:%.*]] 99 // CHECK5: .execute.parallel: 100 // CHECK5-NEXT: [[TMP4:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1:[0-9]+]]) 101 // CHECK5-NEXT: [[TMP5:%.*]] = bitcast i8* [[TMP2]] to void (i16, i32)* 102 // CHECK5-NEXT: call void [[TMP5]](i16 0, i32 [[TMP4]]) 103 // CHECK5-NEXT: br label [[DOTTERMINATE_PARALLEL:%.*]] 104 // CHECK5: .terminate.parallel: 105 // CHECK5-NEXT: call void @__kmpc_kernel_end_parallel() 106 // CHECK5-NEXT: br label [[DOTBARRIER_PARALLEL]] 107 // CHECK5: .barrier.parallel: 108 // CHECK5-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0) 109 // CHECK5-NEXT: br label [[DOTAWAIT_WORK]] 110 // CHECK5: .exit: 111 // CHECK5-NEXT: ret void 112 // CHECK5-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l68 113 // CHECK5-SAME: (i64 [[A:%.*]], i64 [[B:%.*]], i64 [[ARGC:%.*]]) #[[ATTR1:[0-9]+]] { 114 // CHECK5-NEXT: entry: 115 // CHECK5-NEXT: [[A_ADDR:%.*]] = alloca i64, align 8 116 // CHECK5-NEXT: [[B_ADDR:%.*]] = alloca i64, align 8 117 // CHECK5-NEXT: [[ARGC_ADDR:%.*]] = alloca i64, align 8 118 // CHECK5-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4 119 // CHECK5-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 4 120 // CHECK5-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4 121 // CHECK5-NEXT: store i64 [[A]], i64* [[A_ADDR]], align 8 122 // CHECK5-NEXT: store i64 [[B]], i64* [[B_ADDR]], align 8 123 // CHECK5-NEXT: store i64 [[ARGC]], i64* [[ARGC_ADDR]], align 8 124 // CHECK5-NEXT: [[CONV:%.*]] = bitcast i64* [[A_ADDR]] to i32* 125 // CHECK5-NEXT: [[CONV1:%.*]] = bitcast i64* [[B_ADDR]] to i32* 126 // CHECK5-NEXT: [[CONV2:%.*]] = bitcast i64* [[ARGC_ADDR]] to i32* 127 // CHECK5-NEXT: [[NVPTX_TID:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x() 128 // CHECK5-NEXT: [[NVPTX_NUM_THREADS:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() 129 // CHECK5-NEXT: [[NVPTX_WARP_SIZE:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize() 130 // CHECK5-NEXT: [[THREAD_LIMIT:%.*]] = sub nuw i32 [[NVPTX_NUM_THREADS]], [[NVPTX_WARP_SIZE]] 131 // CHECK5-NEXT: [[TMP0:%.*]] = icmp ult i32 [[NVPTX_TID]], [[THREAD_LIMIT]] 132 // CHECK5-NEXT: br i1 [[TMP0]], label [[DOTWORKER:%.*]], label [[DOTMASTERCHECK:%.*]] 133 // CHECK5: .worker: 134 // CHECK5-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l68_worker() #[[ATTR3:[0-9]+]] 135 // CHECK5-NEXT: br label [[DOTEXIT:%.*]] 136 // CHECK5: .mastercheck: 137 // CHECK5-NEXT: [[NVPTX_TID3:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x() 138 // CHECK5-NEXT: [[NVPTX_NUM_THREADS4:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() 139 // CHECK5-NEXT: [[NVPTX_WARP_SIZE5:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize() 140 // CHECK5-NEXT: [[TMP1:%.*]] = sub nuw i32 [[NVPTX_WARP_SIZE5]], 1 141 // CHECK5-NEXT: [[TMP2:%.*]] = sub nuw i32 [[NVPTX_NUM_THREADS4]], 1 142 // CHECK5-NEXT: [[TMP3:%.*]] = xor i32 [[TMP1]], -1 143 // CHECK5-NEXT: [[MASTER_TID:%.*]] = and i32 [[TMP2]], [[TMP3]] 144 // CHECK5-NEXT: [[TMP4:%.*]] = icmp eq i32 [[NVPTX_TID3]], [[MASTER_TID]] 145 // CHECK5-NEXT: br i1 [[TMP4]], label [[DOTMASTER:%.*]], label [[DOTEXIT]] 146 // CHECK5: .master: 147 // CHECK5-NEXT: [[NVPTX_NUM_THREADS6:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() 148 // CHECK5-NEXT: [[NVPTX_WARP_SIZE7:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize() 149 // CHECK5-NEXT: [[THREAD_LIMIT8:%.*]] = sub nuw i32 [[NVPTX_NUM_THREADS6]], [[NVPTX_WARP_SIZE7]] 150 // CHECK5-NEXT: call void @__kmpc_kernel_init(i32 [[THREAD_LIMIT8]], i16 1) 151 // CHECK5-NEXT: call void @__kmpc_data_sharing_init_stack() 152 // CHECK5-NEXT: [[TMP5:%.*]] = load i16, i16* @"_openmp_static_kernel$is_shared", align 2 153 // CHECK5-NEXT: [[TMP6:%.*]] = load i64, i64* @"_openmp_static_kernel$size", align 8 154 // CHECK5-NEXT: call void @__kmpc_get_team_static_memory(i16 0, i8* addrspacecast (i8 addrspace(3)* getelementptr inbounds (%"union._shared_openmp_static_memory_type_$_", %"union._shared_openmp_static_memory_type_$_" addrspace(3)* @"_openmp_shared_static_glob_rd_$_", i32 0, i32 0, i32 0) to i8*), i64 [[TMP6]], i16 [[TMP5]], i8** addrspacecast (i8* addrspace(3)* @"_openmp_kernel_static_glob_rd$ptr" to i8**)) 155 // CHECK5-NEXT: [[TMP7:%.*]] = load i8*, i8* addrspace(3)* @"_openmp_kernel_static_glob_rd$ptr", align 8 156 // CHECK5-NEXT: [[TMP8:%.*]] = getelementptr inbounds i8, i8* [[TMP7]], i64 0 157 // CHECK5-NEXT: [[TMP9:%.*]] = bitcast i8* [[TMP8]] to %struct._globalized_locals_ty* 158 // CHECK5-NEXT: [[TMP10:%.*]] = load i32, i32* [[CONV2]], align 8 159 // CHECK5-NEXT: [[ARGC9:%.*]] = getelementptr inbounds [[STRUCT__GLOBALIZED_LOCALS_TY:%.*]], %struct._globalized_locals_ty* [[TMP9]], i32 0, i32 0 160 // CHECK5-NEXT: store i32 [[TMP10]], i32* [[ARGC9]], align 4 161 // CHECK5-NEXT: [[TMP11:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1]]) 162 // CHECK5-NEXT: store i32 [[TMP11]], i32* [[DOTTHREADID_TEMP_]], align 4 163 // CHECK5-NEXT: call void @__omp_outlined__(i32* [[DOTTHREADID_TEMP_]], i32* [[DOTZERO_ADDR]], i32* [[ARGC9]]) #[[ATTR3]] 164 // CHECK5-NEXT: [[TMP12:%.*]] = load i16, i16* @"_openmp_static_kernel$is_shared", align 2 165 // CHECK5-NEXT: call void @__kmpc_restore_team_static_memory(i16 0, i16 [[TMP12]]) 166 // CHECK5-NEXT: br label [[DOTTERMINATION_NOTIFIER:%.*]] 167 // CHECK5: .termination.notifier: 168 // CHECK5-NEXT: call void @__kmpc_kernel_deinit(i16 1) 169 // CHECK5-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0) 170 // CHECK5-NEXT: br label [[DOTEXIT]] 171 // CHECK5: .exit: 172 // CHECK5-NEXT: ret void 173 // CHECK5-LABEL: define {{[^@]+}}@__omp_outlined__ 174 // CHECK5-SAME: (i32* noalias [[DOTGLOBAL_TID_:%.*]], i32* noalias [[DOTBOUND_TID_:%.*]], i32* nonnull align 4 dereferenceable(4) [[ARGC:%.*]]) #[[ATTR1]] { 175 // CHECK5-NEXT: entry: 176 // CHECK5-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8 177 // CHECK5-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8 178 // CHECK5-NEXT: [[ARGC_ADDR:%.*]] = alloca i32*, align 8 179 // CHECK5-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 8 180 // CHECK5-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 8 181 // CHECK5-NEXT: store i32* [[ARGC]], i32** [[ARGC_ADDR]], align 8 182 // CHECK5-NEXT: [[TMP0:%.*]] = load i32*, i32** [[ARGC_ADDR]], align 8 183 // CHECK5-NEXT: store i32 0, i32* [[TMP0]], align 4 184 // CHECK5-NEXT: ret void 185 // CHECK5-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIPPcEiT__l57_worker 186 // CHECK5-SAME: () #[[ATTR0]] { 187 // CHECK5-NEXT: entry: 188 // CHECK5-NEXT: [[WORK_FN:%.*]] = alloca i8*, align 8 189 // CHECK5-NEXT: [[EXEC_STATUS:%.*]] = alloca i8, align 1 190 // CHECK5-NEXT: store i8* null, i8** [[WORK_FN]], align 8 191 // CHECK5-NEXT: store i8 0, i8* [[EXEC_STATUS]], align 1 192 // CHECK5-NEXT: br label [[DOTAWAIT_WORK:%.*]] 193 // CHECK5: .await.work: 194 // CHECK5-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0) 195 // CHECK5-NEXT: [[TMP0:%.*]] = call i1 @__kmpc_kernel_parallel(i8** [[WORK_FN]]) 196 // CHECK5-NEXT: [[TMP1:%.*]] = zext i1 [[TMP0]] to i8 197 // CHECK5-NEXT: store i8 [[TMP1]], i8* [[EXEC_STATUS]], align 1 198 // CHECK5-NEXT: [[TMP2:%.*]] = load i8*, i8** [[WORK_FN]], align 8 199 // CHECK5-NEXT: [[SHOULD_TERMINATE:%.*]] = icmp eq i8* [[TMP2]], null 200 // CHECK5-NEXT: br i1 [[SHOULD_TERMINATE]], label [[DOTEXIT:%.*]], label [[DOTSELECT_WORKERS:%.*]] 201 // CHECK5: .select.workers: 202 // CHECK5-NEXT: [[TMP3:%.*]] = load i8, i8* [[EXEC_STATUS]], align 1 203 // CHECK5-NEXT: [[IS_ACTIVE:%.*]] = icmp ne i8 [[TMP3]], 0 204 // CHECK5-NEXT: br i1 [[IS_ACTIVE]], label [[DOTEXECUTE_PARALLEL:%.*]], label [[DOTBARRIER_PARALLEL:%.*]] 205 // CHECK5: .execute.parallel: 206 // CHECK5-NEXT: [[TMP4:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1]]) 207 // CHECK5-NEXT: [[TMP5:%.*]] = bitcast i8* [[TMP2]] to void (i16, i32)* 208 // CHECK5-NEXT: call void [[TMP5]](i16 0, i32 [[TMP4]]) 209 // CHECK5-NEXT: br label [[DOTTERMINATE_PARALLEL:%.*]] 210 // CHECK5: .terminate.parallel: 211 // CHECK5-NEXT: call void @__kmpc_kernel_end_parallel() 212 // CHECK5-NEXT: br label [[DOTBARRIER_PARALLEL]] 213 // CHECK5: .barrier.parallel: 214 // CHECK5-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0) 215 // CHECK5-NEXT: br label [[DOTAWAIT_WORK]] 216 // CHECK5: .exit: 217 // CHECK5-NEXT: ret void 218 // CHECK5-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIPPcEiT__l57 219 // CHECK5-SAME: (i64 [[A:%.*]], i64 [[B:%.*]], i8** [[ARGC:%.*]]) #[[ATTR1]] { 220 // CHECK5-NEXT: entry: 221 // CHECK5-NEXT: [[A_ADDR:%.*]] = alloca i64, align 8 222 // CHECK5-NEXT: [[B_ADDR:%.*]] = alloca i64, align 8 223 // CHECK5-NEXT: [[ARGC_ADDR:%.*]] = alloca i8**, align 8 224 // CHECK5-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4 225 // CHECK5-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 4 226 // CHECK5-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4 227 // CHECK5-NEXT: store i64 [[A]], i64* [[A_ADDR]], align 8 228 // CHECK5-NEXT: store i64 [[B]], i64* [[B_ADDR]], align 8 229 // CHECK5-NEXT: store i8** [[ARGC]], i8*** [[ARGC_ADDR]], align 8 230 // CHECK5-NEXT: [[CONV:%.*]] = bitcast i64* [[A_ADDR]] to i32* 231 // CHECK5-NEXT: [[CONV1:%.*]] = bitcast i64* [[B_ADDR]] to i32* 232 // CHECK5-NEXT: [[NVPTX_TID:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x() 233 // CHECK5-NEXT: [[NVPTX_NUM_THREADS:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() 234 // CHECK5-NEXT: [[NVPTX_WARP_SIZE:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize() 235 // CHECK5-NEXT: [[THREAD_LIMIT:%.*]] = sub nuw i32 [[NVPTX_NUM_THREADS]], [[NVPTX_WARP_SIZE]] 236 // CHECK5-NEXT: [[TMP0:%.*]] = icmp ult i32 [[NVPTX_TID]], [[THREAD_LIMIT]] 237 // CHECK5-NEXT: br i1 [[TMP0]], label [[DOTWORKER:%.*]], label [[DOTMASTERCHECK:%.*]] 238 // CHECK5: .worker: 239 // CHECK5-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIPPcEiT__l57_worker() #[[ATTR3]] 240 // CHECK5-NEXT: br label [[DOTEXIT:%.*]] 241 // CHECK5: .mastercheck: 242 // CHECK5-NEXT: [[NVPTX_TID2:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x() 243 // CHECK5-NEXT: [[NVPTX_NUM_THREADS3:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() 244 // CHECK5-NEXT: [[NVPTX_WARP_SIZE4:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize() 245 // CHECK5-NEXT: [[TMP1:%.*]] = sub nuw i32 [[NVPTX_WARP_SIZE4]], 1 246 // CHECK5-NEXT: [[TMP2:%.*]] = sub nuw i32 [[NVPTX_NUM_THREADS3]], 1 247 // CHECK5-NEXT: [[TMP3:%.*]] = xor i32 [[TMP1]], -1 248 // CHECK5-NEXT: [[MASTER_TID:%.*]] = and i32 [[TMP2]], [[TMP3]] 249 // CHECK5-NEXT: [[TMP4:%.*]] = icmp eq i32 [[NVPTX_TID2]], [[MASTER_TID]] 250 // CHECK5-NEXT: br i1 [[TMP4]], label [[DOTMASTER:%.*]], label [[DOTEXIT]] 251 // CHECK5: .master: 252 // CHECK5-NEXT: [[NVPTX_NUM_THREADS5:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() 253 // CHECK5-NEXT: [[NVPTX_WARP_SIZE6:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize() 254 // CHECK5-NEXT: [[THREAD_LIMIT7:%.*]] = sub nuw i32 [[NVPTX_NUM_THREADS5]], [[NVPTX_WARP_SIZE6]] 255 // CHECK5-NEXT: call void @__kmpc_kernel_init(i32 [[THREAD_LIMIT7]], i16 1) 256 // CHECK5-NEXT: call void @__kmpc_data_sharing_init_stack() 257 // CHECK5-NEXT: [[TMP5:%.*]] = load i16, i16* @"_openmp_static_kernel$is_shared1", align 2 258 // CHECK5-NEXT: [[TMP6:%.*]] = load i64, i64* @"_openmp_static_kernel$size2", align 8 259 // CHECK5-NEXT: call void @__kmpc_get_team_static_memory(i16 0, i8* addrspacecast (i8 addrspace(3)* getelementptr inbounds (%"union._shared_openmp_static_memory_type_$_", %"union._shared_openmp_static_memory_type_$_" addrspace(3)* @"_openmp_shared_static_glob_rd_$_", i32 0, i32 0, i32 0) to i8*), i64 [[TMP6]], i16 [[TMP5]], i8** addrspacecast (i8* addrspace(3)* @"_openmp_kernel_static_glob_rd$ptr" to i8**)) 260 // CHECK5-NEXT: [[TMP7:%.*]] = load i8*, i8* addrspace(3)* @"_openmp_kernel_static_glob_rd$ptr", align 8 261 // CHECK5-NEXT: [[TMP8:%.*]] = getelementptr inbounds i8, i8* [[TMP7]], i64 0 262 // CHECK5-NEXT: [[TMP9:%.*]] = bitcast i8* [[TMP8]] to %struct._globalized_locals_ty.0* 263 // CHECK5-NEXT: [[TMP10:%.*]] = load i8**, i8*** [[ARGC_ADDR]], align 8 264 // CHECK5-NEXT: [[ARGC8:%.*]] = getelementptr inbounds [[STRUCT__GLOBALIZED_LOCALS_TY_0:%.*]], %struct._globalized_locals_ty.0* [[TMP9]], i32 0, i32 0 265 // CHECK5-NEXT: store i8** [[TMP10]], i8*** [[ARGC8]], align 8 266 // CHECK5-NEXT: [[TMP11:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1]]) 267 // CHECK5-NEXT: store i32 [[TMP11]], i32* [[DOTTHREADID_TEMP_]], align 4 268 // CHECK5-NEXT: call void @__omp_outlined__3(i32* [[DOTTHREADID_TEMP_]], i32* [[DOTZERO_ADDR]], i8*** [[ARGC8]]) #[[ATTR3]] 269 // CHECK5-NEXT: [[TMP12:%.*]] = load i16, i16* @"_openmp_static_kernel$is_shared1", align 2 270 // CHECK5-NEXT: call void @__kmpc_restore_team_static_memory(i16 0, i16 [[TMP12]]) 271 // CHECK5-NEXT: br label [[DOTTERMINATION_NOTIFIER:%.*]] 272 // CHECK5: .termination.notifier: 273 // CHECK5-NEXT: call void @__kmpc_kernel_deinit(i16 1) 274 // CHECK5-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0) 275 // CHECK5-NEXT: br label [[DOTEXIT]] 276 // CHECK5: .exit: 277 // CHECK5-NEXT: ret void 278 // CHECK5-LABEL: define {{[^@]+}}@__omp_outlined__3 279 // CHECK5-SAME: (i32* noalias [[DOTGLOBAL_TID_:%.*]], i32* noalias [[DOTBOUND_TID_:%.*]], i8*** nonnull align 8 dereferenceable(8) [[ARGC:%.*]]) #[[ATTR1]] { 280 // CHECK5-NEXT: entry: 281 // CHECK5-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8 282 // CHECK5-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8 283 // CHECK5-NEXT: [[ARGC_ADDR:%.*]] = alloca i8***, align 8 284 // CHECK5-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 8 285 // CHECK5-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 8 286 // CHECK5-NEXT: store i8*** [[ARGC]], i8**** [[ARGC_ADDR]], align 8 287 // CHECK5-NEXT: [[TMP0:%.*]] = load i8***, i8**** [[ARGC_ADDR]], align 8 288 // CHECK5-NEXT: store i8** null, i8*** [[TMP0]], align 8 289 // CHECK5-NEXT: ret void 290 // CHECK6-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l68_worker 291 // CHECK6-SAME: () #[[ATTR0:[0-9]+]] { 292 // CHECK6-NEXT: entry: 293 // CHECK6-NEXT: [[WORK_FN:%.*]] = alloca i8*, align 8 294 // CHECK6-NEXT: [[EXEC_STATUS:%.*]] = alloca i8, align 1 295 // CHECK6-NEXT: store i8* null, i8** [[WORK_FN]], align 8 296 // CHECK6-NEXT: store i8 0, i8* [[EXEC_STATUS]], align 1 297 // CHECK6-NEXT: br label [[DOTAWAIT_WORK:%.*]] 298 // CHECK6: .await.work: 299 // CHECK6-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0) 300 // CHECK6-NEXT: [[TMP0:%.*]] = call i1 @__kmpc_kernel_parallel(i8** [[WORK_FN]]) 301 // CHECK6-NEXT: [[TMP1:%.*]] = zext i1 [[TMP0]] to i8 302 // CHECK6-NEXT: store i8 [[TMP1]], i8* [[EXEC_STATUS]], align 1 303 // CHECK6-NEXT: [[TMP2:%.*]] = load i8*, i8** [[WORK_FN]], align 8 304 // CHECK6-NEXT: [[SHOULD_TERMINATE:%.*]] = icmp eq i8* [[TMP2]], null 305 // CHECK6-NEXT: br i1 [[SHOULD_TERMINATE]], label [[DOTEXIT:%.*]], label [[DOTSELECT_WORKERS:%.*]] 306 // CHECK6: .select.workers: 307 // CHECK6-NEXT: [[TMP3:%.*]] = load i8, i8* [[EXEC_STATUS]], align 1 308 // CHECK6-NEXT: [[IS_ACTIVE:%.*]] = icmp ne i8 [[TMP3]], 0 309 // CHECK6-NEXT: br i1 [[IS_ACTIVE]], label [[DOTEXECUTE_PARALLEL:%.*]], label [[DOTBARRIER_PARALLEL:%.*]] 310 // CHECK6: .execute.parallel: 311 // CHECK6-NEXT: [[TMP4:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1:[0-9]+]]) 312 // CHECK6-NEXT: [[TMP5:%.*]] = bitcast i8* [[TMP2]] to void (i16, i32)* 313 // CHECK6-NEXT: call void [[TMP5]](i16 0, i32 [[TMP4]]) 314 // CHECK6-NEXT: br label [[DOTTERMINATE_PARALLEL:%.*]] 315 // CHECK6: .terminate.parallel: 316 // CHECK6-NEXT: call void @__kmpc_kernel_end_parallel() 317 // CHECK6-NEXT: br label [[DOTBARRIER_PARALLEL]] 318 // CHECK6: .barrier.parallel: 319 // CHECK6-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0) 320 // CHECK6-NEXT: br label [[DOTAWAIT_WORK]] 321 // CHECK6: .exit: 322 // CHECK6-NEXT: ret void 323 // CHECK6-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l68 324 // CHECK6-SAME: (i64 [[A:%.*]], i64 [[B:%.*]], i64 [[ARGC:%.*]]) #[[ATTR1:[0-9]+]] { 325 // CHECK6-NEXT: entry: 326 // CHECK6-NEXT: [[A_ADDR:%.*]] = alloca i64, align 8 327 // CHECK6-NEXT: [[B_ADDR:%.*]] = alloca i64, align 8 328 // CHECK6-NEXT: [[ARGC_ADDR:%.*]] = alloca i64, align 8 329 // CHECK6-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4 330 // CHECK6-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 4 331 // CHECK6-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4 332 // CHECK6-NEXT: store i64 [[A]], i64* [[A_ADDR]], align 8 333 // CHECK6-NEXT: store i64 [[B]], i64* [[B_ADDR]], align 8 334 // CHECK6-NEXT: store i64 [[ARGC]], i64* [[ARGC_ADDR]], align 8 335 // CHECK6-NEXT: [[CONV:%.*]] = bitcast i64* [[A_ADDR]] to i32* 336 // CHECK6-NEXT: [[CONV1:%.*]] = bitcast i64* [[B_ADDR]] to i32* 337 // CHECK6-NEXT: [[CONV2:%.*]] = bitcast i64* [[ARGC_ADDR]] to i32* 338 // CHECK6-NEXT: [[NVPTX_TID:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x() 339 // CHECK6-NEXT: [[NVPTX_NUM_THREADS:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() 340 // CHECK6-NEXT: [[NVPTX_WARP_SIZE:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize() 341 // CHECK6-NEXT: [[THREAD_LIMIT:%.*]] = sub nuw i32 [[NVPTX_NUM_THREADS]], [[NVPTX_WARP_SIZE]] 342 // CHECK6-NEXT: [[TMP0:%.*]] = icmp ult i32 [[NVPTX_TID]], [[THREAD_LIMIT]] 343 // CHECK6-NEXT: br i1 [[TMP0]], label [[DOTWORKER:%.*]], label [[DOTMASTERCHECK:%.*]] 344 // CHECK6: .worker: 345 // CHECK6-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l68_worker() #[[ATTR3:[0-9]+]] 346 // CHECK6-NEXT: br label [[DOTEXIT:%.*]] 347 // CHECK6: .mastercheck: 348 // CHECK6-NEXT: [[NVPTX_TID3:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x() 349 // CHECK6-NEXT: [[NVPTX_NUM_THREADS4:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() 350 // CHECK6-NEXT: [[NVPTX_WARP_SIZE5:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize() 351 // CHECK6-NEXT: [[TMP1:%.*]] = sub nuw i32 [[NVPTX_WARP_SIZE5]], 1 352 // CHECK6-NEXT: [[TMP2:%.*]] = sub nuw i32 [[NVPTX_NUM_THREADS4]], 1 353 // CHECK6-NEXT: [[TMP3:%.*]] = xor i32 [[TMP1]], -1 354 // CHECK6-NEXT: [[MASTER_TID:%.*]] = and i32 [[TMP2]], [[TMP3]] 355 // CHECK6-NEXT: [[TMP4:%.*]] = icmp eq i32 [[NVPTX_TID3]], [[MASTER_TID]] 356 // CHECK6-NEXT: br i1 [[TMP4]], label [[DOTMASTER:%.*]], label [[DOTEXIT]] 357 // CHECK6: .master: 358 // CHECK6-NEXT: [[NVPTX_NUM_THREADS6:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() 359 // CHECK6-NEXT: [[NVPTX_WARP_SIZE7:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize() 360 // CHECK6-NEXT: [[THREAD_LIMIT8:%.*]] = sub nuw i32 [[NVPTX_NUM_THREADS6]], [[NVPTX_WARP_SIZE7]] 361 // CHECK6-NEXT: call void @__kmpc_kernel_init(i32 [[THREAD_LIMIT8]], i16 1) 362 // CHECK6-NEXT: call void @__kmpc_data_sharing_init_stack() 363 // CHECK6-NEXT: [[TMP5:%.*]] = call i8* @__kmpc_data_sharing_push_stack(i64 4, i16 1) 364 // CHECK6-NEXT: [[TMP6:%.*]] = bitcast i8* [[TMP5]] to %struct._globalized_locals_ty* 365 // CHECK6-NEXT: [[TMP7:%.*]] = load i32, i32* [[CONV2]], align 8 366 // CHECK6-NEXT: [[ARGC9:%.*]] = getelementptr inbounds [[STRUCT__GLOBALIZED_LOCALS_TY:%.*]], %struct._globalized_locals_ty* [[TMP6]], i32 0, i32 0 367 // CHECK6-NEXT: store i32 [[TMP7]], i32* [[ARGC9]], align 4 368 // CHECK6-NEXT: [[TMP8:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1]]) 369 // CHECK6-NEXT: store i32 [[TMP8]], i32* [[DOTTHREADID_TEMP_]], align 4 370 // CHECK6-NEXT: call void @__omp_outlined__(i32* [[DOTTHREADID_TEMP_]], i32* [[DOTZERO_ADDR]], i32* [[ARGC9]]) #[[ATTR3]] 371 // CHECK6-NEXT: call void @__kmpc_data_sharing_pop_stack(i8* [[TMP5]]) 372 // CHECK6-NEXT: br label [[DOTTERMINATION_NOTIFIER:%.*]] 373 // CHECK6: .termination.notifier: 374 // CHECK6-NEXT: call void @__kmpc_kernel_deinit(i16 1) 375 // CHECK6-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0) 376 // CHECK6-NEXT: br label [[DOTEXIT]] 377 // CHECK6: .exit: 378 // CHECK6-NEXT: ret void 379 // CHECK6-LABEL: define {{[^@]+}}@__omp_outlined__ 380 // CHECK6-SAME: (i32* noalias [[DOTGLOBAL_TID_:%.*]], i32* noalias [[DOTBOUND_TID_:%.*]], i32* nonnull align 4 dereferenceable(4) [[ARGC:%.*]]) #[[ATTR1]] { 381 // CHECK6-NEXT: entry: 382 // CHECK6-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8 383 // CHECK6-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8 384 // CHECK6-NEXT: [[ARGC_ADDR:%.*]] = alloca i32*, align 8 385 // CHECK6-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 8 386 // CHECK6-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 8 387 // CHECK6-NEXT: store i32* [[ARGC]], i32** [[ARGC_ADDR]], align 8 388 // CHECK6-NEXT: [[TMP0:%.*]] = load i32*, i32** [[ARGC_ADDR]], align 8 389 // CHECK6-NEXT: store i32 0, i32* [[TMP0]], align 4 390 // CHECK6-NEXT: ret void 391 // CHECK6-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIPPcEiT__l57_worker 392 // CHECK6-SAME: () #[[ATTR0]] { 393 // CHECK6-NEXT: entry: 394 // CHECK6-NEXT: [[WORK_FN:%.*]] = alloca i8*, align 8 395 // CHECK6-NEXT: [[EXEC_STATUS:%.*]] = alloca i8, align 1 396 // CHECK6-NEXT: store i8* null, i8** [[WORK_FN]], align 8 397 // CHECK6-NEXT: store i8 0, i8* [[EXEC_STATUS]], align 1 398 // CHECK6-NEXT: br label [[DOTAWAIT_WORK:%.*]] 399 // CHECK6: .await.work: 400 // CHECK6-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0) 401 // CHECK6-NEXT: [[TMP0:%.*]] = call i1 @__kmpc_kernel_parallel(i8** [[WORK_FN]]) 402 // CHECK6-NEXT: [[TMP1:%.*]] = zext i1 [[TMP0]] to i8 403 // CHECK6-NEXT: store i8 [[TMP1]], i8* [[EXEC_STATUS]], align 1 404 // CHECK6-NEXT: [[TMP2:%.*]] = load i8*, i8** [[WORK_FN]], align 8 405 // CHECK6-NEXT: [[SHOULD_TERMINATE:%.*]] = icmp eq i8* [[TMP2]], null 406 // CHECK6-NEXT: br i1 [[SHOULD_TERMINATE]], label [[DOTEXIT:%.*]], label [[DOTSELECT_WORKERS:%.*]] 407 // CHECK6: .select.workers: 408 // CHECK6-NEXT: [[TMP3:%.*]] = load i8, i8* [[EXEC_STATUS]], align 1 409 // CHECK6-NEXT: [[IS_ACTIVE:%.*]] = icmp ne i8 [[TMP3]], 0 410 // CHECK6-NEXT: br i1 [[IS_ACTIVE]], label [[DOTEXECUTE_PARALLEL:%.*]], label [[DOTBARRIER_PARALLEL:%.*]] 411 // CHECK6: .execute.parallel: 412 // CHECK6-NEXT: [[TMP4:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1]]) 413 // CHECK6-NEXT: [[TMP5:%.*]] = bitcast i8* [[TMP2]] to void (i16, i32)* 414 // CHECK6-NEXT: call void [[TMP5]](i16 0, i32 [[TMP4]]) 415 // CHECK6-NEXT: br label [[DOTTERMINATE_PARALLEL:%.*]] 416 // CHECK6: .terminate.parallel: 417 // CHECK6-NEXT: call void @__kmpc_kernel_end_parallel() 418 // CHECK6-NEXT: br label [[DOTBARRIER_PARALLEL]] 419 // CHECK6: .barrier.parallel: 420 // CHECK6-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0) 421 // CHECK6-NEXT: br label [[DOTAWAIT_WORK]] 422 // CHECK6: .exit: 423 // CHECK6-NEXT: ret void 424 // CHECK6-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIPPcEiT__l57 425 // CHECK6-SAME: (i64 [[A:%.*]], i64 [[B:%.*]], i8** [[ARGC:%.*]]) #[[ATTR1]] { 426 // CHECK6-NEXT: entry: 427 // CHECK6-NEXT: [[A_ADDR:%.*]] = alloca i64, align 8 428 // CHECK6-NEXT: [[B_ADDR:%.*]] = alloca i64, align 8 429 // CHECK6-NEXT: [[ARGC_ADDR:%.*]] = alloca i8**, align 8 430 // CHECK6-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4 431 // CHECK6-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 4 432 // CHECK6-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4 433 // CHECK6-NEXT: store i64 [[A]], i64* [[A_ADDR]], align 8 434 // CHECK6-NEXT: store i64 [[B]], i64* [[B_ADDR]], align 8 435 // CHECK6-NEXT: store i8** [[ARGC]], i8*** [[ARGC_ADDR]], align 8 436 // CHECK6-NEXT: [[CONV:%.*]] = bitcast i64* [[A_ADDR]] to i32* 437 // CHECK6-NEXT: [[CONV1:%.*]] = bitcast i64* [[B_ADDR]] to i32* 438 // CHECK6-NEXT: [[NVPTX_TID:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x() 439 // CHECK6-NEXT: [[NVPTX_NUM_THREADS:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() 440 // CHECK6-NEXT: [[NVPTX_WARP_SIZE:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize() 441 // CHECK6-NEXT: [[THREAD_LIMIT:%.*]] = sub nuw i32 [[NVPTX_NUM_THREADS]], [[NVPTX_WARP_SIZE]] 442 // CHECK6-NEXT: [[TMP0:%.*]] = icmp ult i32 [[NVPTX_TID]], [[THREAD_LIMIT]] 443 // CHECK6-NEXT: br i1 [[TMP0]], label [[DOTWORKER:%.*]], label [[DOTMASTERCHECK:%.*]] 444 // CHECK6: .worker: 445 // CHECK6-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIPPcEiT__l57_worker() #[[ATTR3]] 446 // CHECK6-NEXT: br label [[DOTEXIT:%.*]] 447 // CHECK6: .mastercheck: 448 // CHECK6-NEXT: [[NVPTX_TID2:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x() 449 // CHECK6-NEXT: [[NVPTX_NUM_THREADS3:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() 450 // CHECK6-NEXT: [[NVPTX_WARP_SIZE4:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize() 451 // CHECK6-NEXT: [[TMP1:%.*]] = sub nuw i32 [[NVPTX_WARP_SIZE4]], 1 452 // CHECK6-NEXT: [[TMP2:%.*]] = sub nuw i32 [[NVPTX_NUM_THREADS3]], 1 453 // CHECK6-NEXT: [[TMP3:%.*]] = xor i32 [[TMP1]], -1 454 // CHECK6-NEXT: [[MASTER_TID:%.*]] = and i32 [[TMP2]], [[TMP3]] 455 // CHECK6-NEXT: [[TMP4:%.*]] = icmp eq i32 [[NVPTX_TID2]], [[MASTER_TID]] 456 // CHECK6-NEXT: br i1 [[TMP4]], label [[DOTMASTER:%.*]], label [[DOTEXIT]] 457 // CHECK6: .master: 458 // CHECK6-NEXT: [[NVPTX_NUM_THREADS5:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() 459 // CHECK6-NEXT: [[NVPTX_WARP_SIZE6:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize() 460 // CHECK6-NEXT: [[THREAD_LIMIT7:%.*]] = sub nuw i32 [[NVPTX_NUM_THREADS5]], [[NVPTX_WARP_SIZE6]] 461 // CHECK6-NEXT: call void @__kmpc_kernel_init(i32 [[THREAD_LIMIT7]], i16 1) 462 // CHECK6-NEXT: call void @__kmpc_data_sharing_init_stack() 463 // CHECK6-NEXT: [[TMP5:%.*]] = call i8* @__kmpc_data_sharing_push_stack(i64 8, i16 1) 464 // CHECK6-NEXT: [[TMP6:%.*]] = bitcast i8* [[TMP5]] to %struct._globalized_locals_ty.0* 465 // CHECK6-NEXT: [[TMP7:%.*]] = load i8**, i8*** [[ARGC_ADDR]], align 8 466 // CHECK6-NEXT: [[ARGC8:%.*]] = getelementptr inbounds [[STRUCT__GLOBALIZED_LOCALS_TY_0:%.*]], %struct._globalized_locals_ty.0* [[TMP6]], i32 0, i32 0 467 // CHECK6-NEXT: store i8** [[TMP7]], i8*** [[ARGC8]], align 8 468 // CHECK6-NEXT: [[TMP8:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1]]) 469 // CHECK6-NEXT: store i32 [[TMP8]], i32* [[DOTTHREADID_TEMP_]], align 4 470 // CHECK6-NEXT: call void @__omp_outlined__1(i32* [[DOTTHREADID_TEMP_]], i32* [[DOTZERO_ADDR]], i8*** [[ARGC8]]) #[[ATTR3]] 471 // CHECK6-NEXT: call void @__kmpc_data_sharing_pop_stack(i8* [[TMP5]]) 472 // CHECK6-NEXT: br label [[DOTTERMINATION_NOTIFIER:%.*]] 473 // CHECK6: .termination.notifier: 474 // CHECK6-NEXT: call void @__kmpc_kernel_deinit(i16 1) 475 // CHECK6-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0) 476 // CHECK6-NEXT: br label [[DOTEXIT]] 477 // CHECK6: .exit: 478 // CHECK6-NEXT: ret void 479 // CHECK6-LABEL: define {{[^@]+}}@__omp_outlined__1 480 // CHECK6-SAME: (i32* noalias [[DOTGLOBAL_TID_:%.*]], i32* noalias [[DOTBOUND_TID_:%.*]], i8*** nonnull align 8 dereferenceable(8) [[ARGC:%.*]]) #[[ATTR1]] { 481 // CHECK6-NEXT: entry: 482 // CHECK6-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8 483 // CHECK6-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8 484 // CHECK6-NEXT: [[ARGC_ADDR:%.*]] = alloca i8***, align 8 485 // CHECK6-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 8 486 // CHECK6-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 8 487 // CHECK6-NEXT: store i8*** [[ARGC]], i8**** [[ARGC_ADDR]], align 8 488 // CHECK6-NEXT: [[TMP0:%.*]] = load i8***, i8**** [[ARGC_ADDR]], align 8 489 // CHECK6-NEXT: store i8** null, i8*** [[TMP0]], align 8 490 // CHECK6-NEXT: ret void 491 // CHECK7-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l68_worker 492 // CHECK7-SAME: () #[[ATTR0:[0-9]+]] { 493 // CHECK7-NEXT: entry: 494 // CHECK7-NEXT: [[WORK_FN:%.*]] = alloca i8*, align 4 495 // CHECK7-NEXT: [[EXEC_STATUS:%.*]] = alloca i8, align 1 496 // CHECK7-NEXT: store i8* null, i8** [[WORK_FN]], align 4 497 // CHECK7-NEXT: store i8 0, i8* [[EXEC_STATUS]], align 1 498 // CHECK7-NEXT: br label [[DOTAWAIT_WORK:%.*]] 499 // CHECK7: .await.work: 500 // CHECK7-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0) 501 // CHECK7-NEXT: [[TMP0:%.*]] = call i1 @__kmpc_kernel_parallel(i8** [[WORK_FN]]) 502 // CHECK7-NEXT: [[TMP1:%.*]] = zext i1 [[TMP0]] to i8 503 // CHECK7-NEXT: store i8 [[TMP1]], i8* [[EXEC_STATUS]], align 1 504 // CHECK7-NEXT: [[TMP2:%.*]] = load i8*, i8** [[WORK_FN]], align 4 505 // CHECK7-NEXT: [[SHOULD_TERMINATE:%.*]] = icmp eq i8* [[TMP2]], null 506 // CHECK7-NEXT: br i1 [[SHOULD_TERMINATE]], label [[DOTEXIT:%.*]], label [[DOTSELECT_WORKERS:%.*]] 507 // CHECK7: .select.workers: 508 // CHECK7-NEXT: [[TMP3:%.*]] = load i8, i8* [[EXEC_STATUS]], align 1 509 // CHECK7-NEXT: [[IS_ACTIVE:%.*]] = icmp ne i8 [[TMP3]], 0 510 // CHECK7-NEXT: br i1 [[IS_ACTIVE]], label [[DOTEXECUTE_PARALLEL:%.*]], label [[DOTBARRIER_PARALLEL:%.*]] 511 // CHECK7: .execute.parallel: 512 // CHECK7-NEXT: [[TMP4:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1:[0-9]+]]) 513 // CHECK7-NEXT: [[TMP5:%.*]] = bitcast i8* [[TMP2]] to void (i16, i32)* 514 // CHECK7-NEXT: call void [[TMP5]](i16 0, i32 [[TMP4]]) 515 // CHECK7-NEXT: br label [[DOTTERMINATE_PARALLEL:%.*]] 516 // CHECK7: .terminate.parallel: 517 // CHECK7-NEXT: call void @__kmpc_kernel_end_parallel() 518 // CHECK7-NEXT: br label [[DOTBARRIER_PARALLEL]] 519 // CHECK7: .barrier.parallel: 520 // CHECK7-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0) 521 // CHECK7-NEXT: br label [[DOTAWAIT_WORK]] 522 // CHECK7: .exit: 523 // CHECK7-NEXT: ret void 524 // CHECK7-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l68 525 // CHECK7-SAME: (i32 [[A:%.*]], i32 [[B:%.*]], i32 [[ARGC:%.*]]) #[[ATTR1:[0-9]+]] { 526 // CHECK7-NEXT: entry: 527 // CHECK7-NEXT: [[A_ADDR:%.*]] = alloca i32, align 4 528 // CHECK7-NEXT: [[B_ADDR:%.*]] = alloca i32, align 4 529 // CHECK7-NEXT: [[ARGC_ADDR:%.*]] = alloca i32, align 4 530 // CHECK7-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4 531 // CHECK7-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 4 532 // CHECK7-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4 533 // CHECK7-NEXT: store i32 [[A]], i32* [[A_ADDR]], align 4 534 // CHECK7-NEXT: store i32 [[B]], i32* [[B_ADDR]], align 4 535 // CHECK7-NEXT: store i32 [[ARGC]], i32* [[ARGC_ADDR]], align 4 536 // CHECK7-NEXT: [[NVPTX_TID:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x() 537 // CHECK7-NEXT: [[NVPTX_NUM_THREADS:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() 538 // CHECK7-NEXT: [[NVPTX_WARP_SIZE:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize() 539 // CHECK7-NEXT: [[THREAD_LIMIT:%.*]] = sub nuw i32 [[NVPTX_NUM_THREADS]], [[NVPTX_WARP_SIZE]] 540 // CHECK7-NEXT: [[TMP0:%.*]] = icmp ult i32 [[NVPTX_TID]], [[THREAD_LIMIT]] 541 // CHECK7-NEXT: br i1 [[TMP0]], label [[DOTWORKER:%.*]], label [[DOTMASTERCHECK:%.*]] 542 // CHECK7: .worker: 543 // CHECK7-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l68_worker() #[[ATTR3:[0-9]+]] 544 // CHECK7-NEXT: br label [[DOTEXIT:%.*]] 545 // CHECK7: .mastercheck: 546 // CHECK7-NEXT: [[NVPTX_TID1:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x() 547 // CHECK7-NEXT: [[NVPTX_NUM_THREADS2:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() 548 // CHECK7-NEXT: [[NVPTX_WARP_SIZE3:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize() 549 // CHECK7-NEXT: [[TMP1:%.*]] = sub nuw i32 [[NVPTX_WARP_SIZE3]], 1 550 // CHECK7-NEXT: [[TMP2:%.*]] = sub nuw i32 [[NVPTX_NUM_THREADS2]], 1 551 // CHECK7-NEXT: [[TMP3:%.*]] = xor i32 [[TMP1]], -1 552 // CHECK7-NEXT: [[MASTER_TID:%.*]] = and i32 [[TMP2]], [[TMP3]] 553 // CHECK7-NEXT: [[TMP4:%.*]] = icmp eq i32 [[NVPTX_TID1]], [[MASTER_TID]] 554 // CHECK7-NEXT: br i1 [[TMP4]], label [[DOTMASTER:%.*]], label [[DOTEXIT]] 555 // CHECK7: .master: 556 // CHECK7-NEXT: [[NVPTX_NUM_THREADS4:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() 557 // CHECK7-NEXT: [[NVPTX_WARP_SIZE5:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize() 558 // CHECK7-NEXT: [[THREAD_LIMIT6:%.*]] = sub nuw i32 [[NVPTX_NUM_THREADS4]], [[NVPTX_WARP_SIZE5]] 559 // CHECK7-NEXT: call void @__kmpc_kernel_init(i32 [[THREAD_LIMIT6]], i16 1) 560 // CHECK7-NEXT: call void @__kmpc_data_sharing_init_stack() 561 // CHECK7-NEXT: [[TMP5:%.*]] = load i16, i16* @"_openmp_static_kernel$is_shared", align 2 562 // CHECK7-NEXT: [[TMP6:%.*]] = load i32, i32* @"_openmp_static_kernel$size", align 4 563 // CHECK7-NEXT: call void @__kmpc_get_team_static_memory(i16 0, i8* addrspacecast (i8 addrspace(3)* getelementptr inbounds (%"union._shared_openmp_static_memory_type_$_", %"union._shared_openmp_static_memory_type_$_" addrspace(3)* @"_openmp_shared_static_glob_rd_$_", i32 0, i32 0, i32 0) to i8*), i32 [[TMP6]], i16 [[TMP5]], i8** addrspacecast (i8* addrspace(3)* @"_openmp_kernel_static_glob_rd$ptr" to i8**)) 564 // CHECK7-NEXT: [[TMP7:%.*]] = load i8*, i8* addrspace(3)* @"_openmp_kernel_static_glob_rd$ptr", align 4 565 // CHECK7-NEXT: [[TMP8:%.*]] = getelementptr inbounds i8, i8* [[TMP7]], i32 0 566 // CHECK7-NEXT: [[TMP9:%.*]] = bitcast i8* [[TMP8]] to %struct._globalized_locals_ty* 567 // CHECK7-NEXT: [[TMP10:%.*]] = load i32, i32* [[ARGC_ADDR]], align 4 568 // CHECK7-NEXT: [[ARGC7:%.*]] = getelementptr inbounds [[STRUCT__GLOBALIZED_LOCALS_TY:%.*]], %struct._globalized_locals_ty* [[TMP9]], i32 0, i32 0 569 // CHECK7-NEXT: store i32 [[TMP10]], i32* [[ARGC7]], align 4 570 // CHECK7-NEXT: [[TMP11:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1]]) 571 // CHECK7-NEXT: store i32 [[TMP11]], i32* [[DOTTHREADID_TEMP_]], align 4 572 // CHECK7-NEXT: call void @__omp_outlined__(i32* [[DOTTHREADID_TEMP_]], i32* [[DOTZERO_ADDR]], i32* [[ARGC7]]) #[[ATTR3]] 573 // CHECK7-NEXT: [[TMP12:%.*]] = load i16, i16* @"_openmp_static_kernel$is_shared", align 2 574 // CHECK7-NEXT: call void @__kmpc_restore_team_static_memory(i16 0, i16 [[TMP12]]) 575 // CHECK7-NEXT: br label [[DOTTERMINATION_NOTIFIER:%.*]] 576 // CHECK7: .termination.notifier: 577 // CHECK7-NEXT: call void @__kmpc_kernel_deinit(i16 1) 578 // CHECK7-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0) 579 // CHECK7-NEXT: br label [[DOTEXIT]] 580 // CHECK7: .exit: 581 // CHECK7-NEXT: ret void 582 // CHECK7-LABEL: define {{[^@]+}}@__omp_outlined__ 583 // CHECK7-SAME: (i32* noalias [[DOTGLOBAL_TID_:%.*]], i32* noalias [[DOTBOUND_TID_:%.*]], i32* nonnull align 4 dereferenceable(4) [[ARGC:%.*]]) #[[ATTR1]] { 584 // CHECK7-NEXT: entry: 585 // CHECK7-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 4 586 // CHECK7-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 4 587 // CHECK7-NEXT: [[ARGC_ADDR:%.*]] = alloca i32*, align 4 588 // CHECK7-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 4 589 // CHECK7-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 4 590 // CHECK7-NEXT: store i32* [[ARGC]], i32** [[ARGC_ADDR]], align 4 591 // CHECK7-NEXT: [[TMP0:%.*]] = load i32*, i32** [[ARGC_ADDR]], align 4 592 // CHECK7-NEXT: store i32 0, i32* [[TMP0]], align 4 593 // CHECK7-NEXT: ret void 594 // CHECK7-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIPPcEiT__l57_worker 595 // CHECK7-SAME: () #[[ATTR0]] { 596 // CHECK7-NEXT: entry: 597 // CHECK7-NEXT: [[WORK_FN:%.*]] = alloca i8*, align 4 598 // CHECK7-NEXT: [[EXEC_STATUS:%.*]] = alloca i8, align 1 599 // CHECK7-NEXT: store i8* null, i8** [[WORK_FN]], align 4 600 // CHECK7-NEXT: store i8 0, i8* [[EXEC_STATUS]], align 1 601 // CHECK7-NEXT: br label [[DOTAWAIT_WORK:%.*]] 602 // CHECK7: .await.work: 603 // CHECK7-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0) 604 // CHECK7-NEXT: [[TMP0:%.*]] = call i1 @__kmpc_kernel_parallel(i8** [[WORK_FN]]) 605 // CHECK7-NEXT: [[TMP1:%.*]] = zext i1 [[TMP0]] to i8 606 // CHECK7-NEXT: store i8 [[TMP1]], i8* [[EXEC_STATUS]], align 1 607 // CHECK7-NEXT: [[TMP2:%.*]] = load i8*, i8** [[WORK_FN]], align 4 608 // CHECK7-NEXT: [[SHOULD_TERMINATE:%.*]] = icmp eq i8* [[TMP2]], null 609 // CHECK7-NEXT: br i1 [[SHOULD_TERMINATE]], label [[DOTEXIT:%.*]], label [[DOTSELECT_WORKERS:%.*]] 610 // CHECK7: .select.workers: 611 // CHECK7-NEXT: [[TMP3:%.*]] = load i8, i8* [[EXEC_STATUS]], align 1 612 // CHECK7-NEXT: [[IS_ACTIVE:%.*]] = icmp ne i8 [[TMP3]], 0 613 // CHECK7-NEXT: br i1 [[IS_ACTIVE]], label [[DOTEXECUTE_PARALLEL:%.*]], label [[DOTBARRIER_PARALLEL:%.*]] 614 // CHECK7: .execute.parallel: 615 // CHECK7-NEXT: [[TMP4:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1]]) 616 // CHECK7-NEXT: [[TMP5:%.*]] = bitcast i8* [[TMP2]] to void (i16, i32)* 617 // CHECK7-NEXT: call void [[TMP5]](i16 0, i32 [[TMP4]]) 618 // CHECK7-NEXT: br label [[DOTTERMINATE_PARALLEL:%.*]] 619 // CHECK7: .terminate.parallel: 620 // CHECK7-NEXT: call void @__kmpc_kernel_end_parallel() 621 // CHECK7-NEXT: br label [[DOTBARRIER_PARALLEL]] 622 // CHECK7: .barrier.parallel: 623 // CHECK7-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0) 624 // CHECK7-NEXT: br label [[DOTAWAIT_WORK]] 625 // CHECK7: .exit: 626 // CHECK7-NEXT: ret void 627 // CHECK7-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIPPcEiT__l57 628 // CHECK7-SAME: (i32 [[A:%.*]], i32 [[B:%.*]], i8** [[ARGC:%.*]]) #[[ATTR1]] { 629 // CHECK7-NEXT: entry: 630 // CHECK7-NEXT: [[A_ADDR:%.*]] = alloca i32, align 4 631 // CHECK7-NEXT: [[B_ADDR:%.*]] = alloca i32, align 4 632 // CHECK7-NEXT: [[ARGC_ADDR:%.*]] = alloca i8**, align 4 633 // CHECK7-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4 634 // CHECK7-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 4 635 // CHECK7-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4 636 // CHECK7-NEXT: store i32 [[A]], i32* [[A_ADDR]], align 4 637 // CHECK7-NEXT: store i32 [[B]], i32* [[B_ADDR]], align 4 638 // CHECK7-NEXT: store i8** [[ARGC]], i8*** [[ARGC_ADDR]], align 4 639 // CHECK7-NEXT: [[NVPTX_TID:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x() 640 // CHECK7-NEXT: [[NVPTX_NUM_THREADS:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() 641 // CHECK7-NEXT: [[NVPTX_WARP_SIZE:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize() 642 // CHECK7-NEXT: [[THREAD_LIMIT:%.*]] = sub nuw i32 [[NVPTX_NUM_THREADS]], [[NVPTX_WARP_SIZE]] 643 // CHECK7-NEXT: [[TMP0:%.*]] = icmp ult i32 [[NVPTX_TID]], [[THREAD_LIMIT]] 644 // CHECK7-NEXT: br i1 [[TMP0]], label [[DOTWORKER:%.*]], label [[DOTMASTERCHECK:%.*]] 645 // CHECK7: .worker: 646 // CHECK7-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIPPcEiT__l57_worker() #[[ATTR3]] 647 // CHECK7-NEXT: br label [[DOTEXIT:%.*]] 648 // CHECK7: .mastercheck: 649 // CHECK7-NEXT: [[NVPTX_TID1:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x() 650 // CHECK7-NEXT: [[NVPTX_NUM_THREADS2:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() 651 // CHECK7-NEXT: [[NVPTX_WARP_SIZE3:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize() 652 // CHECK7-NEXT: [[TMP1:%.*]] = sub nuw i32 [[NVPTX_WARP_SIZE3]], 1 653 // CHECK7-NEXT: [[TMP2:%.*]] = sub nuw i32 [[NVPTX_NUM_THREADS2]], 1 654 // CHECK7-NEXT: [[TMP3:%.*]] = xor i32 [[TMP1]], -1 655 // CHECK7-NEXT: [[MASTER_TID:%.*]] = and i32 [[TMP2]], [[TMP3]] 656 // CHECK7-NEXT: [[TMP4:%.*]] = icmp eq i32 [[NVPTX_TID1]], [[MASTER_TID]] 657 // CHECK7-NEXT: br i1 [[TMP4]], label [[DOTMASTER:%.*]], label [[DOTEXIT]] 658 // CHECK7: .master: 659 // CHECK7-NEXT: [[NVPTX_NUM_THREADS4:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() 660 // CHECK7-NEXT: [[NVPTX_WARP_SIZE5:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize() 661 // CHECK7-NEXT: [[THREAD_LIMIT6:%.*]] = sub nuw i32 [[NVPTX_NUM_THREADS4]], [[NVPTX_WARP_SIZE5]] 662 // CHECK7-NEXT: call void @__kmpc_kernel_init(i32 [[THREAD_LIMIT6]], i16 1) 663 // CHECK7-NEXT: call void @__kmpc_data_sharing_init_stack() 664 // CHECK7-NEXT: [[TMP5:%.*]] = load i16, i16* @"_openmp_static_kernel$is_shared1", align 2 665 // CHECK7-NEXT: [[TMP6:%.*]] = load i32, i32* @"_openmp_static_kernel$size2", align 4 666 // CHECK7-NEXT: call void @__kmpc_get_team_static_memory(i16 0, i8* addrspacecast (i8 addrspace(3)* getelementptr inbounds (%"union._shared_openmp_static_memory_type_$_", %"union._shared_openmp_static_memory_type_$_" addrspace(3)* @"_openmp_shared_static_glob_rd_$_", i32 0, i32 0, i32 0) to i8*), i32 [[TMP6]], i16 [[TMP5]], i8** addrspacecast (i8* addrspace(3)* @"_openmp_kernel_static_glob_rd$ptr" to i8**)) 667 // CHECK7-NEXT: [[TMP7:%.*]] = load i8*, i8* addrspace(3)* @"_openmp_kernel_static_glob_rd$ptr", align 4 668 // CHECK7-NEXT: [[TMP8:%.*]] = getelementptr inbounds i8, i8* [[TMP7]], i32 0 669 // CHECK7-NEXT: [[TMP9:%.*]] = bitcast i8* [[TMP8]] to %struct._globalized_locals_ty.0* 670 // CHECK7-NEXT: [[TMP10:%.*]] = load i8**, i8*** [[ARGC_ADDR]], align 4 671 // CHECK7-NEXT: [[ARGC7:%.*]] = getelementptr inbounds [[STRUCT__GLOBALIZED_LOCALS_TY_0:%.*]], %struct._globalized_locals_ty.0* [[TMP9]], i32 0, i32 0 672 // CHECK7-NEXT: store i8** [[TMP10]], i8*** [[ARGC7]], align 4 673 // CHECK7-NEXT: [[TMP11:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1]]) 674 // CHECK7-NEXT: store i32 [[TMP11]], i32* [[DOTTHREADID_TEMP_]], align 4 675 // CHECK7-NEXT: call void @__omp_outlined__3(i32* [[DOTTHREADID_TEMP_]], i32* [[DOTZERO_ADDR]], i8*** [[ARGC7]]) #[[ATTR3]] 676 // CHECK7-NEXT: [[TMP12:%.*]] = load i16, i16* @"_openmp_static_kernel$is_shared1", align 2 677 // CHECK7-NEXT: call void @__kmpc_restore_team_static_memory(i16 0, i16 [[TMP12]]) 678 // CHECK7-NEXT: br label [[DOTTERMINATION_NOTIFIER:%.*]] 679 // CHECK7: .termination.notifier: 680 // CHECK7-NEXT: call void @__kmpc_kernel_deinit(i16 1) 681 // CHECK7-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0) 682 // CHECK7-NEXT: br label [[DOTEXIT]] 683 // CHECK7: .exit: 684 // CHECK7-NEXT: ret void 685 // CHECK7-LABEL: define {{[^@]+}}@__omp_outlined__3 686 // CHECK7-SAME: (i32* noalias [[DOTGLOBAL_TID_:%.*]], i32* noalias [[DOTBOUND_TID_:%.*]], i8*** nonnull align 4 dereferenceable(4) [[ARGC:%.*]]) #[[ATTR1]] { 687 // CHECK7-NEXT: entry: 688 // CHECK7-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 4 689 // CHECK7-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 4 690 // CHECK7-NEXT: [[ARGC_ADDR:%.*]] = alloca i8***, align 4 691 // CHECK7-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 4 692 // CHECK7-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 4 693 // CHECK7-NEXT: store i8*** [[ARGC]], i8**** [[ARGC_ADDR]], align 4 694 // CHECK7-NEXT: [[TMP0:%.*]] = load i8***, i8**** [[ARGC_ADDR]], align 4 695 // CHECK7-NEXT: store i8** null, i8*** [[TMP0]], align 4 696 // CHECK7-NEXT: ret void 697 // CHECK8-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l68_worker 698 // CHECK8-SAME: () #[[ATTR0:[0-9]+]] { 699 // CHECK8-NEXT: entry: 700 // CHECK8-NEXT: [[WORK_FN:%.*]] = alloca i8*, align 4 701 // CHECK8-NEXT: [[EXEC_STATUS:%.*]] = alloca i8, align 1 702 // CHECK8-NEXT: store i8* null, i8** [[WORK_FN]], align 4 703 // CHECK8-NEXT: store i8 0, i8* [[EXEC_STATUS]], align 1 704 // CHECK8-NEXT: br label [[DOTAWAIT_WORK:%.*]] 705 // CHECK8: .await.work: 706 // CHECK8-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0) 707 // CHECK8-NEXT: [[TMP0:%.*]] = call i1 @__kmpc_kernel_parallel(i8** [[WORK_FN]]) 708 // CHECK8-NEXT: [[TMP1:%.*]] = zext i1 [[TMP0]] to i8 709 // CHECK8-NEXT: store i8 [[TMP1]], i8* [[EXEC_STATUS]], align 1 710 // CHECK8-NEXT: [[TMP2:%.*]] = load i8*, i8** [[WORK_FN]], align 4 711 // CHECK8-NEXT: [[SHOULD_TERMINATE:%.*]] = icmp eq i8* [[TMP2]], null 712 // CHECK8-NEXT: br i1 [[SHOULD_TERMINATE]], label [[DOTEXIT:%.*]], label [[DOTSELECT_WORKERS:%.*]] 713 // CHECK8: .select.workers: 714 // CHECK8-NEXT: [[TMP3:%.*]] = load i8, i8* [[EXEC_STATUS]], align 1 715 // CHECK8-NEXT: [[IS_ACTIVE:%.*]] = icmp ne i8 [[TMP3]], 0 716 // CHECK8-NEXT: br i1 [[IS_ACTIVE]], label [[DOTEXECUTE_PARALLEL:%.*]], label [[DOTBARRIER_PARALLEL:%.*]] 717 // CHECK8: .execute.parallel: 718 // CHECK8-NEXT: [[TMP4:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1:[0-9]+]]) 719 // CHECK8-NEXT: [[TMP5:%.*]] = bitcast i8* [[TMP2]] to void (i16, i32)* 720 // CHECK8-NEXT: call void [[TMP5]](i16 0, i32 [[TMP4]]) 721 // CHECK8-NEXT: br label [[DOTTERMINATE_PARALLEL:%.*]] 722 // CHECK8: .terminate.parallel: 723 // CHECK8-NEXT: call void @__kmpc_kernel_end_parallel() 724 // CHECK8-NEXT: br label [[DOTBARRIER_PARALLEL]] 725 // CHECK8: .barrier.parallel: 726 // CHECK8-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0) 727 // CHECK8-NEXT: br label [[DOTAWAIT_WORK]] 728 // CHECK8: .exit: 729 // CHECK8-NEXT: ret void 730 // CHECK8-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l68 731 // CHECK8-SAME: (i32 [[A:%.*]], i32 [[B:%.*]], i32 [[ARGC:%.*]]) #[[ATTR1:[0-9]+]] { 732 // CHECK8-NEXT: entry: 733 // CHECK8-NEXT: [[A_ADDR:%.*]] = alloca i32, align 4 734 // CHECK8-NEXT: [[B_ADDR:%.*]] = alloca i32, align 4 735 // CHECK8-NEXT: [[ARGC_ADDR:%.*]] = alloca i32, align 4 736 // CHECK8-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4 737 // CHECK8-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 4 738 // CHECK8-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4 739 // CHECK8-NEXT: store i32 [[A]], i32* [[A_ADDR]], align 4 740 // CHECK8-NEXT: store i32 [[B]], i32* [[B_ADDR]], align 4 741 // CHECK8-NEXT: store i32 [[ARGC]], i32* [[ARGC_ADDR]], align 4 742 // CHECK8-NEXT: [[NVPTX_TID:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x() 743 // CHECK8-NEXT: [[NVPTX_NUM_THREADS:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() 744 // CHECK8-NEXT: [[NVPTX_WARP_SIZE:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize() 745 // CHECK8-NEXT: [[THREAD_LIMIT:%.*]] = sub nuw i32 [[NVPTX_NUM_THREADS]], [[NVPTX_WARP_SIZE]] 746 // CHECK8-NEXT: [[TMP0:%.*]] = icmp ult i32 [[NVPTX_TID]], [[THREAD_LIMIT]] 747 // CHECK8-NEXT: br i1 [[TMP0]], label [[DOTWORKER:%.*]], label [[DOTMASTERCHECK:%.*]] 748 // CHECK8: .worker: 749 // CHECK8-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l68_worker() #[[ATTR3:[0-9]+]] 750 // CHECK8-NEXT: br label [[DOTEXIT:%.*]] 751 // CHECK8: .mastercheck: 752 // CHECK8-NEXT: [[NVPTX_TID1:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x() 753 // CHECK8-NEXT: [[NVPTX_NUM_THREADS2:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() 754 // CHECK8-NEXT: [[NVPTX_WARP_SIZE3:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize() 755 // CHECK8-NEXT: [[TMP1:%.*]] = sub nuw i32 [[NVPTX_WARP_SIZE3]], 1 756 // CHECK8-NEXT: [[TMP2:%.*]] = sub nuw i32 [[NVPTX_NUM_THREADS2]], 1 757 // CHECK8-NEXT: [[TMP3:%.*]] = xor i32 [[TMP1]], -1 758 // CHECK8-NEXT: [[MASTER_TID:%.*]] = and i32 [[TMP2]], [[TMP3]] 759 // CHECK8-NEXT: [[TMP4:%.*]] = icmp eq i32 [[NVPTX_TID1]], [[MASTER_TID]] 760 // CHECK8-NEXT: br i1 [[TMP4]], label [[DOTMASTER:%.*]], label [[DOTEXIT]] 761 // CHECK8: .master: 762 // CHECK8-NEXT: [[NVPTX_NUM_THREADS4:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() 763 // CHECK8-NEXT: [[NVPTX_WARP_SIZE5:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize() 764 // CHECK8-NEXT: [[THREAD_LIMIT6:%.*]] = sub nuw i32 [[NVPTX_NUM_THREADS4]], [[NVPTX_WARP_SIZE5]] 765 // CHECK8-NEXT: call void @__kmpc_kernel_init(i32 [[THREAD_LIMIT6]], i16 1) 766 // CHECK8-NEXT: call void @__kmpc_data_sharing_init_stack() 767 // CHECK8-NEXT: [[TMP5:%.*]] = call i8* @__kmpc_data_sharing_push_stack(i32 4, i16 1) 768 // CHECK8-NEXT: [[TMP6:%.*]] = bitcast i8* [[TMP5]] to %struct._globalized_locals_ty* 769 // CHECK8-NEXT: [[TMP7:%.*]] = load i32, i32* [[ARGC_ADDR]], align 4 770 // CHECK8-NEXT: [[ARGC7:%.*]] = getelementptr inbounds [[STRUCT__GLOBALIZED_LOCALS_TY:%.*]], %struct._globalized_locals_ty* [[TMP6]], i32 0, i32 0 771 // CHECK8-NEXT: store i32 [[TMP7]], i32* [[ARGC7]], align 4 772 // CHECK8-NEXT: [[TMP8:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1]]) 773 // CHECK8-NEXT: store i32 [[TMP8]], i32* [[DOTTHREADID_TEMP_]], align 4 774 // CHECK8-NEXT: call void @__omp_outlined__(i32* [[DOTTHREADID_TEMP_]], i32* [[DOTZERO_ADDR]], i32* [[ARGC7]]) #[[ATTR3]] 775 // CHECK8-NEXT: call void @__kmpc_data_sharing_pop_stack(i8* [[TMP5]]) 776 // CHECK8-NEXT: br label [[DOTTERMINATION_NOTIFIER:%.*]] 777 // CHECK8: .termination.notifier: 778 // CHECK8-NEXT: call void @__kmpc_kernel_deinit(i16 1) 779 // CHECK8-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0) 780 // CHECK8-NEXT: br label [[DOTEXIT]] 781 // CHECK8: .exit: 782 // CHECK8-NEXT: ret void 783 // CHECK8-LABEL: define {{[^@]+}}@__omp_outlined__ 784 // CHECK8-SAME: (i32* noalias [[DOTGLOBAL_TID_:%.*]], i32* noalias [[DOTBOUND_TID_:%.*]], i32* nonnull align 4 dereferenceable(4) [[ARGC:%.*]]) #[[ATTR1]] { 785 // CHECK8-NEXT: entry: 786 // CHECK8-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 4 787 // CHECK8-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 4 788 // CHECK8-NEXT: [[ARGC_ADDR:%.*]] = alloca i32*, align 4 789 // CHECK8-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 4 790 // CHECK8-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 4 791 // CHECK8-NEXT: store i32* [[ARGC]], i32** [[ARGC_ADDR]], align 4 792 // CHECK8-NEXT: [[TMP0:%.*]] = load i32*, i32** [[ARGC_ADDR]], align 4 793 // CHECK8-NEXT: store i32 0, i32* [[TMP0]], align 4 794 // CHECK8-NEXT: ret void 795 // CHECK8-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIPPcEiT__l57_worker 796 // CHECK8-SAME: () #[[ATTR0]] { 797 // CHECK8-NEXT: entry: 798 // CHECK8-NEXT: [[WORK_FN:%.*]] = alloca i8*, align 4 799 // CHECK8-NEXT: [[EXEC_STATUS:%.*]] = alloca i8, align 1 800 // CHECK8-NEXT: store i8* null, i8** [[WORK_FN]], align 4 801 // CHECK8-NEXT: store i8 0, i8* [[EXEC_STATUS]], align 1 802 // CHECK8-NEXT: br label [[DOTAWAIT_WORK:%.*]] 803 // CHECK8: .await.work: 804 // CHECK8-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0) 805 // CHECK8-NEXT: [[TMP0:%.*]] = call i1 @__kmpc_kernel_parallel(i8** [[WORK_FN]]) 806 // CHECK8-NEXT: [[TMP1:%.*]] = zext i1 [[TMP0]] to i8 807 // CHECK8-NEXT: store i8 [[TMP1]], i8* [[EXEC_STATUS]], align 1 808 // CHECK8-NEXT: [[TMP2:%.*]] = load i8*, i8** [[WORK_FN]], align 4 809 // CHECK8-NEXT: [[SHOULD_TERMINATE:%.*]] = icmp eq i8* [[TMP2]], null 810 // CHECK8-NEXT: br i1 [[SHOULD_TERMINATE]], label [[DOTEXIT:%.*]], label [[DOTSELECT_WORKERS:%.*]] 811 // CHECK8: .select.workers: 812 // CHECK8-NEXT: [[TMP3:%.*]] = load i8, i8* [[EXEC_STATUS]], align 1 813 // CHECK8-NEXT: [[IS_ACTIVE:%.*]] = icmp ne i8 [[TMP3]], 0 814 // CHECK8-NEXT: br i1 [[IS_ACTIVE]], label [[DOTEXECUTE_PARALLEL:%.*]], label [[DOTBARRIER_PARALLEL:%.*]] 815 // CHECK8: .execute.parallel: 816 // CHECK8-NEXT: [[TMP4:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1]]) 817 // CHECK8-NEXT: [[TMP5:%.*]] = bitcast i8* [[TMP2]] to void (i16, i32)* 818 // CHECK8-NEXT: call void [[TMP5]](i16 0, i32 [[TMP4]]) 819 // CHECK8-NEXT: br label [[DOTTERMINATE_PARALLEL:%.*]] 820 // CHECK8: .terminate.parallel: 821 // CHECK8-NEXT: call void @__kmpc_kernel_end_parallel() 822 // CHECK8-NEXT: br label [[DOTBARRIER_PARALLEL]] 823 // CHECK8: .barrier.parallel: 824 // CHECK8-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0) 825 // CHECK8-NEXT: br label [[DOTAWAIT_WORK]] 826 // CHECK8: .exit: 827 // CHECK8-NEXT: ret void 828 // CHECK8-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIPPcEiT__l57 829 // CHECK8-SAME: (i32 [[A:%.*]], i32 [[B:%.*]], i8** [[ARGC:%.*]]) #[[ATTR1]] { 830 // CHECK8-NEXT: entry: 831 // CHECK8-NEXT: [[A_ADDR:%.*]] = alloca i32, align 4 832 // CHECK8-NEXT: [[B_ADDR:%.*]] = alloca i32, align 4 833 // CHECK8-NEXT: [[ARGC_ADDR:%.*]] = alloca i8**, align 4 834 // CHECK8-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4 835 // CHECK8-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 4 836 // CHECK8-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4 837 // CHECK8-NEXT: store i32 [[A]], i32* [[A_ADDR]], align 4 838 // CHECK8-NEXT: store i32 [[B]], i32* [[B_ADDR]], align 4 839 // CHECK8-NEXT: store i8** [[ARGC]], i8*** [[ARGC_ADDR]], align 4 840 // CHECK8-NEXT: [[NVPTX_TID:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x() 841 // CHECK8-NEXT: [[NVPTX_NUM_THREADS:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() 842 // CHECK8-NEXT: [[NVPTX_WARP_SIZE:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize() 843 // CHECK8-NEXT: [[THREAD_LIMIT:%.*]] = sub nuw i32 [[NVPTX_NUM_THREADS]], [[NVPTX_WARP_SIZE]] 844 // CHECK8-NEXT: [[TMP0:%.*]] = icmp ult i32 [[NVPTX_TID]], [[THREAD_LIMIT]] 845 // CHECK8-NEXT: br i1 [[TMP0]], label [[DOTWORKER:%.*]], label [[DOTMASTERCHECK:%.*]] 846 // CHECK8: .worker: 847 // CHECK8-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIPPcEiT__l57_worker() #[[ATTR3]] 848 // CHECK8-NEXT: br label [[DOTEXIT:%.*]] 849 // CHECK8: .mastercheck: 850 // CHECK8-NEXT: [[NVPTX_TID1:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x() 851 // CHECK8-NEXT: [[NVPTX_NUM_THREADS2:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() 852 // CHECK8-NEXT: [[NVPTX_WARP_SIZE3:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize() 853 // CHECK8-NEXT: [[TMP1:%.*]] = sub nuw i32 [[NVPTX_WARP_SIZE3]], 1 854 // CHECK8-NEXT: [[TMP2:%.*]] = sub nuw i32 [[NVPTX_NUM_THREADS2]], 1 855 // CHECK8-NEXT: [[TMP3:%.*]] = xor i32 [[TMP1]], -1 856 // CHECK8-NEXT: [[MASTER_TID:%.*]] = and i32 [[TMP2]], [[TMP3]] 857 // CHECK8-NEXT: [[TMP4:%.*]] = icmp eq i32 [[NVPTX_TID1]], [[MASTER_TID]] 858 // CHECK8-NEXT: br i1 [[TMP4]], label [[DOTMASTER:%.*]], label [[DOTEXIT]] 859 // CHECK8: .master: 860 // CHECK8-NEXT: [[NVPTX_NUM_THREADS4:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() 861 // CHECK8-NEXT: [[NVPTX_WARP_SIZE5:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize() 862 // CHECK8-NEXT: [[THREAD_LIMIT6:%.*]] = sub nuw i32 [[NVPTX_NUM_THREADS4]], [[NVPTX_WARP_SIZE5]] 863 // CHECK8-NEXT: call void @__kmpc_kernel_init(i32 [[THREAD_LIMIT6]], i16 1) 864 // CHECK8-NEXT: call void @__kmpc_data_sharing_init_stack() 865 // CHECK8-NEXT: [[TMP5:%.*]] = call i8* @__kmpc_data_sharing_push_stack(i32 4, i16 1) 866 // CHECK8-NEXT: [[TMP6:%.*]] = bitcast i8* [[TMP5]] to %struct._globalized_locals_ty.0* 867 // CHECK8-NEXT: [[TMP7:%.*]] = load i8**, i8*** [[ARGC_ADDR]], align 4 868 // CHECK8-NEXT: [[ARGC7:%.*]] = getelementptr inbounds [[STRUCT__GLOBALIZED_LOCALS_TY_0:%.*]], %struct._globalized_locals_ty.0* [[TMP6]], i32 0, i32 0 869 // CHECK8-NEXT: store i8** [[TMP7]], i8*** [[ARGC7]], align 4 870 // CHECK8-NEXT: [[TMP8:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1]]) 871 // CHECK8-NEXT: store i32 [[TMP8]], i32* [[DOTTHREADID_TEMP_]], align 4 872 // CHECK8-NEXT: call void @__omp_outlined__1(i32* [[DOTTHREADID_TEMP_]], i32* [[DOTZERO_ADDR]], i8*** [[ARGC7]]) #[[ATTR3]] 873 // CHECK8-NEXT: call void @__kmpc_data_sharing_pop_stack(i8* [[TMP5]]) 874 // CHECK8-NEXT: br label [[DOTTERMINATION_NOTIFIER:%.*]] 875 // CHECK8: .termination.notifier: 876 // CHECK8-NEXT: call void @__kmpc_kernel_deinit(i16 1) 877 // CHECK8-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0) 878 // CHECK8-NEXT: br label [[DOTEXIT]] 879 // CHECK8: .exit: 880 // CHECK8-NEXT: ret void 881 // CHECK8-LABEL: define {{[^@]+}}@__omp_outlined__1 882 // CHECK8-SAME: (i32* noalias [[DOTGLOBAL_TID_:%.*]], i32* noalias [[DOTBOUND_TID_:%.*]], i8*** nonnull align 4 dereferenceable(4) [[ARGC:%.*]]) #[[ATTR1]] { 883 // CHECK8-NEXT: entry: 884 // CHECK8-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 4 885 // CHECK8-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 4 886 // CHECK8-NEXT: [[ARGC_ADDR:%.*]] = alloca i8***, align 4 887 // CHECK8-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 4 888 // CHECK8-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 4 889 // CHECK8-NEXT: store i8*** [[ARGC]], i8**** [[ARGC_ADDR]], align 4 890 // CHECK8-NEXT: [[TMP0:%.*]] = load i8***, i8**** [[ARGC_ADDR]], align 4 891 // CHECK8-NEXT: store i8** null, i8*** [[TMP0]], align 4 892 // CHECK8-NEXT: ret void 893 // CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l23_worker 894 // CHECK1-SAME: () #[[ATTR0:[0-9]+]] { 895 // CHECK1-NEXT: entry: 896 // CHECK1-NEXT: [[WORK_FN:%.*]] = alloca i8*, align 8 897 // CHECK1-NEXT: [[EXEC_STATUS:%.*]] = alloca i8, align 1 898 // CHECK1-NEXT: store i8* null, i8** [[WORK_FN]], align 8 899 // CHECK1-NEXT: store i8 0, i8* [[EXEC_STATUS]], align 1 900 // CHECK1-NEXT: br label [[DOTAWAIT_WORK:%.*]] 901 // CHECK1: .await.work: 902 // CHECK1-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0) 903 // CHECK1-NEXT: [[TMP0:%.*]] = call i1 @__kmpc_kernel_parallel(i8** [[WORK_FN]]) 904 // CHECK1-NEXT: [[TMP1:%.*]] = zext i1 [[TMP0]] to i8 905 // CHECK1-NEXT: store i8 [[TMP1]], i8* [[EXEC_STATUS]], align 1 906 // CHECK1-NEXT: [[TMP2:%.*]] = load i8*, i8** [[WORK_FN]], align 8 907 // CHECK1-NEXT: [[SHOULD_TERMINATE:%.*]] = icmp eq i8* [[TMP2]], null 908 // CHECK1-NEXT: br i1 [[SHOULD_TERMINATE]], label [[DOTEXIT:%.*]], label [[DOTSELECT_WORKERS:%.*]] 909 // CHECK1: .select.workers: 910 // CHECK1-NEXT: [[TMP3:%.*]] = load i8, i8* [[EXEC_STATUS]], align 1 911 // CHECK1-NEXT: [[IS_ACTIVE:%.*]] = icmp ne i8 [[TMP3]], 0 912 // CHECK1-NEXT: br i1 [[IS_ACTIVE]], label [[DOTEXECUTE_PARALLEL:%.*]], label [[DOTBARRIER_PARALLEL:%.*]] 913 // CHECK1: .execute.parallel: 914 // CHECK1-NEXT: [[TMP4:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1:[0-9]+]]) 915 // CHECK1-NEXT: [[TMP5:%.*]] = bitcast i8* [[TMP2]] to void (i16, i32)* 916 // CHECK1-NEXT: call void [[TMP5]](i16 0, i32 [[TMP4]]) 917 // CHECK1-NEXT: br label [[DOTTERMINATE_PARALLEL:%.*]] 918 // CHECK1: .terminate.parallel: 919 // CHECK1-NEXT: call void @__kmpc_kernel_end_parallel() 920 // CHECK1-NEXT: br label [[DOTBARRIER_PARALLEL]] 921 // CHECK1: .barrier.parallel: 922 // CHECK1-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0) 923 // CHECK1-NEXT: br label [[DOTAWAIT_WORK]] 924 // CHECK1: .exit: 925 // CHECK1-NEXT: ret void 926 // 927 // 928 // CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l23 929 // CHECK1-SAME: (i64 [[ARGC:%.*]]) #[[ATTR1:[0-9]+]] { 930 // CHECK1-NEXT: entry: 931 // CHECK1-NEXT: [[ARGC_ADDR:%.*]] = alloca i64, align 8 932 // CHECK1-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4 933 // CHECK1-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 4 934 // CHECK1-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4 935 // CHECK1-NEXT: store i64 [[ARGC]], i64* [[ARGC_ADDR]], align 8 936 // CHECK1-NEXT: [[CONV:%.*]] = bitcast i64* [[ARGC_ADDR]] to i32* 937 // CHECK1-NEXT: [[NVPTX_TID:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x() 938 // CHECK1-NEXT: [[NVPTX_NUM_THREADS:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() 939 // CHECK1-NEXT: [[NVPTX_WARP_SIZE:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize() 940 // CHECK1-NEXT: [[THREAD_LIMIT:%.*]] = sub nuw i32 [[NVPTX_NUM_THREADS]], [[NVPTX_WARP_SIZE]] 941 // CHECK1-NEXT: [[TMP0:%.*]] = icmp ult i32 [[NVPTX_TID]], [[THREAD_LIMIT]] 942 // CHECK1-NEXT: br i1 [[TMP0]], label [[DOTWORKER:%.*]], label [[DOTMASTERCHECK:%.*]] 943 // CHECK1: .worker: 944 // CHECK1-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l23_worker() #[[ATTR3:[0-9]+]] 945 // CHECK1-NEXT: br label [[DOTEXIT:%.*]] 946 // CHECK1: .mastercheck: 947 // CHECK1-NEXT: [[NVPTX_TID1:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x() 948 // CHECK1-NEXT: [[NVPTX_NUM_THREADS2:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() 949 // CHECK1-NEXT: [[NVPTX_WARP_SIZE3:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize() 950 // CHECK1-NEXT: [[TMP1:%.*]] = sub nuw i32 [[NVPTX_WARP_SIZE3]], 1 951 // CHECK1-NEXT: [[TMP2:%.*]] = sub nuw i32 [[NVPTX_NUM_THREADS2]], 1 952 // CHECK1-NEXT: [[TMP3:%.*]] = xor i32 [[TMP1]], -1 953 // CHECK1-NEXT: [[MASTER_TID:%.*]] = and i32 [[TMP2]], [[TMP3]] 954 // CHECK1-NEXT: [[TMP4:%.*]] = icmp eq i32 [[NVPTX_TID1]], [[MASTER_TID]] 955 // CHECK1-NEXT: br i1 [[TMP4]], label [[DOTMASTER:%.*]], label [[DOTEXIT]] 956 // CHECK1: .master: 957 // CHECK1-NEXT: [[NVPTX_NUM_THREADS4:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() 958 // CHECK1-NEXT: [[NVPTX_WARP_SIZE5:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize() 959 // CHECK1-NEXT: [[THREAD_LIMIT6:%.*]] = sub nuw i32 [[NVPTX_NUM_THREADS4]], [[NVPTX_WARP_SIZE5]] 960 // CHECK1-NEXT: call void @__kmpc_kernel_init(i32 [[THREAD_LIMIT6]], i16 1) 961 // CHECK1-NEXT: [[TMP5:%.*]] = load i32, i32* [[CONV]], align 8 962 // CHECK1-NEXT: [[ARGC7:%.*]] = call i8* @__kmpc_alloc_shared(i64 4) 963 // CHECK1-NEXT: [[ARGC_ON_STACK:%.*]] = bitcast i8* [[ARGC7]] to i32* 964 // CHECK1-NEXT: store i32 [[TMP5]], i32* [[ARGC_ON_STACK]], align 4 965 // CHECK1-NEXT: [[TMP6:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1]]) 966 // CHECK1-NEXT: store i32 [[TMP6]], i32* [[DOTTHREADID_TEMP_]], align 4 967 // CHECK1-NEXT: call void @__omp_outlined__(i32* [[DOTTHREADID_TEMP_]], i32* [[DOTZERO_ADDR]], i32* [[ARGC_ON_STACK]]) #[[ATTR3]] 968 // CHECK1-NEXT: call void @__kmpc_free_shared(i8* [[ARGC7]]) 969 // CHECK1-NEXT: br label [[DOTTERMINATION_NOTIFIER:%.*]] 970 // CHECK1: .termination.notifier: 971 // CHECK1-NEXT: call void @__kmpc_kernel_deinit(i16 1) 972 // CHECK1-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0) 973 // CHECK1-NEXT: br label [[DOTEXIT]] 974 // CHECK1: .exit: 975 // CHECK1-NEXT: ret void 976 // 977 // 978 // CHECK1-LABEL: define {{[^@]+}}@__omp_outlined__ 979 // CHECK1-SAME: (i32* noalias [[DOTGLOBAL_TID_:%.*]], i32* noalias [[DOTBOUND_TID_:%.*]], i32* nonnull align 4 dereferenceable(4) [[ARGC:%.*]]) #[[ATTR1]] { 980 // CHECK1-NEXT: entry: 981 // CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8 982 // CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8 983 // CHECK1-NEXT: [[ARGC_ADDR:%.*]] = alloca i32*, align 8 984 // CHECK1-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 8 985 // CHECK1-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 8 986 // CHECK1-NEXT: store i32* [[ARGC]], i32** [[ARGC_ADDR]], align 8 987 // CHECK1-NEXT: [[TMP0:%.*]] = load i32*, i32** [[ARGC_ADDR]], align 8 988 // CHECK1-NEXT: store i32 0, i32* [[TMP0]], align 4 989 // CHECK1-NEXT: ret void 990 // 991 // 992 // CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIPPcEiT__l15_worker 993 // CHECK1-SAME: () #[[ATTR0]] { 994 // CHECK1-NEXT: entry: 995 // CHECK1-NEXT: [[WORK_FN:%.*]] = alloca i8*, align 8 996 // CHECK1-NEXT: [[EXEC_STATUS:%.*]] = alloca i8, align 1 997 // CHECK1-NEXT: store i8* null, i8** [[WORK_FN]], align 8 998 // CHECK1-NEXT: store i8 0, i8* [[EXEC_STATUS]], align 1 999 // CHECK1-NEXT: br label [[DOTAWAIT_WORK:%.*]] 1000 // CHECK1: .await.work: 1001 // CHECK1-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0) 1002 // CHECK1-NEXT: [[TMP0:%.*]] = call i1 @__kmpc_kernel_parallel(i8** [[WORK_FN]]) 1003 // CHECK1-NEXT: [[TMP1:%.*]] = zext i1 [[TMP0]] to i8 1004 // CHECK1-NEXT: store i8 [[TMP1]], i8* [[EXEC_STATUS]], align 1 1005 // CHECK1-NEXT: [[TMP2:%.*]] = load i8*, i8** [[WORK_FN]], align 8 1006 // CHECK1-NEXT: [[SHOULD_TERMINATE:%.*]] = icmp eq i8* [[TMP2]], null 1007 // CHECK1-NEXT: br i1 [[SHOULD_TERMINATE]], label [[DOTEXIT:%.*]], label [[DOTSELECT_WORKERS:%.*]] 1008 // CHECK1: .select.workers: 1009 // CHECK1-NEXT: [[TMP3:%.*]] = load i8, i8* [[EXEC_STATUS]], align 1 1010 // CHECK1-NEXT: [[IS_ACTIVE:%.*]] = icmp ne i8 [[TMP3]], 0 1011 // CHECK1-NEXT: br i1 [[IS_ACTIVE]], label [[DOTEXECUTE_PARALLEL:%.*]], label [[DOTBARRIER_PARALLEL:%.*]] 1012 // CHECK1: .execute.parallel: 1013 // CHECK1-NEXT: [[TMP4:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1]]) 1014 // CHECK1-NEXT: [[TMP5:%.*]] = bitcast i8* [[TMP2]] to void (i16, i32)* 1015 // CHECK1-NEXT: call void [[TMP5]](i16 0, i32 [[TMP4]]) 1016 // CHECK1-NEXT: br label [[DOTTERMINATE_PARALLEL:%.*]] 1017 // CHECK1: .terminate.parallel: 1018 // CHECK1-NEXT: call void @__kmpc_kernel_end_parallel() 1019 // CHECK1-NEXT: br label [[DOTBARRIER_PARALLEL]] 1020 // CHECK1: .barrier.parallel: 1021 // CHECK1-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0) 1022 // CHECK1-NEXT: br label [[DOTAWAIT_WORK]] 1023 // CHECK1: .exit: 1024 // CHECK1-NEXT: ret void 1025 // 1026 // 1027 // CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIPPcEiT__l15 1028 // CHECK1-SAME: (i8** [[ARGC:%.*]]) #[[ATTR1]] { 1029 // CHECK1-NEXT: entry: 1030 // CHECK1-NEXT: [[ARGC_ADDR:%.*]] = alloca i8**, align 8 1031 // CHECK1-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4 1032 // CHECK1-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 4 1033 // CHECK1-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4 1034 // CHECK1-NEXT: store i8** [[ARGC]], i8*** [[ARGC_ADDR]], align 8 1035 // CHECK1-NEXT: [[NVPTX_TID:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x() 1036 // CHECK1-NEXT: [[NVPTX_NUM_THREADS:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() 1037 // CHECK1-NEXT: [[NVPTX_WARP_SIZE:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize() 1038 // CHECK1-NEXT: [[THREAD_LIMIT:%.*]] = sub nuw i32 [[NVPTX_NUM_THREADS]], [[NVPTX_WARP_SIZE]] 1039 // CHECK1-NEXT: [[TMP0:%.*]] = icmp ult i32 [[NVPTX_TID]], [[THREAD_LIMIT]] 1040 // CHECK1-NEXT: br i1 [[TMP0]], label [[DOTWORKER:%.*]], label [[DOTMASTERCHECK:%.*]] 1041 // CHECK1: .worker: 1042 // CHECK1-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIPPcEiT__l15_worker() #[[ATTR3]] 1043 // CHECK1-NEXT: br label [[DOTEXIT:%.*]] 1044 // CHECK1: .mastercheck: 1045 // CHECK1-NEXT: [[NVPTX_TID1:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x() 1046 // CHECK1-NEXT: [[NVPTX_NUM_THREADS2:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() 1047 // CHECK1-NEXT: [[NVPTX_WARP_SIZE3:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize() 1048 // CHECK1-NEXT: [[TMP1:%.*]] = sub nuw i32 [[NVPTX_WARP_SIZE3]], 1 1049 // CHECK1-NEXT: [[TMP2:%.*]] = sub nuw i32 [[NVPTX_NUM_THREADS2]], 1 1050 // CHECK1-NEXT: [[TMP3:%.*]] = xor i32 [[TMP1]], -1 1051 // CHECK1-NEXT: [[MASTER_TID:%.*]] = and i32 [[TMP2]], [[TMP3]] 1052 // CHECK1-NEXT: [[TMP4:%.*]] = icmp eq i32 [[NVPTX_TID1]], [[MASTER_TID]] 1053 // CHECK1-NEXT: br i1 [[TMP4]], label [[DOTMASTER:%.*]], label [[DOTEXIT]] 1054 // CHECK1: .master: 1055 // CHECK1-NEXT: [[NVPTX_NUM_THREADS4:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() 1056 // CHECK1-NEXT: [[NVPTX_WARP_SIZE5:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize() 1057 // CHECK1-NEXT: [[THREAD_LIMIT6:%.*]] = sub nuw i32 [[NVPTX_NUM_THREADS4]], [[NVPTX_WARP_SIZE5]] 1058 // CHECK1-NEXT: call void @__kmpc_kernel_init(i32 [[THREAD_LIMIT6]], i16 1) 1059 // CHECK1-NEXT: [[TMP5:%.*]] = load i8**, i8*** [[ARGC_ADDR]], align 8 1060 // CHECK1-NEXT: [[ARGC7:%.*]] = call i8* @__kmpc_alloc_shared(i64 8) 1061 // CHECK1-NEXT: [[ARGC_ON_STACK:%.*]] = bitcast i8* [[ARGC7]] to i8*** 1062 // CHECK1-NEXT: store i8** [[TMP5]], i8*** [[ARGC_ON_STACK]], align 8 1063 // CHECK1-NEXT: [[TMP6:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1]]) 1064 // CHECK1-NEXT: store i32 [[TMP6]], i32* [[DOTTHREADID_TEMP_]], align 4 1065 // CHECK1-NEXT: call void @__omp_outlined__1(i32* [[DOTTHREADID_TEMP_]], i32* [[DOTZERO_ADDR]], i8*** [[ARGC_ON_STACK]]) #[[ATTR3]] 1066 // CHECK1-NEXT: call void @__kmpc_free_shared(i8* [[ARGC7]]) 1067 // CHECK1-NEXT: br label [[DOTTERMINATION_NOTIFIER:%.*]] 1068 // CHECK1: .termination.notifier: 1069 // CHECK1-NEXT: call void @__kmpc_kernel_deinit(i16 1) 1070 // CHECK1-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0) 1071 // CHECK1-NEXT: br label [[DOTEXIT]] 1072 // CHECK1: .exit: 1073 // CHECK1-NEXT: ret void 1074 // 1075 // 1076 // CHECK1-LABEL: define {{[^@]+}}@__omp_outlined__1 1077 // CHECK1-SAME: (i32* noalias [[DOTGLOBAL_TID_:%.*]], i32* noalias [[DOTBOUND_TID_:%.*]], i8*** nonnull align 8 dereferenceable(8) [[ARGC:%.*]]) #[[ATTR1]] { 1078 // CHECK1-NEXT: entry: 1079 // CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8 1080 // CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8 1081 // CHECK1-NEXT: [[ARGC_ADDR:%.*]] = alloca i8***, align 8 1082 // CHECK1-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 8 1083 // CHECK1-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 8 1084 // CHECK1-NEXT: store i8*** [[ARGC]], i8**** [[ARGC_ADDR]], align 8 1085 // CHECK1-NEXT: [[TMP0:%.*]] = load i8***, i8**** [[ARGC_ADDR]], align 8 1086 // CHECK1-NEXT: store i8** null, i8*** [[TMP0]], align 8 1087 // CHECK1-NEXT: ret void 1088 // 1089 // 1090 // CHECK2-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l23_worker 1091 // CHECK2-SAME: () #[[ATTR0:[0-9]+]] { 1092 // CHECK2-NEXT: entry: 1093 // CHECK2-NEXT: [[WORK_FN:%.*]] = alloca i8*, align 4 1094 // CHECK2-NEXT: [[EXEC_STATUS:%.*]] = alloca i8, align 1 1095 // CHECK2-NEXT: store i8* null, i8** [[WORK_FN]], align 4 1096 // CHECK2-NEXT: store i8 0, i8* [[EXEC_STATUS]], align 1 1097 // CHECK2-NEXT: br label [[DOTAWAIT_WORK:%.*]] 1098 // CHECK2: .await.work: 1099 // CHECK2-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0) 1100 // CHECK2-NEXT: [[TMP0:%.*]] = call i1 @__kmpc_kernel_parallel(i8** [[WORK_FN]]) 1101 // CHECK2-NEXT: [[TMP1:%.*]] = zext i1 [[TMP0]] to i8 1102 // CHECK2-NEXT: store i8 [[TMP1]], i8* [[EXEC_STATUS]], align 1 1103 // CHECK2-NEXT: [[TMP2:%.*]] = load i8*, i8** [[WORK_FN]], align 4 1104 // CHECK2-NEXT: [[SHOULD_TERMINATE:%.*]] = icmp eq i8* [[TMP2]], null 1105 // CHECK2-NEXT: br i1 [[SHOULD_TERMINATE]], label [[DOTEXIT:%.*]], label [[DOTSELECT_WORKERS:%.*]] 1106 // CHECK2: .select.workers: 1107 // CHECK2-NEXT: [[TMP3:%.*]] = load i8, i8* [[EXEC_STATUS]], align 1 1108 // CHECK2-NEXT: [[IS_ACTIVE:%.*]] = icmp ne i8 [[TMP3]], 0 1109 // CHECK2-NEXT: br i1 [[IS_ACTIVE]], label [[DOTEXECUTE_PARALLEL:%.*]], label [[DOTBARRIER_PARALLEL:%.*]] 1110 // CHECK2: .execute.parallel: 1111 // CHECK2-NEXT: [[TMP4:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1:[0-9]+]]) 1112 // CHECK2-NEXT: [[TMP5:%.*]] = bitcast i8* [[TMP2]] to void (i16, i32)* 1113 // CHECK2-NEXT: call void [[TMP5]](i16 0, i32 [[TMP4]]) 1114 // CHECK2-NEXT: br label [[DOTTERMINATE_PARALLEL:%.*]] 1115 // CHECK2: .terminate.parallel: 1116 // CHECK2-NEXT: call void @__kmpc_kernel_end_parallel() 1117 // CHECK2-NEXT: br label [[DOTBARRIER_PARALLEL]] 1118 // CHECK2: .barrier.parallel: 1119 // CHECK2-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0) 1120 // CHECK2-NEXT: br label [[DOTAWAIT_WORK]] 1121 // CHECK2: .exit: 1122 // CHECK2-NEXT: ret void 1123 // 1124 // 1125 // CHECK2-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l23 1126 // CHECK2-SAME: (i32 [[ARGC:%.*]]) #[[ATTR1:[0-9]+]] { 1127 // CHECK2-NEXT: entry: 1128 // CHECK2-NEXT: [[ARGC_ADDR:%.*]] = alloca i32, align 4 1129 // CHECK2-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4 1130 // CHECK2-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 4 1131 // CHECK2-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4 1132 // CHECK2-NEXT: store i32 [[ARGC]], i32* [[ARGC_ADDR]], align 4 1133 // CHECK2-NEXT: [[NVPTX_TID:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x() 1134 // CHECK2-NEXT: [[NVPTX_NUM_THREADS:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() 1135 // CHECK2-NEXT: [[NVPTX_WARP_SIZE:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize() 1136 // CHECK2-NEXT: [[THREAD_LIMIT:%.*]] = sub nuw i32 [[NVPTX_NUM_THREADS]], [[NVPTX_WARP_SIZE]] 1137 // CHECK2-NEXT: [[TMP0:%.*]] = icmp ult i32 [[NVPTX_TID]], [[THREAD_LIMIT]] 1138 // CHECK2-NEXT: br i1 [[TMP0]], label [[DOTWORKER:%.*]], label [[DOTMASTERCHECK:%.*]] 1139 // CHECK2: .worker: 1140 // CHECK2-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l23_worker() #[[ATTR3:[0-9]+]] 1141 // CHECK2-NEXT: br label [[DOTEXIT:%.*]] 1142 // CHECK2: .mastercheck: 1143 // CHECK2-NEXT: [[NVPTX_TID1:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x() 1144 // CHECK2-NEXT: [[NVPTX_NUM_THREADS2:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() 1145 // CHECK2-NEXT: [[NVPTX_WARP_SIZE3:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize() 1146 // CHECK2-NEXT: [[TMP1:%.*]] = sub nuw i32 [[NVPTX_WARP_SIZE3]], 1 1147 // CHECK2-NEXT: [[TMP2:%.*]] = sub nuw i32 [[NVPTX_NUM_THREADS2]], 1 1148 // CHECK2-NEXT: [[TMP3:%.*]] = xor i32 [[TMP1]], -1 1149 // CHECK2-NEXT: [[MASTER_TID:%.*]] = and i32 [[TMP2]], [[TMP3]] 1150 // CHECK2-NEXT: [[TMP4:%.*]] = icmp eq i32 [[NVPTX_TID1]], [[MASTER_TID]] 1151 // CHECK2-NEXT: br i1 [[TMP4]], label [[DOTMASTER:%.*]], label [[DOTEXIT]] 1152 // CHECK2: .master: 1153 // CHECK2-NEXT: [[NVPTX_NUM_THREADS4:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() 1154 // CHECK2-NEXT: [[NVPTX_WARP_SIZE5:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize() 1155 // CHECK2-NEXT: [[THREAD_LIMIT6:%.*]] = sub nuw i32 [[NVPTX_NUM_THREADS4]], [[NVPTX_WARP_SIZE5]] 1156 // CHECK2-NEXT: call void @__kmpc_kernel_init(i32 [[THREAD_LIMIT6]], i16 1) 1157 // CHECK2-NEXT: [[TMP5:%.*]] = load i32, i32* [[ARGC_ADDR]], align 4 1158 // CHECK2-NEXT: [[ARGC7:%.*]] = call i8* @__kmpc_alloc_shared(i32 4) 1159 // CHECK2-NEXT: [[ARGC_ON_STACK:%.*]] = bitcast i8* [[ARGC7]] to i32* 1160 // CHECK2-NEXT: store i32 [[TMP5]], i32* [[ARGC_ON_STACK]], align 4 1161 // CHECK2-NEXT: [[TMP6:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1]]) 1162 // CHECK2-NEXT: store i32 [[TMP6]], i32* [[DOTTHREADID_TEMP_]], align 4 1163 // CHECK2-NEXT: call void @__omp_outlined__(i32* [[DOTTHREADID_TEMP_]], i32* [[DOTZERO_ADDR]], i32* [[ARGC_ON_STACK]]) #[[ATTR3]] 1164 // CHECK2-NEXT: call void @__kmpc_free_shared(i8* [[ARGC7]]) 1165 // CHECK2-NEXT: br label [[DOTTERMINATION_NOTIFIER:%.*]] 1166 // CHECK2: .termination.notifier: 1167 // CHECK2-NEXT: call void @__kmpc_kernel_deinit(i16 1) 1168 // CHECK2-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0) 1169 // CHECK2-NEXT: br label [[DOTEXIT]] 1170 // CHECK2: .exit: 1171 // CHECK2-NEXT: ret void 1172 // 1173 // 1174 // CHECK2-LABEL: define {{[^@]+}}@__omp_outlined__ 1175 // CHECK2-SAME: (i32* noalias [[DOTGLOBAL_TID_:%.*]], i32* noalias [[DOTBOUND_TID_:%.*]], i32* nonnull align 4 dereferenceable(4) [[ARGC:%.*]]) #[[ATTR1]] { 1176 // CHECK2-NEXT: entry: 1177 // CHECK2-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 4 1178 // CHECK2-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 4 1179 // CHECK2-NEXT: [[ARGC_ADDR:%.*]] = alloca i32*, align 4 1180 // CHECK2-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 4 1181 // CHECK2-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 4 1182 // CHECK2-NEXT: store i32* [[ARGC]], i32** [[ARGC_ADDR]], align 4 1183 // CHECK2-NEXT: [[TMP0:%.*]] = load i32*, i32** [[ARGC_ADDR]], align 4 1184 // CHECK2-NEXT: store i32 0, i32* [[TMP0]], align 4 1185 // CHECK2-NEXT: ret void 1186 // 1187 // 1188 // CHECK2-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIPPcEiT__l15_worker 1189 // CHECK2-SAME: () #[[ATTR0]] { 1190 // CHECK2-NEXT: entry: 1191 // CHECK2-NEXT: [[WORK_FN:%.*]] = alloca i8*, align 4 1192 // CHECK2-NEXT: [[EXEC_STATUS:%.*]] = alloca i8, align 1 1193 // CHECK2-NEXT: store i8* null, i8** [[WORK_FN]], align 4 1194 // CHECK2-NEXT: store i8 0, i8* [[EXEC_STATUS]], align 1 1195 // CHECK2-NEXT: br label [[DOTAWAIT_WORK:%.*]] 1196 // CHECK2: .await.work: 1197 // CHECK2-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0) 1198 // CHECK2-NEXT: [[TMP0:%.*]] = call i1 @__kmpc_kernel_parallel(i8** [[WORK_FN]]) 1199 // CHECK2-NEXT: [[TMP1:%.*]] = zext i1 [[TMP0]] to i8 1200 // CHECK2-NEXT: store i8 [[TMP1]], i8* [[EXEC_STATUS]], align 1 1201 // CHECK2-NEXT: [[TMP2:%.*]] = load i8*, i8** [[WORK_FN]], align 4 1202 // CHECK2-NEXT: [[SHOULD_TERMINATE:%.*]] = icmp eq i8* [[TMP2]], null 1203 // CHECK2-NEXT: br i1 [[SHOULD_TERMINATE]], label [[DOTEXIT:%.*]], label [[DOTSELECT_WORKERS:%.*]] 1204 // CHECK2: .select.workers: 1205 // CHECK2-NEXT: [[TMP3:%.*]] = load i8, i8* [[EXEC_STATUS]], align 1 1206 // CHECK2-NEXT: [[IS_ACTIVE:%.*]] = icmp ne i8 [[TMP3]], 0 1207 // CHECK2-NEXT: br i1 [[IS_ACTIVE]], label [[DOTEXECUTE_PARALLEL:%.*]], label [[DOTBARRIER_PARALLEL:%.*]] 1208 // CHECK2: .execute.parallel: 1209 // CHECK2-NEXT: [[TMP4:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1]]) 1210 // CHECK2-NEXT: [[TMP5:%.*]] = bitcast i8* [[TMP2]] to void (i16, i32)* 1211 // CHECK2-NEXT: call void [[TMP5]](i16 0, i32 [[TMP4]]) 1212 // CHECK2-NEXT: br label [[DOTTERMINATE_PARALLEL:%.*]] 1213 // CHECK2: .terminate.parallel: 1214 // CHECK2-NEXT: call void @__kmpc_kernel_end_parallel() 1215 // CHECK2-NEXT: br label [[DOTBARRIER_PARALLEL]] 1216 // CHECK2: .barrier.parallel: 1217 // CHECK2-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0) 1218 // CHECK2-NEXT: br label [[DOTAWAIT_WORK]] 1219 // CHECK2: .exit: 1220 // CHECK2-NEXT: ret void 1221 // 1222 // 1223 // CHECK2-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIPPcEiT__l15 1224 // CHECK2-SAME: (i8** [[ARGC:%.*]]) #[[ATTR1]] { 1225 // CHECK2-NEXT: entry: 1226 // CHECK2-NEXT: [[ARGC_ADDR:%.*]] = alloca i8**, align 4 1227 // CHECK2-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4 1228 // CHECK2-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 4 1229 // CHECK2-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4 1230 // CHECK2-NEXT: store i8** [[ARGC]], i8*** [[ARGC_ADDR]], align 4 1231 // CHECK2-NEXT: [[NVPTX_TID:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x() 1232 // CHECK2-NEXT: [[NVPTX_NUM_THREADS:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() 1233 // CHECK2-NEXT: [[NVPTX_WARP_SIZE:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize() 1234 // CHECK2-NEXT: [[THREAD_LIMIT:%.*]] = sub nuw i32 [[NVPTX_NUM_THREADS]], [[NVPTX_WARP_SIZE]] 1235 // CHECK2-NEXT: [[TMP0:%.*]] = icmp ult i32 [[NVPTX_TID]], [[THREAD_LIMIT]] 1236 // CHECK2-NEXT: br i1 [[TMP0]], label [[DOTWORKER:%.*]], label [[DOTMASTERCHECK:%.*]] 1237 // CHECK2: .worker: 1238 // CHECK2-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIPPcEiT__l15_worker() #[[ATTR3]] 1239 // CHECK2-NEXT: br label [[DOTEXIT:%.*]] 1240 // CHECK2: .mastercheck: 1241 // CHECK2-NEXT: [[NVPTX_TID1:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x() 1242 // CHECK2-NEXT: [[NVPTX_NUM_THREADS2:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() 1243 // CHECK2-NEXT: [[NVPTX_WARP_SIZE3:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize() 1244 // CHECK2-NEXT: [[TMP1:%.*]] = sub nuw i32 [[NVPTX_WARP_SIZE3]], 1 1245 // CHECK2-NEXT: [[TMP2:%.*]] = sub nuw i32 [[NVPTX_NUM_THREADS2]], 1 1246 // CHECK2-NEXT: [[TMP3:%.*]] = xor i32 [[TMP1]], -1 1247 // CHECK2-NEXT: [[MASTER_TID:%.*]] = and i32 [[TMP2]], [[TMP3]] 1248 // CHECK2-NEXT: [[TMP4:%.*]] = icmp eq i32 [[NVPTX_TID1]], [[MASTER_TID]] 1249 // CHECK2-NEXT: br i1 [[TMP4]], label [[DOTMASTER:%.*]], label [[DOTEXIT]] 1250 // CHECK2: .master: 1251 // CHECK2-NEXT: [[NVPTX_NUM_THREADS4:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() 1252 // CHECK2-NEXT: [[NVPTX_WARP_SIZE5:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize() 1253 // CHECK2-NEXT: [[THREAD_LIMIT6:%.*]] = sub nuw i32 [[NVPTX_NUM_THREADS4]], [[NVPTX_WARP_SIZE5]] 1254 // CHECK2-NEXT: call void @__kmpc_kernel_init(i32 [[THREAD_LIMIT6]], i16 1) 1255 // CHECK2-NEXT: [[TMP5:%.*]] = load i8**, i8*** [[ARGC_ADDR]], align 4 1256 // CHECK2-NEXT: [[ARGC7:%.*]] = call i8* @__kmpc_alloc_shared(i32 4) 1257 // CHECK2-NEXT: [[ARGC_ON_STACK:%.*]] = bitcast i8* [[ARGC7]] to i8*** 1258 // CHECK2-NEXT: store i8** [[TMP5]], i8*** [[ARGC_ON_STACK]], align 4 1259 // CHECK2-NEXT: [[TMP6:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1]]) 1260 // CHECK2-NEXT: store i32 [[TMP6]], i32* [[DOTTHREADID_TEMP_]], align 4 1261 // CHECK2-NEXT: call void @__omp_outlined__1(i32* [[DOTTHREADID_TEMP_]], i32* [[DOTZERO_ADDR]], i8*** [[ARGC_ON_STACK]]) #[[ATTR3]] 1262 // CHECK2-NEXT: call void @__kmpc_free_shared(i8* [[ARGC7]]) 1263 // CHECK2-NEXT: br label [[DOTTERMINATION_NOTIFIER:%.*]] 1264 // CHECK2: .termination.notifier: 1265 // CHECK2-NEXT: call void @__kmpc_kernel_deinit(i16 1) 1266 // CHECK2-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0) 1267 // CHECK2-NEXT: br label [[DOTEXIT]] 1268 // CHECK2: .exit: 1269 // CHECK2-NEXT: ret void 1270 // 1271 // 1272 // CHECK2-LABEL: define {{[^@]+}}@__omp_outlined__1 1273 // CHECK2-SAME: (i32* noalias [[DOTGLOBAL_TID_:%.*]], i32* noalias [[DOTBOUND_TID_:%.*]], i8*** nonnull align 4 dereferenceable(4) [[ARGC:%.*]]) #[[ATTR1]] { 1274 // CHECK2-NEXT: entry: 1275 // CHECK2-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 4 1276 // CHECK2-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 4 1277 // CHECK2-NEXT: [[ARGC_ADDR:%.*]] = alloca i8***, align 4 1278 // CHECK2-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 4 1279 // CHECK2-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 4 1280 // CHECK2-NEXT: store i8*** [[ARGC]], i8**** [[ARGC_ADDR]], align 4 1281 // CHECK2-NEXT: [[TMP0:%.*]] = load i8***, i8**** [[ARGC_ADDR]], align 4 1282 // CHECK2-NEXT: store i8** null, i8*** [[TMP0]], align 4 1283 // CHECK2-NEXT: ret void 1284 // 1285 // 1286 // CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l64_worker 1287 // CHECK3-SAME: () #[[ATTR0:[0-9]+]] { 1288 // CHECK3-NEXT: entry: 1289 // CHECK3-NEXT: [[WORK_FN:%.*]] = alloca i8*, align 8 1290 // CHECK3-NEXT: [[EXEC_STATUS:%.*]] = alloca i8, align 1 1291 // CHECK3-NEXT: store i8* null, i8** [[WORK_FN]], align 8 1292 // CHECK3-NEXT: store i8 0, i8* [[EXEC_STATUS]], align 1 1293 // CHECK3-NEXT: br label [[DOTAWAIT_WORK:%.*]] 1294 // CHECK3: .await.work: 1295 // CHECK3-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0) 1296 // CHECK3-NEXT: [[TMP0:%.*]] = call i1 @__kmpc_kernel_parallel(i8** [[WORK_FN]]) 1297 // CHECK3-NEXT: [[TMP1:%.*]] = zext i1 [[TMP0]] to i8 1298 // CHECK3-NEXT: store i8 [[TMP1]], i8* [[EXEC_STATUS]], align 1 1299 // CHECK3-NEXT: [[TMP2:%.*]] = load i8*, i8** [[WORK_FN]], align 8 1300 // CHECK3-NEXT: [[SHOULD_TERMINATE:%.*]] = icmp eq i8* [[TMP2]], null 1301 // CHECK3-NEXT: br i1 [[SHOULD_TERMINATE]], label [[DOTEXIT:%.*]], label [[DOTSELECT_WORKERS:%.*]] 1302 // CHECK3: .select.workers: 1303 // CHECK3-NEXT: [[TMP3:%.*]] = load i8, i8* [[EXEC_STATUS]], align 1 1304 // CHECK3-NEXT: [[IS_ACTIVE:%.*]] = icmp ne i8 [[TMP3]], 0 1305 // CHECK3-NEXT: br i1 [[IS_ACTIVE]], label [[DOTEXECUTE_PARALLEL:%.*]], label [[DOTBARRIER_PARALLEL:%.*]] 1306 // CHECK3: .execute.parallel: 1307 // CHECK3-NEXT: [[TMP4:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1:[0-9]+]]) 1308 // CHECK3-NEXT: [[TMP5:%.*]] = bitcast i8* [[TMP2]] to void (i16, i32)* 1309 // CHECK3-NEXT: call void [[TMP5]](i16 0, i32 [[TMP4]]) 1310 // CHECK3-NEXT: br label [[DOTTERMINATE_PARALLEL:%.*]] 1311 // CHECK3: .terminate.parallel: 1312 // CHECK3-NEXT: call void @__kmpc_kernel_end_parallel() 1313 // CHECK3-NEXT: br label [[DOTBARRIER_PARALLEL]] 1314 // CHECK3: .barrier.parallel: 1315 // CHECK3-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0) 1316 // CHECK3-NEXT: br label [[DOTAWAIT_WORK]] 1317 // CHECK3: .exit: 1318 // CHECK3-NEXT: ret void 1319 // 1320 // 1321 // CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l64 1322 // CHECK3-SAME: (i64 [[A:%.*]], i64 [[B:%.*]], i64 [[ARGC:%.*]]) #[[ATTR1:[0-9]+]] { 1323 // CHECK3-NEXT: entry: 1324 // CHECK3-NEXT: [[A_ADDR:%.*]] = alloca i64, align 8 1325 // CHECK3-NEXT: [[B_ADDR:%.*]] = alloca i64, align 8 1326 // CHECK3-NEXT: [[ARGC_ADDR:%.*]] = alloca i64, align 8 1327 // CHECK3-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4 1328 // CHECK3-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 4 1329 // CHECK3-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4 1330 // CHECK3-NEXT: store i64 [[A]], i64* [[A_ADDR]], align 8 1331 // CHECK3-NEXT: store i64 [[B]], i64* [[B_ADDR]], align 8 1332 // CHECK3-NEXT: store i64 [[ARGC]], i64* [[ARGC_ADDR]], align 8 1333 // CHECK3-NEXT: [[CONV:%.*]] = bitcast i64* [[A_ADDR]] to i32* 1334 // CHECK3-NEXT: [[CONV1:%.*]] = bitcast i64* [[B_ADDR]] to i32* 1335 // CHECK3-NEXT: [[CONV2:%.*]] = bitcast i64* [[ARGC_ADDR]] to i32* 1336 // CHECK3-NEXT: [[NVPTX_TID:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x() 1337 // CHECK3-NEXT: [[NVPTX_NUM_THREADS:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() 1338 // CHECK3-NEXT: [[NVPTX_WARP_SIZE:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize() 1339 // CHECK3-NEXT: [[THREAD_LIMIT:%.*]] = sub nuw i32 [[NVPTX_NUM_THREADS]], [[NVPTX_WARP_SIZE]] 1340 // CHECK3-NEXT: [[TMP0:%.*]] = icmp ult i32 [[NVPTX_TID]], [[THREAD_LIMIT]] 1341 // CHECK3-NEXT: br i1 [[TMP0]], label [[DOTWORKER:%.*]], label [[DOTMASTERCHECK:%.*]] 1342 // CHECK3: .worker: 1343 // CHECK3-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l64_worker() #[[ATTR3:[0-9]+]] 1344 // CHECK3-NEXT: br label [[DOTEXIT:%.*]] 1345 // CHECK3: .mastercheck: 1346 // CHECK3-NEXT: [[NVPTX_TID3:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x() 1347 // CHECK3-NEXT: [[NVPTX_NUM_THREADS4:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() 1348 // CHECK3-NEXT: [[NVPTX_WARP_SIZE5:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize() 1349 // CHECK3-NEXT: [[TMP1:%.*]] = sub nuw i32 [[NVPTX_WARP_SIZE5]], 1 1350 // CHECK3-NEXT: [[TMP2:%.*]] = sub nuw i32 [[NVPTX_NUM_THREADS4]], 1 1351 // CHECK3-NEXT: [[TMP3:%.*]] = xor i32 [[TMP1]], -1 1352 // CHECK3-NEXT: [[MASTER_TID:%.*]] = and i32 [[TMP2]], [[TMP3]] 1353 // CHECK3-NEXT: [[TMP4:%.*]] = icmp eq i32 [[NVPTX_TID3]], [[MASTER_TID]] 1354 // CHECK3-NEXT: br i1 [[TMP4]], label [[DOTMASTER:%.*]], label [[DOTEXIT]] 1355 // CHECK3: .master: 1356 // CHECK3-NEXT: [[NVPTX_NUM_THREADS6:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() 1357 // CHECK3-NEXT: [[NVPTX_WARP_SIZE7:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize() 1358 // CHECK3-NEXT: [[THREAD_LIMIT8:%.*]] = sub nuw i32 [[NVPTX_NUM_THREADS6]], [[NVPTX_WARP_SIZE7]] 1359 // CHECK3-NEXT: call void @__kmpc_kernel_init(i32 [[THREAD_LIMIT8]], i16 1) 1360 // CHECK3-NEXT: [[TMP5:%.*]] = load i32, i32* [[CONV2]], align 8 1361 // CHECK3-NEXT: [[ARGC9:%.*]] = call i8* @__kmpc_alloc_shared(i64 4) 1362 // CHECK3-NEXT: [[ARGC_ON_STACK:%.*]] = bitcast i8* [[ARGC9]] to i32* 1363 // CHECK3-NEXT: store i32 [[TMP5]], i32* [[ARGC_ON_STACK]], align 4 1364 // CHECK3-NEXT: [[TMP6:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1]]) 1365 // CHECK3-NEXT: store i32 [[TMP6]], i32* [[DOTTHREADID_TEMP_]], align 4 1366 // CHECK3-NEXT: call void @__omp_outlined__(i32* [[DOTTHREADID_TEMP_]], i32* [[DOTZERO_ADDR]], i32* [[ARGC_ON_STACK]]) #[[ATTR3]] 1367 // CHECK3-NEXT: call void @__kmpc_free_shared(i8* [[ARGC9]]) 1368 // CHECK3-NEXT: br label [[DOTTERMINATION_NOTIFIER:%.*]] 1369 // CHECK3: .termination.notifier: 1370 // CHECK3-NEXT: call void @__kmpc_kernel_deinit(i16 1) 1371 // CHECK3-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0) 1372 // CHECK3-NEXT: br label [[DOTEXIT]] 1373 // CHECK3: .exit: 1374 // CHECK3-NEXT: ret void 1375 // 1376 // 1377 // CHECK3-LABEL: define {{[^@]+}}@__omp_outlined__ 1378 // CHECK3-SAME: (i32* noalias [[DOTGLOBAL_TID_:%.*]], i32* noalias [[DOTBOUND_TID_:%.*]], i32* nonnull align 4 dereferenceable(4) [[ARGC:%.*]]) #[[ATTR1]] { 1379 // CHECK3-NEXT: entry: 1380 // CHECK3-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8 1381 // CHECK3-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8 1382 // CHECK3-NEXT: [[ARGC_ADDR:%.*]] = alloca i32*, align 8 1383 // CHECK3-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 8 1384 // CHECK3-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 8 1385 // CHECK3-NEXT: store i32* [[ARGC]], i32** [[ARGC_ADDR]], align 8 1386 // CHECK3-NEXT: [[TMP0:%.*]] = load i32*, i32** [[ARGC_ADDR]], align 8 1387 // CHECK3-NEXT: store i32 0, i32* [[TMP0]], align 4 1388 // CHECK3-NEXT: ret void 1389 // 1390 // 1391 // CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIPPcEiT__l53_worker 1392 // CHECK3-SAME: () #[[ATTR0]] { 1393 // CHECK3-NEXT: entry: 1394 // CHECK3-NEXT: [[WORK_FN:%.*]] = alloca i8*, align 8 1395 // CHECK3-NEXT: [[EXEC_STATUS:%.*]] = alloca i8, align 1 1396 // CHECK3-NEXT: store i8* null, i8** [[WORK_FN]], align 8 1397 // CHECK3-NEXT: store i8 0, i8* [[EXEC_STATUS]], align 1 1398 // CHECK3-NEXT: br label [[DOTAWAIT_WORK:%.*]] 1399 // CHECK3: .await.work: 1400 // CHECK3-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0) 1401 // CHECK3-NEXT: [[TMP0:%.*]] = call i1 @__kmpc_kernel_parallel(i8** [[WORK_FN]]) 1402 // CHECK3-NEXT: [[TMP1:%.*]] = zext i1 [[TMP0]] to i8 1403 // CHECK3-NEXT: store i8 [[TMP1]], i8* [[EXEC_STATUS]], align 1 1404 // CHECK3-NEXT: [[TMP2:%.*]] = load i8*, i8** [[WORK_FN]], align 8 1405 // CHECK3-NEXT: [[SHOULD_TERMINATE:%.*]] = icmp eq i8* [[TMP2]], null 1406 // CHECK3-NEXT: br i1 [[SHOULD_TERMINATE]], label [[DOTEXIT:%.*]], label [[DOTSELECT_WORKERS:%.*]] 1407 // CHECK3: .select.workers: 1408 // CHECK3-NEXT: [[TMP3:%.*]] = load i8, i8* [[EXEC_STATUS]], align 1 1409 // CHECK3-NEXT: [[IS_ACTIVE:%.*]] = icmp ne i8 [[TMP3]], 0 1410 // CHECK3-NEXT: br i1 [[IS_ACTIVE]], label [[DOTEXECUTE_PARALLEL:%.*]], label [[DOTBARRIER_PARALLEL:%.*]] 1411 // CHECK3: .execute.parallel: 1412 // CHECK3-NEXT: [[TMP4:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1]]) 1413 // CHECK3-NEXT: [[TMP5:%.*]] = bitcast i8* [[TMP2]] to void (i16, i32)* 1414 // CHECK3-NEXT: call void [[TMP5]](i16 0, i32 [[TMP4]]) 1415 // CHECK3-NEXT: br label [[DOTTERMINATE_PARALLEL:%.*]] 1416 // CHECK3: .terminate.parallel: 1417 // CHECK3-NEXT: call void @__kmpc_kernel_end_parallel() 1418 // CHECK3-NEXT: br label [[DOTBARRIER_PARALLEL]] 1419 // CHECK3: .barrier.parallel: 1420 // CHECK3-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0) 1421 // CHECK3-NEXT: br label [[DOTAWAIT_WORK]] 1422 // CHECK3: .exit: 1423 // CHECK3-NEXT: ret void 1424 // 1425 // 1426 // CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIPPcEiT__l53 1427 // CHECK3-SAME: (i64 [[A:%.*]], i64 [[B:%.*]], i8** [[ARGC:%.*]]) #[[ATTR1]] { 1428 // CHECK3-NEXT: entry: 1429 // CHECK3-NEXT: [[A_ADDR:%.*]] = alloca i64, align 8 1430 // CHECK3-NEXT: [[B_ADDR:%.*]] = alloca i64, align 8 1431 // CHECK3-NEXT: [[ARGC_ADDR:%.*]] = alloca i8**, align 8 1432 // CHECK3-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4 1433 // CHECK3-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 4 1434 // CHECK3-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4 1435 // CHECK3-NEXT: store i64 [[A]], i64* [[A_ADDR]], align 8 1436 // CHECK3-NEXT: store i64 [[B]], i64* [[B_ADDR]], align 8 1437 // CHECK3-NEXT: store i8** [[ARGC]], i8*** [[ARGC_ADDR]], align 8 1438 // CHECK3-NEXT: [[CONV:%.*]] = bitcast i64* [[A_ADDR]] to i32* 1439 // CHECK3-NEXT: [[CONV1:%.*]] = bitcast i64* [[B_ADDR]] to i32* 1440 // CHECK3-NEXT: [[NVPTX_TID:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x() 1441 // CHECK3-NEXT: [[NVPTX_NUM_THREADS:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() 1442 // CHECK3-NEXT: [[NVPTX_WARP_SIZE:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize() 1443 // CHECK3-NEXT: [[THREAD_LIMIT:%.*]] = sub nuw i32 [[NVPTX_NUM_THREADS]], [[NVPTX_WARP_SIZE]] 1444 // CHECK3-NEXT: [[TMP0:%.*]] = icmp ult i32 [[NVPTX_TID]], [[THREAD_LIMIT]] 1445 // CHECK3-NEXT: br i1 [[TMP0]], label [[DOTWORKER:%.*]], label [[DOTMASTERCHECK:%.*]] 1446 // CHECK3: .worker: 1447 // CHECK3-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIPPcEiT__l53_worker() #[[ATTR3]] 1448 // CHECK3-NEXT: br label [[DOTEXIT:%.*]] 1449 // CHECK3: .mastercheck: 1450 // CHECK3-NEXT: [[NVPTX_TID2:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x() 1451 // CHECK3-NEXT: [[NVPTX_NUM_THREADS3:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() 1452 // CHECK3-NEXT: [[NVPTX_WARP_SIZE4:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize() 1453 // CHECK3-NEXT: [[TMP1:%.*]] = sub nuw i32 [[NVPTX_WARP_SIZE4]], 1 1454 // CHECK3-NEXT: [[TMP2:%.*]] = sub nuw i32 [[NVPTX_NUM_THREADS3]], 1 1455 // CHECK3-NEXT: [[TMP3:%.*]] = xor i32 [[TMP1]], -1 1456 // CHECK3-NEXT: [[MASTER_TID:%.*]] = and i32 [[TMP2]], [[TMP3]] 1457 // CHECK3-NEXT: [[TMP4:%.*]] = icmp eq i32 [[NVPTX_TID2]], [[MASTER_TID]] 1458 // CHECK3-NEXT: br i1 [[TMP4]], label [[DOTMASTER:%.*]], label [[DOTEXIT]] 1459 // CHECK3: .master: 1460 // CHECK3-NEXT: [[NVPTX_NUM_THREADS5:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() 1461 // CHECK3-NEXT: [[NVPTX_WARP_SIZE6:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize() 1462 // CHECK3-NEXT: [[THREAD_LIMIT7:%.*]] = sub nuw i32 [[NVPTX_NUM_THREADS5]], [[NVPTX_WARP_SIZE6]] 1463 // CHECK3-NEXT: call void @__kmpc_kernel_init(i32 [[THREAD_LIMIT7]], i16 1) 1464 // CHECK3-NEXT: [[TMP5:%.*]] = load i8**, i8*** [[ARGC_ADDR]], align 8 1465 // CHECK3-NEXT: [[ARGC8:%.*]] = call i8* @__kmpc_alloc_shared(i64 8) 1466 // CHECK3-NEXT: [[ARGC_ON_STACK:%.*]] = bitcast i8* [[ARGC8]] to i8*** 1467 // CHECK3-NEXT: store i8** [[TMP5]], i8*** [[ARGC_ON_STACK]], align 8 1468 // CHECK3-NEXT: [[TMP6:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1]]) 1469 // CHECK3-NEXT: store i32 [[TMP6]], i32* [[DOTTHREADID_TEMP_]], align 4 1470 // CHECK3-NEXT: call void @__omp_outlined__1(i32* [[DOTTHREADID_TEMP_]], i32* [[DOTZERO_ADDR]], i8*** [[ARGC_ON_STACK]]) #[[ATTR3]] 1471 // CHECK3-NEXT: call void @__kmpc_free_shared(i8* [[ARGC8]]) 1472 // CHECK3-NEXT: br label [[DOTTERMINATION_NOTIFIER:%.*]] 1473 // CHECK3: .termination.notifier: 1474 // CHECK3-NEXT: call void @__kmpc_kernel_deinit(i16 1) 1475 // CHECK3-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0) 1476 // CHECK3-NEXT: br label [[DOTEXIT]] 1477 // CHECK3: .exit: 1478 // CHECK3-NEXT: ret void 1479 // 1480 // 1481 // CHECK3-LABEL: define {{[^@]+}}@__omp_outlined__1 1482 // CHECK3-SAME: (i32* noalias [[DOTGLOBAL_TID_:%.*]], i32* noalias [[DOTBOUND_TID_:%.*]], i8*** nonnull align 8 dereferenceable(8) [[ARGC:%.*]]) #[[ATTR1]] { 1483 // CHECK3-NEXT: entry: 1484 // CHECK3-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8 1485 // CHECK3-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8 1486 // CHECK3-NEXT: [[ARGC_ADDR:%.*]] = alloca i8***, align 8 1487 // CHECK3-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 8 1488 // CHECK3-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 8 1489 // CHECK3-NEXT: store i8*** [[ARGC]], i8**** [[ARGC_ADDR]], align 8 1490 // CHECK3-NEXT: [[TMP0:%.*]] = load i8***, i8**** [[ARGC_ADDR]], align 8 1491 // CHECK3-NEXT: store i8** null, i8*** [[TMP0]], align 8 1492 // CHECK3-NEXT: ret void 1493 // 1494 // 1495 // CHECK4-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l64_worker 1496 // CHECK4-SAME: () #[[ATTR0:[0-9]+]] { 1497 // CHECK4-NEXT: entry: 1498 // CHECK4-NEXT: [[WORK_FN:%.*]] = alloca i8*, align 4 1499 // CHECK4-NEXT: [[EXEC_STATUS:%.*]] = alloca i8, align 1 1500 // CHECK4-NEXT: store i8* null, i8** [[WORK_FN]], align 4 1501 // CHECK4-NEXT: store i8 0, i8* [[EXEC_STATUS]], align 1 1502 // CHECK4-NEXT: br label [[DOTAWAIT_WORK:%.*]] 1503 // CHECK4: .await.work: 1504 // CHECK4-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0) 1505 // CHECK4-NEXT: [[TMP0:%.*]] = call i1 @__kmpc_kernel_parallel(i8** [[WORK_FN]]) 1506 // CHECK4-NEXT: [[TMP1:%.*]] = zext i1 [[TMP0]] to i8 1507 // CHECK4-NEXT: store i8 [[TMP1]], i8* [[EXEC_STATUS]], align 1 1508 // CHECK4-NEXT: [[TMP2:%.*]] = load i8*, i8** [[WORK_FN]], align 4 1509 // CHECK4-NEXT: [[SHOULD_TERMINATE:%.*]] = icmp eq i8* [[TMP2]], null 1510 // CHECK4-NEXT: br i1 [[SHOULD_TERMINATE]], label [[DOTEXIT:%.*]], label [[DOTSELECT_WORKERS:%.*]] 1511 // CHECK4: .select.workers: 1512 // CHECK4-NEXT: [[TMP3:%.*]] = load i8, i8* [[EXEC_STATUS]], align 1 1513 // CHECK4-NEXT: [[IS_ACTIVE:%.*]] = icmp ne i8 [[TMP3]], 0 1514 // CHECK4-NEXT: br i1 [[IS_ACTIVE]], label [[DOTEXECUTE_PARALLEL:%.*]], label [[DOTBARRIER_PARALLEL:%.*]] 1515 // CHECK4: .execute.parallel: 1516 // CHECK4-NEXT: [[TMP4:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1:[0-9]+]]) 1517 // CHECK4-NEXT: [[TMP5:%.*]] = bitcast i8* [[TMP2]] to void (i16, i32)* 1518 // CHECK4-NEXT: call void [[TMP5]](i16 0, i32 [[TMP4]]) 1519 // CHECK4-NEXT: br label [[DOTTERMINATE_PARALLEL:%.*]] 1520 // CHECK4: .terminate.parallel: 1521 // CHECK4-NEXT: call void @__kmpc_kernel_end_parallel() 1522 // CHECK4-NEXT: br label [[DOTBARRIER_PARALLEL]] 1523 // CHECK4: .barrier.parallel: 1524 // CHECK4-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0) 1525 // CHECK4-NEXT: br label [[DOTAWAIT_WORK]] 1526 // CHECK4: .exit: 1527 // CHECK4-NEXT: ret void 1528 // 1529 // 1530 // CHECK4-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l64 1531 // CHECK4-SAME: (i32 [[A:%.*]], i32 [[B:%.*]], i32 [[ARGC:%.*]]) #[[ATTR1:[0-9]+]] { 1532 // CHECK4-NEXT: entry: 1533 // CHECK4-NEXT: [[A_ADDR:%.*]] = alloca i32, align 4 1534 // CHECK4-NEXT: [[B_ADDR:%.*]] = alloca i32, align 4 1535 // CHECK4-NEXT: [[ARGC_ADDR:%.*]] = alloca i32, align 4 1536 // CHECK4-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4 1537 // CHECK4-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 4 1538 // CHECK4-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4 1539 // CHECK4-NEXT: store i32 [[A]], i32* [[A_ADDR]], align 4 1540 // CHECK4-NEXT: store i32 [[B]], i32* [[B_ADDR]], align 4 1541 // CHECK4-NEXT: store i32 [[ARGC]], i32* [[ARGC_ADDR]], align 4 1542 // CHECK4-NEXT: [[NVPTX_TID:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x() 1543 // CHECK4-NEXT: [[NVPTX_NUM_THREADS:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() 1544 // CHECK4-NEXT: [[NVPTX_WARP_SIZE:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize() 1545 // CHECK4-NEXT: [[THREAD_LIMIT:%.*]] = sub nuw i32 [[NVPTX_NUM_THREADS]], [[NVPTX_WARP_SIZE]] 1546 // CHECK4-NEXT: [[TMP0:%.*]] = icmp ult i32 [[NVPTX_TID]], [[THREAD_LIMIT]] 1547 // CHECK4-NEXT: br i1 [[TMP0]], label [[DOTWORKER:%.*]], label [[DOTMASTERCHECK:%.*]] 1548 // CHECK4: .worker: 1549 // CHECK4-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l64_worker() #[[ATTR3:[0-9]+]] 1550 // CHECK4-NEXT: br label [[DOTEXIT:%.*]] 1551 // CHECK4: .mastercheck: 1552 // CHECK4-NEXT: [[NVPTX_TID1:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x() 1553 // CHECK4-NEXT: [[NVPTX_NUM_THREADS2:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() 1554 // CHECK4-NEXT: [[NVPTX_WARP_SIZE3:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize() 1555 // CHECK4-NEXT: [[TMP1:%.*]] = sub nuw i32 [[NVPTX_WARP_SIZE3]], 1 1556 // CHECK4-NEXT: [[TMP2:%.*]] = sub nuw i32 [[NVPTX_NUM_THREADS2]], 1 1557 // CHECK4-NEXT: [[TMP3:%.*]] = xor i32 [[TMP1]], -1 1558 // CHECK4-NEXT: [[MASTER_TID:%.*]] = and i32 [[TMP2]], [[TMP3]] 1559 // CHECK4-NEXT: [[TMP4:%.*]] = icmp eq i32 [[NVPTX_TID1]], [[MASTER_TID]] 1560 // CHECK4-NEXT: br i1 [[TMP4]], label [[DOTMASTER:%.*]], label [[DOTEXIT]] 1561 // CHECK4: .master: 1562 // CHECK4-NEXT: [[NVPTX_NUM_THREADS4:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() 1563 // CHECK4-NEXT: [[NVPTX_WARP_SIZE5:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize() 1564 // CHECK4-NEXT: [[THREAD_LIMIT6:%.*]] = sub nuw i32 [[NVPTX_NUM_THREADS4]], [[NVPTX_WARP_SIZE5]] 1565 // CHECK4-NEXT: call void @__kmpc_kernel_init(i32 [[THREAD_LIMIT6]], i16 1) 1566 // CHECK4-NEXT: [[TMP5:%.*]] = load i32, i32* [[ARGC_ADDR]], align 4 1567 // CHECK4-NEXT: [[ARGC7:%.*]] = call i8* @__kmpc_alloc_shared(i32 4) 1568 // CHECK4-NEXT: [[ARGC_ON_STACK:%.*]] = bitcast i8* [[ARGC7]] to i32* 1569 // CHECK4-NEXT: store i32 [[TMP5]], i32* [[ARGC_ON_STACK]], align 4 1570 // CHECK4-NEXT: [[TMP6:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1]]) 1571 // CHECK4-NEXT: store i32 [[TMP6]], i32* [[DOTTHREADID_TEMP_]], align 4 1572 // CHECK4-NEXT: call void @__omp_outlined__(i32* [[DOTTHREADID_TEMP_]], i32* [[DOTZERO_ADDR]], i32* [[ARGC_ON_STACK]]) #[[ATTR3]] 1573 // CHECK4-NEXT: call void @__kmpc_free_shared(i8* [[ARGC7]]) 1574 // CHECK4-NEXT: br label [[DOTTERMINATION_NOTIFIER:%.*]] 1575 // CHECK4: .termination.notifier: 1576 // CHECK4-NEXT: call void @__kmpc_kernel_deinit(i16 1) 1577 // CHECK4-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0) 1578 // CHECK4-NEXT: br label [[DOTEXIT]] 1579 // CHECK4: .exit: 1580 // CHECK4-NEXT: ret void 1581 // 1582 // 1583 // CHECK4-LABEL: define {{[^@]+}}@__omp_outlined__ 1584 // CHECK4-SAME: (i32* noalias [[DOTGLOBAL_TID_:%.*]], i32* noalias [[DOTBOUND_TID_:%.*]], i32* nonnull align 4 dereferenceable(4) [[ARGC:%.*]]) #[[ATTR1]] { 1585 // CHECK4-NEXT: entry: 1586 // CHECK4-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 4 1587 // CHECK4-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 4 1588 // CHECK4-NEXT: [[ARGC_ADDR:%.*]] = alloca i32*, align 4 1589 // CHECK4-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 4 1590 // CHECK4-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 4 1591 // CHECK4-NEXT: store i32* [[ARGC]], i32** [[ARGC_ADDR]], align 4 1592 // CHECK4-NEXT: [[TMP0:%.*]] = load i32*, i32** [[ARGC_ADDR]], align 4 1593 // CHECK4-NEXT: store i32 0, i32* [[TMP0]], align 4 1594 // CHECK4-NEXT: ret void 1595 // 1596 // 1597 // CHECK4-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIPPcEiT__l53_worker 1598 // CHECK4-SAME: () #[[ATTR0]] { 1599 // CHECK4-NEXT: entry: 1600 // CHECK4-NEXT: [[WORK_FN:%.*]] = alloca i8*, align 4 1601 // CHECK4-NEXT: [[EXEC_STATUS:%.*]] = alloca i8, align 1 1602 // CHECK4-NEXT: store i8* null, i8** [[WORK_FN]], align 4 1603 // CHECK4-NEXT: store i8 0, i8* [[EXEC_STATUS]], align 1 1604 // CHECK4-NEXT: br label [[DOTAWAIT_WORK:%.*]] 1605 // CHECK4: .await.work: 1606 // CHECK4-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0) 1607 // CHECK4-NEXT: [[TMP0:%.*]] = call i1 @__kmpc_kernel_parallel(i8** [[WORK_FN]]) 1608 // CHECK4-NEXT: [[TMP1:%.*]] = zext i1 [[TMP0]] to i8 1609 // CHECK4-NEXT: store i8 [[TMP1]], i8* [[EXEC_STATUS]], align 1 1610 // CHECK4-NEXT: [[TMP2:%.*]] = load i8*, i8** [[WORK_FN]], align 4 1611 // CHECK4-NEXT: [[SHOULD_TERMINATE:%.*]] = icmp eq i8* [[TMP2]], null 1612 // CHECK4-NEXT: br i1 [[SHOULD_TERMINATE]], label [[DOTEXIT:%.*]], label [[DOTSELECT_WORKERS:%.*]] 1613 // CHECK4: .select.workers: 1614 // CHECK4-NEXT: [[TMP3:%.*]] = load i8, i8* [[EXEC_STATUS]], align 1 1615 // CHECK4-NEXT: [[IS_ACTIVE:%.*]] = icmp ne i8 [[TMP3]], 0 1616 // CHECK4-NEXT: br i1 [[IS_ACTIVE]], label [[DOTEXECUTE_PARALLEL:%.*]], label [[DOTBARRIER_PARALLEL:%.*]] 1617 // CHECK4: .execute.parallel: 1618 // CHECK4-NEXT: [[TMP4:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1]]) 1619 // CHECK4-NEXT: [[TMP5:%.*]] = bitcast i8* [[TMP2]] to void (i16, i32)* 1620 // CHECK4-NEXT: call void [[TMP5]](i16 0, i32 [[TMP4]]) 1621 // CHECK4-NEXT: br label [[DOTTERMINATE_PARALLEL:%.*]] 1622 // CHECK4: .terminate.parallel: 1623 // CHECK4-NEXT: call void @__kmpc_kernel_end_parallel() 1624 // CHECK4-NEXT: br label [[DOTBARRIER_PARALLEL]] 1625 // CHECK4: .barrier.parallel: 1626 // CHECK4-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0) 1627 // CHECK4-NEXT: br label [[DOTAWAIT_WORK]] 1628 // CHECK4: .exit: 1629 // CHECK4-NEXT: ret void 1630 // 1631 // 1632 // CHECK4-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIPPcEiT__l53 1633 // CHECK4-SAME: (i32 [[A:%.*]], i32 [[B:%.*]], i8** [[ARGC:%.*]]) #[[ATTR1]] { 1634 // CHECK4-NEXT: entry: 1635 // CHECK4-NEXT: [[A_ADDR:%.*]] = alloca i32, align 4 1636 // CHECK4-NEXT: [[B_ADDR:%.*]] = alloca i32, align 4 1637 // CHECK4-NEXT: [[ARGC_ADDR:%.*]] = alloca i8**, align 4 1638 // CHECK4-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4 1639 // CHECK4-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 4 1640 // CHECK4-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4 1641 // CHECK4-NEXT: store i32 [[A]], i32* [[A_ADDR]], align 4 1642 // CHECK4-NEXT: store i32 [[B]], i32* [[B_ADDR]], align 4 1643 // CHECK4-NEXT: store i8** [[ARGC]], i8*** [[ARGC_ADDR]], align 4 1644 // CHECK4-NEXT: [[NVPTX_TID:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x() 1645 // CHECK4-NEXT: [[NVPTX_NUM_THREADS:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() 1646 // CHECK4-NEXT: [[NVPTX_WARP_SIZE:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize() 1647 // CHECK4-NEXT: [[THREAD_LIMIT:%.*]] = sub nuw i32 [[NVPTX_NUM_THREADS]], [[NVPTX_WARP_SIZE]] 1648 // CHECK4-NEXT: [[TMP0:%.*]] = icmp ult i32 [[NVPTX_TID]], [[THREAD_LIMIT]] 1649 // CHECK4-NEXT: br i1 [[TMP0]], label [[DOTWORKER:%.*]], label [[DOTMASTERCHECK:%.*]] 1650 // CHECK4: .worker: 1651 // CHECK4-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIPPcEiT__l53_worker() #[[ATTR3]] 1652 // CHECK4-NEXT: br label [[DOTEXIT:%.*]] 1653 // CHECK4: .mastercheck: 1654 // CHECK4-NEXT: [[NVPTX_TID1:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x() 1655 // CHECK4-NEXT: [[NVPTX_NUM_THREADS2:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() 1656 // CHECK4-NEXT: [[NVPTX_WARP_SIZE3:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize() 1657 // CHECK4-NEXT: [[TMP1:%.*]] = sub nuw i32 [[NVPTX_WARP_SIZE3]], 1 1658 // CHECK4-NEXT: [[TMP2:%.*]] = sub nuw i32 [[NVPTX_NUM_THREADS2]], 1 1659 // CHECK4-NEXT: [[TMP3:%.*]] = xor i32 [[TMP1]], -1 1660 // CHECK4-NEXT: [[MASTER_TID:%.*]] = and i32 [[TMP2]], [[TMP3]] 1661 // CHECK4-NEXT: [[TMP4:%.*]] = icmp eq i32 [[NVPTX_TID1]], [[MASTER_TID]] 1662 // CHECK4-NEXT: br i1 [[TMP4]], label [[DOTMASTER:%.*]], label [[DOTEXIT]] 1663 // CHECK4: .master: 1664 // CHECK4-NEXT: [[NVPTX_NUM_THREADS4:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() 1665 // CHECK4-NEXT: [[NVPTX_WARP_SIZE5:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize() 1666 // CHECK4-NEXT: [[THREAD_LIMIT6:%.*]] = sub nuw i32 [[NVPTX_NUM_THREADS4]], [[NVPTX_WARP_SIZE5]] 1667 // CHECK4-NEXT: call void @__kmpc_kernel_init(i32 [[THREAD_LIMIT6]], i16 1) 1668 // CHECK4-NEXT: [[TMP5:%.*]] = load i8**, i8*** [[ARGC_ADDR]], align 4 1669 // CHECK4-NEXT: [[ARGC7:%.*]] = call i8* @__kmpc_alloc_shared(i32 4) 1670 // CHECK4-NEXT: [[ARGC_ON_STACK:%.*]] = bitcast i8* [[ARGC7]] to i8*** 1671 // CHECK4-NEXT: store i8** [[TMP5]], i8*** [[ARGC_ON_STACK]], align 4 1672 // CHECK4-NEXT: [[TMP6:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1]]) 1673 // CHECK4-NEXT: store i32 [[TMP6]], i32* [[DOTTHREADID_TEMP_]], align 4 1674 // CHECK4-NEXT: call void @__omp_outlined__1(i32* [[DOTTHREADID_TEMP_]], i32* [[DOTZERO_ADDR]], i8*** [[ARGC_ON_STACK]]) #[[ATTR3]] 1675 // CHECK4-NEXT: call void @__kmpc_free_shared(i8* [[ARGC7]]) 1676 // CHECK4-NEXT: br label [[DOTTERMINATION_NOTIFIER:%.*]] 1677 // CHECK4: .termination.notifier: 1678 // CHECK4-NEXT: call void @__kmpc_kernel_deinit(i16 1) 1679 // CHECK4-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0) 1680 // CHECK4-NEXT: br label [[DOTEXIT]] 1681 // CHECK4: .exit: 1682 // CHECK4-NEXT: ret void 1683 // 1684 // 1685 // CHECK4-LABEL: define {{[^@]+}}@__omp_outlined__1 1686 // CHECK4-SAME: (i32* noalias [[DOTGLOBAL_TID_:%.*]], i32* noalias [[DOTBOUND_TID_:%.*]], i8*** nonnull align 4 dereferenceable(4) [[ARGC:%.*]]) #[[ATTR1]] { 1687 // CHECK4-NEXT: entry: 1688 // CHECK4-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 4 1689 // CHECK4-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 4 1690 // CHECK4-NEXT: [[ARGC_ADDR:%.*]] = alloca i8***, align 4 1691 // CHECK4-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 4 1692 // CHECK4-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 4 1693 // CHECK4-NEXT: store i8*** [[ARGC]], i8**** [[ARGC_ADDR]], align 4 1694 // CHECK4-NEXT: [[TMP0:%.*]] = load i8***, i8**** [[ARGC_ADDR]], align 4 1695 // CHECK4-NEXT: store i8** null, i8*** [[TMP0]], align 4 1696 // CHECK4-NEXT: ret void 1697 // 1698