367 lines · cpp
1// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --function-signature --include-generated-funcs --replace-value-regex "__omp_offloading_[0-9a-z]+_[0-9a-z]+" "reduction_size[.].+[.]" "pl_cond[.].+[.|,]" --prefix-filecheck-ir-name _2// Test target codegen - host bc file has to be created first.3// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc4// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple nvptx64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --check-prefix=CHECK15// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm-bc %s -o %t-x86-host.bc6// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple nvptx-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s --check-prefix=CHECK27// RUN: %clang_cc1 -verify -fopenmp -fexceptions -fcxx-exceptions -x c++ -triple nvptx-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s --check-prefix=CHECK28// expected-no-diagnostics9#ifndef HEADER10#define HEADER11 12template<typename tx>13tx ftemplate(int n) {14 tx a = 0;15 short aa = 0;16 tx b[10];17 18 #pragma omp target teams if(0)19 {20 b[2] += 1;21 }22 23 #pragma omp target teams if(1)24 {25 a = '1';26 }27 28 #pragma omp target teams if(n>40)29 {30 aa = 1;31 }32 33 #pragma omp target teams34 {35#pragma omp parallel36#pragma omp parallel37 aa = 1;38 }39 40 return a;41}42 43int bar(int n){44 int a = 0;45 46 a += ftemplate<char>(n);47 48 return a;49}50 51#endif52// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l2353// CHECK1-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], i64 noundef [[A:%.*]]) #[[ATTR0:[0-9]+]] {54// CHECK1-NEXT: entry:55// CHECK1-NEXT: [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 856// CHECK1-NEXT: [[A_ADDR:%.*]] = alloca i64, align 857// CHECK1-NEXT: [[A_CASTED:%.*]] = alloca i64, align 858// CHECK1-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 459// CHECK1-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 460// CHECK1-NEXT: store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 861// CHECK1-NEXT: store i64 [[A]], ptr [[A_ADDR]], align 862// CHECK1-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_target_init(ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l23_kernel_environment, ptr [[DYN_PTR]])63// CHECK1-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP0]], -164// CHECK1-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]65// CHECK1: user_code.entry:66// CHECK1-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1:[0-9]+]])67// CHECK1-NEXT: [[TMP2:%.*]] = load i8, ptr [[A_ADDR]], align 168// CHECK1-NEXT: store i8 [[TMP2]], ptr [[A_CASTED]], align 169// CHECK1-NEXT: [[TMP3:%.*]] = load i64, ptr [[A_CASTED]], align 870// CHECK1-NEXT: store i32 0, ptr [[DOTZERO_ADDR]], align 471// CHECK1-NEXT: store i32 [[TMP1]], ptr [[DOTTHREADID_TEMP_]], align 472// CHECK1-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l23_omp_outlined(ptr [[DOTTHREADID_TEMP_]], ptr [[DOTZERO_ADDR]], i64 [[TMP3]]) #[[ATTR2:[0-9]+]]73// CHECK1-NEXT: call void @__kmpc_target_deinit()74// CHECK1-NEXT: ret void75// CHECK1: worker.exit:76// CHECK1-NEXT: ret void77//78//79// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l23_omp_outlined80// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i64 noundef [[A:%.*]]) #[[ATTR1:[0-9]+]] {81// CHECK1-NEXT: entry:82// CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 883// CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 884// CHECK1-NEXT: [[A_ADDR:%.*]] = alloca i64, align 885// CHECK1-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 886// CHECK1-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 887// CHECK1-NEXT: store i64 [[A]], ptr [[A_ADDR]], align 888// CHECK1-NEXT: store i8 49, ptr [[A_ADDR]], align 189// CHECK1-NEXT: ret void90//91//92// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l2893// CHECK1-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], i64 noundef [[AA:%.*]]) #[[ATTR0]] {94// CHECK1-NEXT: entry:95// CHECK1-NEXT: [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 896// CHECK1-NEXT: [[AA_ADDR:%.*]] = alloca i64, align 897// CHECK1-NEXT: [[AA_CASTED:%.*]] = alloca i64, align 898// CHECK1-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 499// CHECK1-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 4100// CHECK1-NEXT: store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 8101// CHECK1-NEXT: store i64 [[AA]], ptr [[AA_ADDR]], align 8102// CHECK1-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_target_init(ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l28_kernel_environment, ptr [[DYN_PTR]])103// CHECK1-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP0]], -1104// CHECK1-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]105// CHECK1: user_code.entry:106// CHECK1-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])107// CHECK1-NEXT: [[TMP2:%.*]] = load i16, ptr [[AA_ADDR]], align 2108// CHECK1-NEXT: store i16 [[TMP2]], ptr [[AA_CASTED]], align 2109// CHECK1-NEXT: [[TMP3:%.*]] = load i64, ptr [[AA_CASTED]], align 8110// CHECK1-NEXT: store i32 0, ptr [[DOTZERO_ADDR]], align 4111// CHECK1-NEXT: store i32 [[TMP1]], ptr [[DOTTHREADID_TEMP_]], align 4112// CHECK1-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l28_omp_outlined(ptr [[DOTTHREADID_TEMP_]], ptr [[DOTZERO_ADDR]], i64 [[TMP3]]) #[[ATTR2]]113// CHECK1-NEXT: call void @__kmpc_target_deinit()114// CHECK1-NEXT: ret void115// CHECK1: worker.exit:116// CHECK1-NEXT: ret void117//118//119// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l28_omp_outlined120// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i64 noundef [[AA:%.*]]) #[[ATTR1]] {121// CHECK1-NEXT: entry:122// CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8123// CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8124// CHECK1-NEXT: [[AA_ADDR:%.*]] = alloca i64, align 8125// CHECK1-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8126// CHECK1-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8127// CHECK1-NEXT: store i64 [[AA]], ptr [[AA_ADDR]], align 8128// CHECK1-NEXT: store i16 1, ptr [[AA_ADDR]], align 2129// CHECK1-NEXT: ret void130//131//132// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l33133// CHECK1-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], i64 noundef [[AA:%.*]]) #[[ATTR0]] {134// CHECK1-NEXT: entry:135// CHECK1-NEXT: [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 8136// CHECK1-NEXT: [[AA_ADDR:%.*]] = alloca i64, align 8137// CHECK1-NEXT: [[AA_CASTED:%.*]] = alloca i64, align 8138// CHECK1-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4139// CHECK1-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 4140// CHECK1-NEXT: store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 8141// CHECK1-NEXT: store i64 [[AA]], ptr [[AA_ADDR]], align 8142// CHECK1-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_target_init(ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l33_kernel_environment, ptr [[DYN_PTR]])143// CHECK1-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP0]], -1144// CHECK1-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]145// CHECK1: user_code.entry:146// CHECK1-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])147// CHECK1-NEXT: [[TMP2:%.*]] = load i16, ptr [[AA_ADDR]], align 2148// CHECK1-NEXT: store i16 [[TMP2]], ptr [[AA_CASTED]], align 2149// CHECK1-NEXT: [[TMP3:%.*]] = load i64, ptr [[AA_CASTED]], align 8150// CHECK1-NEXT: store i32 0, ptr [[DOTZERO_ADDR]], align 4151// CHECK1-NEXT: store i32 [[TMP1]], ptr [[DOTTHREADID_TEMP_]], align 4152// CHECK1-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l33_omp_outlined(ptr [[DOTTHREADID_TEMP_]], ptr [[DOTZERO_ADDR]], i64 [[TMP3]]) #[[ATTR2]]153// CHECK1-NEXT: call void @__kmpc_target_deinit()154// CHECK1-NEXT: ret void155// CHECK1: worker.exit:156// CHECK1-NEXT: ret void157//158//159// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l33_omp_outlined160// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i64 noundef [[AA:%.*]]) #[[ATTR1]] {161// CHECK1-NEXT: entry:162// CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8163// CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8164// CHECK1-NEXT: [[AA_ADDR:%.*]] = alloca i64, align 8165// CHECK1-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [1 x ptr], align 8166// CHECK1-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8167// CHECK1-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8168// CHECK1-NEXT: store i64 [[AA]], ptr [[AA_ADDR]], align 8169// CHECK1-NEXT: [[TMP0:%.*]] = getelementptr inbounds [1 x ptr], ptr [[CAPTURED_VARS_ADDRS]], i64 0, i64 0170// CHECK1-NEXT: store ptr [[AA_ADDR]], ptr [[TMP0]], align 8171// CHECK1-NEXT: [[TMP1:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8172// CHECK1-NEXT: [[TMP2:%.*]] = load i32, ptr [[TMP1]], align 4173// CHECK1-NEXT: call void @__kmpc_parallel_51(ptr @[[GLOB1]], i32 [[TMP2]], i32 1, i32 -1, i32 -1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l33_omp_outlined_omp_outlined, ptr null, ptr [[CAPTURED_VARS_ADDRS]], i64 1)174// CHECK1-NEXT: ret void175//176//177// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l33_omp_outlined_omp_outlined178// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 2 dereferenceable(2) [[AA:%.*]]) #[[ATTR1]] {179// CHECK1-NEXT: entry:180// CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8181// CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8182// CHECK1-NEXT: [[AA_ADDR:%.*]] = alloca ptr, align 8183// CHECK1-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [1 x ptr], align 8184// CHECK1-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8185// CHECK1-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8186// CHECK1-NEXT: store ptr [[AA]], ptr [[AA_ADDR]], align 8187// CHECK1-NEXT: [[TMP0:%.*]] = load ptr, ptr [[AA_ADDR]], align 8188// CHECK1-NEXT: [[TMP1:%.*]] = getelementptr inbounds [1 x ptr], ptr [[CAPTURED_VARS_ADDRS]], i64 0, i64 0189// CHECK1-NEXT: store ptr [[TMP0]], ptr [[TMP1]], align 8190// CHECK1-NEXT: [[TMP2:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8191// CHECK1-NEXT: [[TMP3:%.*]] = load i32, ptr [[TMP2]], align 4192// CHECK1-NEXT: call void @__kmpc_parallel_51(ptr @[[GLOB1]], i32 [[TMP3]], i32 1, i32 -1, i32 -1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l33_omp_outlined_omp_outlined_omp_outlined, ptr null, ptr [[CAPTURED_VARS_ADDRS]], i64 1)193// CHECK1-NEXT: ret void194//195//196// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l33_omp_outlined_omp_outlined_omp_outlined197// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 2 dereferenceable(2) [[AA:%.*]]) #[[ATTR1]] {198// CHECK1-NEXT: entry:199// CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8200// CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8201// CHECK1-NEXT: [[AA_ADDR:%.*]] = alloca ptr, align 8202// CHECK1-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8203// CHECK1-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8204// CHECK1-NEXT: store ptr [[AA]], ptr [[AA_ADDR]], align 8205// CHECK1-NEXT: [[TMP0:%.*]] = load ptr, ptr [[AA_ADDR]], align 8206// CHECK1-NEXT: store i16 1, ptr [[TMP0]], align 2207// CHECK1-NEXT: ret void208//209//210// CHECK2-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l23211// CHECK2-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], i32 noundef [[A:%.*]]) #[[ATTR0:[0-9]+]] {212// CHECK2-NEXT: entry:213// CHECK2-NEXT: [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 4214// CHECK2-NEXT: [[A_ADDR:%.*]] = alloca i32, align 4215// CHECK2-NEXT: [[A_CASTED:%.*]] = alloca i32, align 4216// CHECK2-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4217// CHECK2-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 4218// CHECK2-NEXT: store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 4219// CHECK2-NEXT: store i32 [[A]], ptr [[A_ADDR]], align 4220// CHECK2-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_target_init(ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l23_kernel_environment, ptr [[DYN_PTR]])221// CHECK2-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP0]], -1222// CHECK2-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]223// CHECK2: user_code.entry:224// CHECK2-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1:[0-9]+]])225// CHECK2-NEXT: [[TMP2:%.*]] = load i8, ptr [[A_ADDR]], align 1226// CHECK2-NEXT: store i8 [[TMP2]], ptr [[A_CASTED]], align 1227// CHECK2-NEXT: [[TMP3:%.*]] = load i32, ptr [[A_CASTED]], align 4228// CHECK2-NEXT: store i32 0, ptr [[DOTZERO_ADDR]], align 4229// CHECK2-NEXT: store i32 [[TMP1]], ptr [[DOTTHREADID_TEMP_]], align 4230// CHECK2-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l23_omp_outlined(ptr [[DOTTHREADID_TEMP_]], ptr [[DOTZERO_ADDR]], i32 [[TMP3]]) #[[ATTR2:[0-9]+]]231// CHECK2-NEXT: call void @__kmpc_target_deinit()232// CHECK2-NEXT: ret void233// CHECK2: worker.exit:234// CHECK2-NEXT: ret void235//236//237// CHECK2-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l23_omp_outlined238// CHECK2-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i32 noundef [[A:%.*]]) #[[ATTR1:[0-9]+]] {239// CHECK2-NEXT: entry:240// CHECK2-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 4241// CHECK2-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 4242// CHECK2-NEXT: [[A_ADDR:%.*]] = alloca i32, align 4243// CHECK2-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 4244// CHECK2-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 4245// CHECK2-NEXT: store i32 [[A]], ptr [[A_ADDR]], align 4246// CHECK2-NEXT: store i8 49, ptr [[A_ADDR]], align 1247// CHECK2-NEXT: ret void248//249//250// CHECK2-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l28251// CHECK2-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], i32 noundef [[AA:%.*]]) #[[ATTR0]] {252// CHECK2-NEXT: entry:253// CHECK2-NEXT: [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 4254// CHECK2-NEXT: [[AA_ADDR:%.*]] = alloca i32, align 4255// CHECK2-NEXT: [[AA_CASTED:%.*]] = alloca i32, align 4256// CHECK2-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4257// CHECK2-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 4258// CHECK2-NEXT: store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 4259// CHECK2-NEXT: store i32 [[AA]], ptr [[AA_ADDR]], align 4260// CHECK2-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_target_init(ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l28_kernel_environment, ptr [[DYN_PTR]])261// CHECK2-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP0]], -1262// CHECK2-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]263// CHECK2: user_code.entry:264// CHECK2-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])265// CHECK2-NEXT: [[TMP2:%.*]] = load i16, ptr [[AA_ADDR]], align 2266// CHECK2-NEXT: store i16 [[TMP2]], ptr [[AA_CASTED]], align 2267// CHECK2-NEXT: [[TMP3:%.*]] = load i32, ptr [[AA_CASTED]], align 4268// CHECK2-NEXT: store i32 0, ptr [[DOTZERO_ADDR]], align 4269// CHECK2-NEXT: store i32 [[TMP1]], ptr [[DOTTHREADID_TEMP_]], align 4270// CHECK2-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l28_omp_outlined(ptr [[DOTTHREADID_TEMP_]], ptr [[DOTZERO_ADDR]], i32 [[TMP3]]) #[[ATTR2]]271// CHECK2-NEXT: call void @__kmpc_target_deinit()272// CHECK2-NEXT: ret void273// CHECK2: worker.exit:274// CHECK2-NEXT: ret void275//276//277// CHECK2-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l28_omp_outlined278// CHECK2-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i32 noundef [[AA:%.*]]) #[[ATTR1]] {279// CHECK2-NEXT: entry:280// CHECK2-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 4281// CHECK2-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 4282// CHECK2-NEXT: [[AA_ADDR:%.*]] = alloca i32, align 4283// CHECK2-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 4284// CHECK2-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 4285// CHECK2-NEXT: store i32 [[AA]], ptr [[AA_ADDR]], align 4286// CHECK2-NEXT: store i16 1, ptr [[AA_ADDR]], align 2287// CHECK2-NEXT: ret void288//289//290// CHECK2-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l33291// CHECK2-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], i32 noundef [[AA:%.*]]) #[[ATTR0]] {292// CHECK2-NEXT: entry:293// CHECK2-NEXT: [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 4294// CHECK2-NEXT: [[AA_ADDR:%.*]] = alloca i32, align 4295// CHECK2-NEXT: [[AA_CASTED:%.*]] = alloca i32, align 4296// CHECK2-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4297// CHECK2-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 4298// CHECK2-NEXT: store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 4299// CHECK2-NEXT: store i32 [[AA]], ptr [[AA_ADDR]], align 4300// CHECK2-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_target_init(ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l33_kernel_environment, ptr [[DYN_PTR]])301// CHECK2-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP0]], -1302// CHECK2-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]303// CHECK2: user_code.entry:304// CHECK2-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])305// CHECK2-NEXT: [[TMP2:%.*]] = load i16, ptr [[AA_ADDR]], align 2306// CHECK2-NEXT: store i16 [[TMP2]], ptr [[AA_CASTED]], align 2307// CHECK2-NEXT: [[TMP3:%.*]] = load i32, ptr [[AA_CASTED]], align 4308// CHECK2-NEXT: store i32 0, ptr [[DOTZERO_ADDR]], align 4309// CHECK2-NEXT: store i32 [[TMP1]], ptr [[DOTTHREADID_TEMP_]], align 4310// CHECK2-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l33_omp_outlined(ptr [[DOTTHREADID_TEMP_]], ptr [[DOTZERO_ADDR]], i32 [[TMP3]]) #[[ATTR2]]311// CHECK2-NEXT: call void @__kmpc_target_deinit()312// CHECK2-NEXT: ret void313// CHECK2: worker.exit:314// CHECK2-NEXT: ret void315//316//317// CHECK2-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l33_omp_outlined318// CHECK2-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i32 noundef [[AA:%.*]]) #[[ATTR1]] {319// CHECK2-NEXT: entry:320// CHECK2-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 4321// CHECK2-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 4322// CHECK2-NEXT: [[AA_ADDR:%.*]] = alloca i32, align 4323// CHECK2-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [1 x ptr], align 4324// CHECK2-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 4325// CHECK2-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 4326// CHECK2-NEXT: store i32 [[AA]], ptr [[AA_ADDR]], align 4327// CHECK2-NEXT: [[TMP0:%.*]] = getelementptr inbounds [1 x ptr], ptr [[CAPTURED_VARS_ADDRS]], i32 0, i32 0328// CHECK2-NEXT: store ptr [[AA_ADDR]], ptr [[TMP0]], align 4329// CHECK2-NEXT: [[TMP1:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 4330// CHECK2-NEXT: [[TMP2:%.*]] = load i32, ptr [[TMP1]], align 4331// CHECK2-NEXT: call void @__kmpc_parallel_51(ptr @[[GLOB1]], i32 [[TMP2]], i32 1, i32 -1, i32 -1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l33_omp_outlined_omp_outlined, ptr null, ptr [[CAPTURED_VARS_ADDRS]], i32 1)332// CHECK2-NEXT: ret void333//334//335// CHECK2-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l33_omp_outlined_omp_outlined336// CHECK2-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 2 dereferenceable(2) [[AA:%.*]]) #[[ATTR1]] {337// CHECK2-NEXT: entry:338// CHECK2-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 4339// CHECK2-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 4340// CHECK2-NEXT: [[AA_ADDR:%.*]] = alloca ptr, align 4341// CHECK2-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [1 x ptr], align 4342// CHECK2-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 4343// CHECK2-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 4344// CHECK2-NEXT: store ptr [[AA]], ptr [[AA_ADDR]], align 4345// CHECK2-NEXT: [[TMP0:%.*]] = load ptr, ptr [[AA_ADDR]], align 4346// CHECK2-NEXT: [[TMP1:%.*]] = getelementptr inbounds [1 x ptr], ptr [[CAPTURED_VARS_ADDRS]], i32 0, i32 0347// CHECK2-NEXT: store ptr [[TMP0]], ptr [[TMP1]], align 4348// CHECK2-NEXT: [[TMP2:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 4349// CHECK2-NEXT: [[TMP3:%.*]] = load i32, ptr [[TMP2]], align 4350// CHECK2-NEXT: call void @__kmpc_parallel_51(ptr @[[GLOB1]], i32 [[TMP3]], i32 1, i32 -1, i32 -1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l33_omp_outlined_omp_outlined_omp_outlined, ptr null, ptr [[CAPTURED_VARS_ADDRS]], i32 1)351// CHECK2-NEXT: ret void352//353//354// CHECK2-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l33_omp_outlined_omp_outlined_omp_outlined355// CHECK2-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 2 dereferenceable(2) [[AA:%.*]]) #[[ATTR1]] {356// CHECK2-NEXT: entry:357// CHECK2-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 4358// CHECK2-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 4359// CHECK2-NEXT: [[AA_ADDR:%.*]] = alloca ptr, align 4360// CHECK2-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 4361// CHECK2-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 4362// CHECK2-NEXT: store ptr [[AA]], ptr [[AA_ADDR]], align 4363// CHECK2-NEXT: [[TMP0:%.*]] = load ptr, ptr [[AA_ADDR]], align 4364// CHECK2-NEXT: store i16 1, ptr [[TMP0]], align 2365// CHECK2-NEXT: ret void366//367