422 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 -DCK1 -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 -DCK1 -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 -DCK1 -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 -DCK1 -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// expected-no-diagnostics8#ifndef HEADER9#define HEADER10 11#ifdef CK112 13template <typename T>14int tmain(T argc) {15#pragma omp target16#pragma omp teams17 argc = 0;18 return 0;19}20 21 22int main (int argc, char **argv) {23#pragma omp target24#pragma omp teams25 {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_teams33 34 35// target region in template36 37 38 39#endif // CK140 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.bc43// RUN: %clang_cc1 -DCK2 -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=CHECK344// 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.bc45// RUN: %clang_cc1 -DCK2 -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=CHECK446// expected-no-diagnostics47#ifdef CK248 49template <typename T>50int tmain(T argc) {51 int a = 10;52 int b = 5;53#pragma omp target54#pragma omp teams num_teams(a) thread_limit(b)55 {56 argc = 0;57 }58 return 0;59}60 61int main (int argc, char **argv) {62 int a = 20;63 int b = 5;64#pragma omp target65#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 // CK278#endif79// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l2380// CHECK1-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], i64 noundef [[ARGC:%.*]]) #[[ATTR0:[0-9]+]] {81// CHECK1-NEXT: entry:82// CHECK1-NEXT: [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 883// CHECK1-NEXT: [[ARGC_ADDR:%.*]] = alloca i64, align 884// CHECK1-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 485// CHECK1-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 486// CHECK1-NEXT: store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 887// CHECK1-NEXT: store i64 [[ARGC]], ptr [[ARGC_ADDR]], align 888// CHECK1-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_target_init(ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l23_kernel_environment, ptr [[DYN_PTR]])89// CHECK1-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP0]], -190// CHECK1-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]91// CHECK1: user_code.entry:92// CHECK1-NEXT: [[TMP1:%.*]] = load i32, ptr [[ARGC_ADDR]], align 493// CHECK1-NEXT: [[ARGC1:%.*]] = call align 8 ptr @__kmpc_alloc_shared(i64 4)94// CHECK1-NEXT: store i32 [[TMP1]], ptr [[ARGC1]], align 495// CHECK1-NEXT: [[TMP2:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1:[0-9]+]])96// CHECK1-NEXT: store i32 0, ptr [[DOTZERO_ADDR]], align 497// CHECK1-NEXT: store i32 [[TMP2]], ptr [[DOTTHREADID_TEMP_]], align 498// CHECK1-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l23_omp_outlined(ptr [[DOTTHREADID_TEMP_]], ptr [[DOTZERO_ADDR]], ptr [[ARGC1]]) #[[ATTR3:[0-9]+]]99// CHECK1-NEXT: call void @__kmpc_free_shared(ptr [[ARGC1]], i64 4)100// CHECK1-NEXT: call void @__kmpc_target_deinit()101// CHECK1-NEXT: ret void102// CHECK1: worker.exit:103// CHECK1-NEXT: ret void104//105//106// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l23_omp_outlined107// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[ARGC:%.*]]) #[[ATTR2:[0-9]+]] {108// CHECK1-NEXT: entry:109// CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8110// CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8111// CHECK1-NEXT: [[ARGC_ADDR:%.*]] = alloca ptr, align 8112// CHECK1-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8113// CHECK1-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8114// CHECK1-NEXT: store ptr [[ARGC]], ptr [[ARGC_ADDR]], align 8115// CHECK1-NEXT: [[TMP0:%.*]] = load ptr, ptr [[ARGC_ADDR]], align 8116// CHECK1-NEXT: store i32 0, ptr [[TMP0]], align 4117// CHECK1-NEXT: ret void118//119//120// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIPPcEiT__l15121// CHECK1-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], ptr noundef [[ARGC:%.*]]) #[[ATTR0]] {122// CHECK1-NEXT: entry:123// CHECK1-NEXT: [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 8124// CHECK1-NEXT: [[ARGC_ADDR:%.*]] = alloca ptr, align 8125// CHECK1-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4126// CHECK1-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 4127// CHECK1-NEXT: store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 8128// CHECK1-NEXT: store ptr [[ARGC]], ptr [[ARGC_ADDR]], align 8129// CHECK1-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_target_init(ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIPPcEiT__l15_kernel_environment, ptr [[DYN_PTR]])130// CHECK1-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP0]], -1131// CHECK1-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]132// CHECK1: user_code.entry:133// CHECK1-NEXT: [[TMP1:%.*]] = load ptr, ptr [[ARGC_ADDR]], align 8134// CHECK1-NEXT: [[ARGC1:%.*]] = call align 8 ptr @__kmpc_alloc_shared(i64 8)135// CHECK1-NEXT: store ptr [[TMP1]], ptr [[ARGC1]], align 8136// CHECK1-NEXT: [[TMP2:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])137// CHECK1-NEXT: store i32 0, ptr [[DOTZERO_ADDR]], align 4138// CHECK1-NEXT: store i32 [[TMP2]], ptr [[DOTTHREADID_TEMP_]], align 4139// CHECK1-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIPPcEiT__l15_omp_outlined(ptr [[DOTTHREADID_TEMP_]], ptr [[DOTZERO_ADDR]], ptr [[ARGC1]]) #[[ATTR3]]140// CHECK1-NEXT: call void @__kmpc_free_shared(ptr [[ARGC1]], i64 8)141// CHECK1-NEXT: call void @__kmpc_target_deinit()142// CHECK1-NEXT: ret void143// CHECK1: worker.exit:144// CHECK1-NEXT: ret void145//146//147// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIPPcEiT__l15_omp_outlined148// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 8 dereferenceable(8) [[ARGC:%.*]]) #[[ATTR2]] {149// CHECK1-NEXT: entry:150// CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8151// CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8152// CHECK1-NEXT: [[ARGC_ADDR:%.*]] = alloca ptr, align 8153// CHECK1-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8154// CHECK1-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8155// CHECK1-NEXT: store ptr [[ARGC]], ptr [[ARGC_ADDR]], align 8156// CHECK1-NEXT: [[TMP0:%.*]] = load ptr, ptr [[ARGC_ADDR]], align 8157// CHECK1-NEXT: store ptr null, ptr [[TMP0]], align 8158// CHECK1-NEXT: ret void159//160//161// CHECK2-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l23162// CHECK2-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], i32 noundef [[ARGC:%.*]]) #[[ATTR0:[0-9]+]] {163// CHECK2-NEXT: entry:164// CHECK2-NEXT: [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 4165// CHECK2-NEXT: [[ARGC_ADDR:%.*]] = alloca i32, align 4166// CHECK2-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4167// CHECK2-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 4168// CHECK2-NEXT: store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 4169// CHECK2-NEXT: store i32 [[ARGC]], ptr [[ARGC_ADDR]], align 4170// CHECK2-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_target_init(ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l23_kernel_environment, ptr [[DYN_PTR]])171// CHECK2-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP0]], -1172// CHECK2-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]173// CHECK2: user_code.entry:174// CHECK2-NEXT: [[TMP1:%.*]] = load i32, ptr [[ARGC_ADDR]], align 4175// CHECK2-NEXT: [[ARGC1:%.*]] = call align 8 ptr @__kmpc_alloc_shared(i32 4)176// CHECK2-NEXT: store i32 [[TMP1]], ptr [[ARGC1]], align 4177// CHECK2-NEXT: [[TMP2:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1:[0-9]+]])178// CHECK2-NEXT: store i32 0, ptr [[DOTZERO_ADDR]], align 4179// CHECK2-NEXT: store i32 [[TMP2]], ptr [[DOTTHREADID_TEMP_]], align 4180// CHECK2-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l23_omp_outlined(ptr [[DOTTHREADID_TEMP_]], ptr [[DOTZERO_ADDR]], ptr [[ARGC1]]) #[[ATTR3:[0-9]+]]181// CHECK2-NEXT: call void @__kmpc_free_shared(ptr [[ARGC1]], i32 4)182// CHECK2-NEXT: call void @__kmpc_target_deinit()183// CHECK2-NEXT: ret void184// CHECK2: worker.exit:185// CHECK2-NEXT: ret void186//187//188// CHECK2-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l23_omp_outlined189// CHECK2-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[ARGC:%.*]]) #[[ATTR2:[0-9]+]] {190// CHECK2-NEXT: entry:191// CHECK2-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 4192// CHECK2-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 4193// CHECK2-NEXT: [[ARGC_ADDR:%.*]] = alloca ptr, align 4194// CHECK2-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 4195// CHECK2-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 4196// CHECK2-NEXT: store ptr [[ARGC]], ptr [[ARGC_ADDR]], align 4197// CHECK2-NEXT: [[TMP0:%.*]] = load ptr, ptr [[ARGC_ADDR]], align 4198// CHECK2-NEXT: store i32 0, ptr [[TMP0]], align 4199// CHECK2-NEXT: ret void200//201//202// CHECK2-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIPPcEiT__l15203// CHECK2-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], ptr noundef [[ARGC:%.*]]) #[[ATTR0]] {204// CHECK2-NEXT: entry:205// CHECK2-NEXT: [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 4206// CHECK2-NEXT: [[ARGC_ADDR:%.*]] = alloca ptr, align 4207// CHECK2-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4208// CHECK2-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 4209// CHECK2-NEXT: store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 4210// CHECK2-NEXT: store ptr [[ARGC]], ptr [[ARGC_ADDR]], align 4211// CHECK2-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_target_init(ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIPPcEiT__l15_kernel_environment, ptr [[DYN_PTR]])212// CHECK2-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP0]], -1213// CHECK2-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]214// CHECK2: user_code.entry:215// CHECK2-NEXT: [[TMP1:%.*]] = load ptr, ptr [[ARGC_ADDR]], align 4216// CHECK2-NEXT: [[ARGC1:%.*]] = call align 8 ptr @__kmpc_alloc_shared(i32 4)217// CHECK2-NEXT: store ptr [[TMP1]], ptr [[ARGC1]], align 4218// CHECK2-NEXT: [[TMP2:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])219// CHECK2-NEXT: store i32 0, ptr [[DOTZERO_ADDR]], align 4220// CHECK2-NEXT: store i32 [[TMP2]], ptr [[DOTTHREADID_TEMP_]], align 4221// CHECK2-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIPPcEiT__l15_omp_outlined(ptr [[DOTTHREADID_TEMP_]], ptr [[DOTZERO_ADDR]], ptr [[ARGC1]]) #[[ATTR3]]222// CHECK2-NEXT: call void @__kmpc_free_shared(ptr [[ARGC1]], i32 4)223// CHECK2-NEXT: call void @__kmpc_target_deinit()224// CHECK2-NEXT: ret void225// CHECK2: worker.exit:226// CHECK2-NEXT: ret void227//228//229// CHECK2-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIPPcEiT__l15_omp_outlined230// CHECK2-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[ARGC:%.*]]) #[[ATTR2]] {231// CHECK2-NEXT: entry:232// CHECK2-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 4233// CHECK2-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 4234// CHECK2-NEXT: [[ARGC_ADDR:%.*]] = alloca ptr, align 4235// CHECK2-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 4236// CHECK2-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 4237// CHECK2-NEXT: store ptr [[ARGC]], ptr [[ARGC_ADDR]], align 4238// CHECK2-NEXT: [[TMP0:%.*]] = load ptr, ptr [[ARGC_ADDR]], align 4239// CHECK2-NEXT: store ptr null, ptr [[TMP0]], align 4240// CHECK2-NEXT: ret void241//242//243// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l64244// CHECK3-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], i64 noundef [[A:%.*]], i64 noundef [[B:%.*]], i64 noundef [[ARGC:%.*]]) #[[ATTR0:[0-9]+]] {245// CHECK3-NEXT: entry:246// CHECK3-NEXT: [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 8247// CHECK3-NEXT: [[A_ADDR:%.*]] = alloca i64, align 8248// CHECK3-NEXT: [[B_ADDR:%.*]] = alloca i64, align 8249// CHECK3-NEXT: [[ARGC_ADDR:%.*]] = alloca i64, align 8250// CHECK3-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4251// CHECK3-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 4252// CHECK3-NEXT: store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 8253// CHECK3-NEXT: store i64 [[A]], ptr [[A_ADDR]], align 8254// CHECK3-NEXT: store i64 [[B]], ptr [[B_ADDR]], align 8255// CHECK3-NEXT: store i64 [[ARGC]], ptr [[ARGC_ADDR]], align 8256// CHECK3-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_target_init(ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l64_kernel_environment, ptr [[DYN_PTR]])257// CHECK3-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP0]], -1258// CHECK3-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]259// CHECK3: user_code.entry:260// CHECK3-NEXT: [[TMP1:%.*]] = load i32, ptr [[ARGC_ADDR]], align 4261// CHECK3-NEXT: [[ARGC1:%.*]] = call align 8 ptr @__kmpc_alloc_shared(i64 4)262// CHECK3-NEXT: store i32 [[TMP1]], ptr [[ARGC1]], align 4263// CHECK3-NEXT: [[TMP2:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1:[0-9]+]])264// CHECK3-NEXT: store i32 0, ptr [[DOTZERO_ADDR]], align 4265// CHECK3-NEXT: store i32 [[TMP2]], ptr [[DOTTHREADID_TEMP_]], align 4266// CHECK3-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l64_omp_outlined(ptr [[DOTTHREADID_TEMP_]], ptr [[DOTZERO_ADDR]], ptr [[ARGC1]]) #[[ATTR3:[0-9]+]]267// CHECK3-NEXT: call void @__kmpc_free_shared(ptr [[ARGC1]], i64 4)268// CHECK3-NEXT: call void @__kmpc_target_deinit()269// CHECK3-NEXT: ret void270// CHECK3: worker.exit:271// CHECK3-NEXT: ret void272//273//274// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l64_omp_outlined275// CHECK3-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[ARGC:%.*]]) #[[ATTR2:[0-9]+]] {276// CHECK3-NEXT: entry:277// CHECK3-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8278// CHECK3-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8279// CHECK3-NEXT: [[ARGC_ADDR:%.*]] = alloca ptr, align 8280// CHECK3-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8281// CHECK3-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8282// CHECK3-NEXT: store ptr [[ARGC]], ptr [[ARGC_ADDR]], align 8283// CHECK3-NEXT: [[TMP0:%.*]] = load ptr, ptr [[ARGC_ADDR]], align 8284// CHECK3-NEXT: store i32 0, ptr [[TMP0]], align 4285// CHECK3-NEXT: ret void286//287//288// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIPPcEiT__l53289// CHECK3-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], i64 noundef [[A:%.*]], i64 noundef [[B:%.*]], ptr noundef [[ARGC:%.*]]) #[[ATTR0]] {290// CHECK3-NEXT: entry:291// CHECK3-NEXT: [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 8292// CHECK3-NEXT: [[A_ADDR:%.*]] = alloca i64, align 8293// CHECK3-NEXT: [[B_ADDR:%.*]] = alloca i64, align 8294// CHECK3-NEXT: [[ARGC_ADDR:%.*]] = alloca ptr, align 8295// CHECK3-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4296// CHECK3-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 4297// CHECK3-NEXT: store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 8298// CHECK3-NEXT: store i64 [[A]], ptr [[A_ADDR]], align 8299// CHECK3-NEXT: store i64 [[B]], ptr [[B_ADDR]], align 8300// CHECK3-NEXT: store ptr [[ARGC]], ptr [[ARGC_ADDR]], align 8301// CHECK3-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_target_init(ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIPPcEiT__l53_kernel_environment, ptr [[DYN_PTR]])302// CHECK3-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP0]], -1303// CHECK3-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]304// CHECK3: user_code.entry:305// CHECK3-NEXT: [[TMP1:%.*]] = load ptr, ptr [[ARGC_ADDR]], align 8306// CHECK3-NEXT: [[ARGC1:%.*]] = call align 8 ptr @__kmpc_alloc_shared(i64 8)307// CHECK3-NEXT: store ptr [[TMP1]], ptr [[ARGC1]], align 8308// CHECK3-NEXT: [[TMP2:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])309// CHECK3-NEXT: store i32 0, ptr [[DOTZERO_ADDR]], align 4310// CHECK3-NEXT: store i32 [[TMP2]], ptr [[DOTTHREADID_TEMP_]], align 4311// CHECK3-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIPPcEiT__l53_omp_outlined(ptr [[DOTTHREADID_TEMP_]], ptr [[DOTZERO_ADDR]], ptr [[ARGC1]]) #[[ATTR3]]312// CHECK3-NEXT: call void @__kmpc_free_shared(ptr [[ARGC1]], i64 8)313// CHECK3-NEXT: call void @__kmpc_target_deinit()314// CHECK3-NEXT: ret void315// CHECK3: worker.exit:316// CHECK3-NEXT: ret void317//318//319// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIPPcEiT__l53_omp_outlined320// CHECK3-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 8 dereferenceable(8) [[ARGC:%.*]]) #[[ATTR2]] {321// CHECK3-NEXT: entry:322// CHECK3-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8323// CHECK3-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8324// CHECK3-NEXT: [[ARGC_ADDR:%.*]] = alloca ptr, align 8325// CHECK3-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8326// CHECK3-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8327// CHECK3-NEXT: store ptr [[ARGC]], ptr [[ARGC_ADDR]], align 8328// CHECK3-NEXT: [[TMP0:%.*]] = load ptr, ptr [[ARGC_ADDR]], align 8329// CHECK3-NEXT: store ptr null, ptr [[TMP0]], align 8330// CHECK3-NEXT: ret void331//332//333// CHECK4-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l64334// CHECK4-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], i32 noundef [[A:%.*]], i32 noundef [[B:%.*]], i32 noundef [[ARGC:%.*]]) #[[ATTR0:[0-9]+]] {335// CHECK4-NEXT: entry:336// CHECK4-NEXT: [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 4337// CHECK4-NEXT: [[A_ADDR:%.*]] = alloca i32, align 4338// CHECK4-NEXT: [[B_ADDR:%.*]] = alloca i32, align 4339// CHECK4-NEXT: [[ARGC_ADDR:%.*]] = alloca i32, align 4340// CHECK4-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4341// CHECK4-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 4342// CHECK4-NEXT: store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 4343// CHECK4-NEXT: store i32 [[A]], ptr [[A_ADDR]], align 4344// CHECK4-NEXT: store i32 [[B]], ptr [[B_ADDR]], align 4345// CHECK4-NEXT: store i32 [[ARGC]], ptr [[ARGC_ADDR]], align 4346// CHECK4-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_target_init(ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l64_kernel_environment, ptr [[DYN_PTR]])347// CHECK4-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP0]], -1348// CHECK4-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]349// CHECK4: user_code.entry:350// CHECK4-NEXT: [[TMP1:%.*]] = load i32, ptr [[ARGC_ADDR]], align 4351// CHECK4-NEXT: [[ARGC1:%.*]] = call align 8 ptr @__kmpc_alloc_shared(i32 4)352// CHECK4-NEXT: store i32 [[TMP1]], ptr [[ARGC1]], align 4353// CHECK4-NEXT: [[TMP2:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1:[0-9]+]])354// CHECK4-NEXT: store i32 0, ptr [[DOTZERO_ADDR]], align 4355// CHECK4-NEXT: store i32 [[TMP2]], ptr [[DOTTHREADID_TEMP_]], align 4356// CHECK4-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l64_omp_outlined(ptr [[DOTTHREADID_TEMP_]], ptr [[DOTZERO_ADDR]], ptr [[ARGC1]]) #[[ATTR3:[0-9]+]]357// CHECK4-NEXT: call void @__kmpc_free_shared(ptr [[ARGC1]], i32 4)358// CHECK4-NEXT: call void @__kmpc_target_deinit()359// CHECK4-NEXT: ret void360// CHECK4: worker.exit:361// CHECK4-NEXT: ret void362//363//364// CHECK4-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l64_omp_outlined365// CHECK4-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[ARGC:%.*]]) #[[ATTR2:[0-9]+]] {366// CHECK4-NEXT: entry:367// CHECK4-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 4368// CHECK4-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 4369// CHECK4-NEXT: [[ARGC_ADDR:%.*]] = alloca ptr, align 4370// CHECK4-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 4371// CHECK4-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 4372// CHECK4-NEXT: store ptr [[ARGC]], ptr [[ARGC_ADDR]], align 4373// CHECK4-NEXT: [[TMP0:%.*]] = load ptr, ptr [[ARGC_ADDR]], align 4374// CHECK4-NEXT: store i32 0, ptr [[TMP0]], align 4375// CHECK4-NEXT: ret void376//377//378// CHECK4-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIPPcEiT__l53379// CHECK4-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], i32 noundef [[A:%.*]], i32 noundef [[B:%.*]], ptr noundef [[ARGC:%.*]]) #[[ATTR0]] {380// CHECK4-NEXT: entry:381// CHECK4-NEXT: [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 4382// CHECK4-NEXT: [[A_ADDR:%.*]] = alloca i32, align 4383// CHECK4-NEXT: [[B_ADDR:%.*]] = alloca i32, align 4384// CHECK4-NEXT: [[ARGC_ADDR:%.*]] = alloca ptr, align 4385// CHECK4-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4386// CHECK4-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 4387// CHECK4-NEXT: store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 4388// CHECK4-NEXT: store i32 [[A]], ptr [[A_ADDR]], align 4389// CHECK4-NEXT: store i32 [[B]], ptr [[B_ADDR]], align 4390// CHECK4-NEXT: store ptr [[ARGC]], ptr [[ARGC_ADDR]], align 4391// CHECK4-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_target_init(ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIPPcEiT__l53_kernel_environment, ptr [[DYN_PTR]])392// CHECK4-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP0]], -1393// CHECK4-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]394// CHECK4: user_code.entry:395// CHECK4-NEXT: [[TMP1:%.*]] = load ptr, ptr [[ARGC_ADDR]], align 4396// CHECK4-NEXT: [[ARGC1:%.*]] = call align 8 ptr @__kmpc_alloc_shared(i32 4)397// CHECK4-NEXT: store ptr [[TMP1]], ptr [[ARGC1]], align 4398// CHECK4-NEXT: [[TMP2:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])399// CHECK4-NEXT: store i32 0, ptr [[DOTZERO_ADDR]], align 4400// CHECK4-NEXT: store i32 [[TMP2]], ptr [[DOTTHREADID_TEMP_]], align 4401// CHECK4-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIPPcEiT__l53_omp_outlined(ptr [[DOTTHREADID_TEMP_]], ptr [[DOTZERO_ADDR]], ptr [[ARGC1]]) #[[ATTR3]]402// CHECK4-NEXT: call void @__kmpc_free_shared(ptr [[ARGC1]], i32 4)403// CHECK4-NEXT: call void @__kmpc_target_deinit()404// CHECK4-NEXT: ret void405// CHECK4: worker.exit:406// CHECK4-NEXT: ret void407//408//409// CHECK4-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIPPcEiT__l53_omp_outlined410// CHECK4-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[ARGC:%.*]]) #[[ATTR2]] {411// CHECK4-NEXT: entry:412// CHECK4-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 4413// CHECK4-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 4414// CHECK4-NEXT: [[ARGC_ADDR:%.*]] = alloca ptr, align 4415// CHECK4-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 4416// CHECK4-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 4417// CHECK4-NEXT: store ptr [[ARGC]], ptr [[ARGC_ADDR]], align 4418// CHECK4-NEXT: [[TMP0:%.*]] = load ptr, ptr [[ARGC_ADDR]], align 4419// CHECK4-NEXT: store ptr null, ptr [[TMP0]], align 4420// CHECK4-NEXT: ret void421//422