918 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 -no-enable-noundef-analysis -verify -Wno-vla -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc4// RUN: %clang_cc1 -no-enable-noundef-analysis -verify -Wno-vla -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 - -disable-llvm-optzns | FileCheck %s --check-prefix=CHECK15// RUN: %clang_cc1 -no-enable-noundef-analysis -verify -Wno-vla -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm-bc %s -o %t-x86-host.bc6// RUN: %clang_cc1 -no-enable-noundef-analysis -verify -Wno-vla -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 - -disable-llvm-optzns | FileCheck %s --check-prefix=CHECK27// RUN: %clang_cc1 -no-enable-noundef-analysis -verify -Wno-vla -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 - -disable-llvm-optzns | FileCheck %s --check-prefix=CHECK28// expected-no-diagnostics9 10#ifndef HEADER11#define HEADER12 13__thread int id;14 15int baz(int f, double &a);16 17template <typename tx, typename ty>18struct TT {19 tx X;20 ty Y;21 tx &operator[](int i) { return X; }22};23 24void targetBar(int *Ptr1, int *Ptr2) {25#pragma omp target map(Ptr1[:0], Ptr2)26#pragma omp parallel num_threads(2)27 *Ptr1 = *Ptr2;28}29 30int foo(int n) {31 int a = 0;32 short aa = 0;33 float b[10];34 float bn[n];35 double c[5][10];36 double cn[5][n];37 TT<long long, char> d;38 39#pragma omp target40 {41 }42 43#pragma omp target if (0)44 {45 }46 47#pragma omp target if (1)48 {49 aa += 1;50 aa += 2;51 }52 53#pragma omp target if (n > 20)54 {55 a += 1;56 b[2] += 1.0;57 bn[3] += 1.0;58 c[1][2] += 1.0;59 cn[1][3] += 1.0;60 d.X += 1;61 d.Y += 1;62 d[0] += 1;63 }64 65 return a;66}67 68template <typename tx>69tx ftemplate(int n) {70 tx a = 0;71 short aa = 0;72 tx b[10];73 74#pragma omp target if (n > 40)75 {76 a += 1;77 aa += 1;78 b[2] += 1;79 }80 81 return a;82}83 84static int fstatic(int n) {85 int a = 0;86 short aa = 0;87 char aaa = 0;88 int b[10];89 90#pragma omp target if (n > 50)91 {92 a += 1;93 aa += 1;94 aaa += 1;95 b[2] += 1;96 }97 98 return a;99}100 101struct S1 {102 double a;103 104 int r1(int n) {105 int b = n + 1;106 short int c[2][n];107 108#pragma omp target if (n > 60)109 {110 this->a = (double)b + 1.5;111 c[1][1] = ++a;112 baz(a, a);113 }114 115 return c[1][1] + (int)b;116 }117};118 119int bar(int n) {120 int a = 0;121 122 a += foo(n);123 124 S1 S;125 a += S.r1(n);126 127 a += fstatic(n);128 129 a += ftemplate<int>(n);130 131 return a;132}133 134int baz(int f, double &a) {135#pragma omp parallel136 f = 2 + a;137 return f;138}139 140extern void assert(int) throw() __attribute__((__noreturn__));141void unreachable_call() {142#pragma omp target143 assert(0);144}145 146#endif147// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9targetBarPiS__l25148// CHECK1-SAME: (ptr noalias [[DYN_PTR:%.*]], ptr [[PTR1:%.*]], ptr nonnull align 8 dereferenceable(8) [[PTR2:%.*]]) #[[ATTR0:[0-9]+]] {149// CHECK1-NEXT: entry:150// CHECK1-NEXT: [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 8151// CHECK1-NEXT: [[PTR1_ADDR:%.*]] = alloca ptr, align 8152// CHECK1-NEXT: [[PTR2_ADDR:%.*]] = alloca ptr, align 8153// CHECK1-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [2 x ptr], align 8154// CHECK1-NEXT: store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 8155// CHECK1-NEXT: store ptr [[PTR1]], ptr [[PTR1_ADDR]], align 8156// CHECK1-NEXT: store ptr [[PTR2]], ptr [[PTR2_ADDR]], align 8157// CHECK1-NEXT: [[TMP0:%.*]] = load ptr, ptr [[PTR2_ADDR]], align 8158// CHECK1-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_target_init(ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9targetBarPiS__l25_kernel_environment, ptr [[DYN_PTR]])159// CHECK1-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP1]], -1160// CHECK1-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]161// CHECK1: user_code.entry:162// CHECK1-NEXT: [[TMP2:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1:[0-9]+]])163// CHECK1-NEXT: [[TMP3:%.*]] = getelementptr inbounds [2 x ptr], ptr [[CAPTURED_VARS_ADDRS]], i64 0, i64 0164// CHECK1-NEXT: store ptr [[PTR1_ADDR]], ptr [[TMP3]], align 8165// CHECK1-NEXT: [[TMP4:%.*]] = getelementptr inbounds [2 x ptr], ptr [[CAPTURED_VARS_ADDRS]], i64 0, i64 1166// CHECK1-NEXT: store ptr [[TMP0]], ptr [[TMP4]], align 8167// CHECK1-NEXT: call void @__kmpc_parallel_51(ptr @[[GLOB1]], i32 [[TMP2]], i32 1, i32 2, i32 -1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9targetBarPiS__l25_omp_outlined, ptr null, ptr [[CAPTURED_VARS_ADDRS]], i64 2)168// CHECK1-NEXT: call void @__kmpc_target_deinit()169// CHECK1-NEXT: ret void170// CHECK1: worker.exit:171// CHECK1-NEXT: ret void172//173//174// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9targetBarPiS__l25_omp_outlined175// CHECK1-SAME: (ptr noalias [[DOTGLOBAL_TID_:%.*]], ptr noalias [[DOTBOUND_TID_:%.*]], ptr nonnull align 8 dereferenceable(8) [[PTR1:%.*]], ptr nonnull align 8 dereferenceable(8) [[PTR2:%.*]]) #[[ATTR1:[0-9]+]] {176// CHECK1-NEXT: entry:177// CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8178// CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8179// CHECK1-NEXT: [[PTR1_ADDR:%.*]] = alloca ptr, align 8180// CHECK1-NEXT: [[PTR2_ADDR:%.*]] = alloca ptr, align 8181// CHECK1-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8182// CHECK1-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8183// CHECK1-NEXT: store ptr [[PTR1]], ptr [[PTR1_ADDR]], align 8184// CHECK1-NEXT: store ptr [[PTR2]], ptr [[PTR2_ADDR]], align 8185// CHECK1-NEXT: [[TMP0:%.*]] = load ptr, ptr [[PTR1_ADDR]], align 8186// CHECK1-NEXT: [[TMP1:%.*]] = load ptr, ptr [[PTR2_ADDR]], align 8187// CHECK1-NEXT: [[TMP2:%.*]] = load ptr, ptr [[TMP1]], align 8188// CHECK1-NEXT: [[TMP3:%.*]] = load i32, ptr [[TMP2]], align 4189// CHECK1-NEXT: [[TMP4:%.*]] = load ptr, ptr [[TMP0]], align 8190// CHECK1-NEXT: store i32 [[TMP3]], ptr [[TMP4]], align 4191// CHECK1-NEXT: ret void192//193//194// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3fooi_l39195// CHECK1-SAME: (ptr noalias [[DYN_PTR:%.*]]) #[[ATTR4:[0-9]+]] {196// CHECK1-NEXT: entry:197// CHECK1-NEXT: [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 8198// CHECK1-NEXT: store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 8199// CHECK1-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_target_init(ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3fooi_l39_kernel_environment, ptr [[DYN_PTR]])200// CHECK1-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP0]], -1201// CHECK1-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]202// CHECK1: user_code.entry:203// CHECK1-NEXT: call void @__kmpc_target_deinit()204// CHECK1-NEXT: ret void205// CHECK1: worker.exit:206// CHECK1-NEXT: ret void207//208//209// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3fooi_l47210// CHECK1-SAME: (ptr noalias [[DYN_PTR:%.*]], i64 [[AA:%.*]]) #[[ATTR4]] {211// CHECK1-NEXT: entry:212// CHECK1-NEXT: [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 8213// CHECK1-NEXT: [[AA_ADDR:%.*]] = alloca i64, align 8214// CHECK1-NEXT: store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 8215// CHECK1-NEXT: store i64 [[AA]], ptr [[AA_ADDR]], align 8216// CHECK1-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_target_init(ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3fooi_l47_kernel_environment, ptr [[DYN_PTR]])217// CHECK1-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP0]], -1218// CHECK1-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]219// CHECK1: user_code.entry:220// CHECK1-NEXT: [[TMP1:%.*]] = load i16, ptr [[AA_ADDR]], align 2221// CHECK1-NEXT: [[CONV:%.*]] = sext i16 [[TMP1]] to i32222// CHECK1-NEXT: [[ADD:%.*]] = add nsw i32 [[CONV]], 1223// CHECK1-NEXT: [[CONV1:%.*]] = trunc i32 [[ADD]] to i16224// CHECK1-NEXT: store i16 [[CONV1]], ptr [[AA_ADDR]], align 2225// CHECK1-NEXT: [[TMP2:%.*]] = load i16, ptr [[AA_ADDR]], align 2226// CHECK1-NEXT: [[CONV2:%.*]] = sext i16 [[TMP2]] to i32227// CHECK1-NEXT: [[ADD3:%.*]] = add nsw i32 [[CONV2]], 2228// CHECK1-NEXT: [[CONV4:%.*]] = trunc i32 [[ADD3]] to i16229// CHECK1-NEXT: store i16 [[CONV4]], ptr [[AA_ADDR]], align 2230// CHECK1-NEXT: call void @__kmpc_target_deinit()231// CHECK1-NEXT: ret void232// CHECK1: worker.exit:233// CHECK1-NEXT: ret void234//235//236// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3fooi_l53237// CHECK1-SAME: (ptr noalias [[DYN_PTR:%.*]], i64 [[A:%.*]], ptr nonnull align 4 dereferenceable(40) [[B:%.*]], i64 [[VLA:%.*]], ptr nonnull align 4 dereferenceable(4) [[BN:%.*]], ptr nonnull align 8 dereferenceable(400) [[C:%.*]], i64 [[VLA1:%.*]], i64 [[VLA3:%.*]], ptr nonnull align 8 dereferenceable(8) [[CN:%.*]], ptr nonnull align 8 dereferenceable(16) [[D:%.*]]) #[[ATTR4]] {238// CHECK1-NEXT: entry:239// CHECK1-NEXT: [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 8240// CHECK1-NEXT: [[A_ADDR:%.*]] = alloca i64, align 8241// CHECK1-NEXT: [[B_ADDR:%.*]] = alloca ptr, align 8242// CHECK1-NEXT: [[VLA_ADDR:%.*]] = alloca i64, align 8243// CHECK1-NEXT: [[BN_ADDR:%.*]] = alloca ptr, align 8244// CHECK1-NEXT: [[C_ADDR:%.*]] = alloca ptr, align 8245// CHECK1-NEXT: [[VLA_ADDR2:%.*]] = alloca i64, align 8246// CHECK1-NEXT: [[VLA_ADDR4:%.*]] = alloca i64, align 8247// CHECK1-NEXT: [[CN_ADDR:%.*]] = alloca ptr, align 8248// CHECK1-NEXT: [[D_ADDR:%.*]] = alloca ptr, align 8249// CHECK1-NEXT: store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 8250// CHECK1-NEXT: store i64 [[A]], ptr [[A_ADDR]], align 8251// CHECK1-NEXT: store ptr [[B]], ptr [[B_ADDR]], align 8252// CHECK1-NEXT: store i64 [[VLA]], ptr [[VLA_ADDR]], align 8253// CHECK1-NEXT: store ptr [[BN]], ptr [[BN_ADDR]], align 8254// CHECK1-NEXT: store ptr [[C]], ptr [[C_ADDR]], align 8255// CHECK1-NEXT: store i64 [[VLA1]], ptr [[VLA_ADDR2]], align 8256// CHECK1-NEXT: store i64 [[VLA3]], ptr [[VLA_ADDR4]], align 8257// CHECK1-NEXT: store ptr [[CN]], ptr [[CN_ADDR]], align 8258// CHECK1-NEXT: store ptr [[D]], ptr [[D_ADDR]], align 8259// CHECK1-NEXT: [[TMP0:%.*]] = load ptr, ptr [[B_ADDR]], align 8260// CHECK1-NEXT: [[TMP1:%.*]] = load i64, ptr [[VLA_ADDR]], align 8261// CHECK1-NEXT: [[TMP2:%.*]] = load ptr, ptr [[BN_ADDR]], align 8262// CHECK1-NEXT: [[TMP3:%.*]] = load ptr, ptr [[C_ADDR]], align 8263// CHECK1-NEXT: [[TMP4:%.*]] = load i64, ptr [[VLA_ADDR2]], align 8264// CHECK1-NEXT: [[TMP5:%.*]] = load i64, ptr [[VLA_ADDR4]], align 8265// CHECK1-NEXT: [[TMP6:%.*]] = load ptr, ptr [[CN_ADDR]], align 8266// CHECK1-NEXT: [[TMP7:%.*]] = load ptr, ptr [[D_ADDR]], align 8267// CHECK1-NEXT: [[TMP8:%.*]] = call i32 @__kmpc_target_init(ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3fooi_l53_kernel_environment, ptr [[DYN_PTR]])268// CHECK1-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP8]], -1269// CHECK1-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]270// CHECK1: user_code.entry:271// CHECK1-NEXT: [[TMP9:%.*]] = load i32, ptr [[A_ADDR]], align 4272// CHECK1-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP9]], 1273// CHECK1-NEXT: store i32 [[ADD]], ptr [[A_ADDR]], align 4274// CHECK1-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x float], ptr [[TMP0]], i64 0, i64 2275// CHECK1-NEXT: [[TMP10:%.*]] = load float, ptr [[ARRAYIDX]], align 4276// CHECK1-NEXT: [[CONV:%.*]] = fpext float [[TMP10]] to double277// CHECK1-NEXT: [[ADD5:%.*]] = fadd double [[CONV]], 1.000000e+00278// CHECK1-NEXT: [[CONV6:%.*]] = fptrunc double [[ADD5]] to float279// CHECK1-NEXT: store float [[CONV6]], ptr [[ARRAYIDX]], align 4280// CHECK1-NEXT: [[ARRAYIDX7:%.*]] = getelementptr inbounds float, ptr [[TMP2]], i64 3281// CHECK1-NEXT: [[TMP11:%.*]] = load float, ptr [[ARRAYIDX7]], align 4282// CHECK1-NEXT: [[CONV8:%.*]] = fpext float [[TMP11]] to double283// CHECK1-NEXT: [[ADD9:%.*]] = fadd double [[CONV8]], 1.000000e+00284// CHECK1-NEXT: [[CONV10:%.*]] = fptrunc double [[ADD9]] to float285// CHECK1-NEXT: store float [[CONV10]], ptr [[ARRAYIDX7]], align 4286// CHECK1-NEXT: [[ARRAYIDX11:%.*]] = getelementptr inbounds [5 x [10 x double]], ptr [[TMP3]], i64 0, i64 1287// CHECK1-NEXT: [[ARRAYIDX12:%.*]] = getelementptr inbounds [10 x double], ptr [[ARRAYIDX11]], i64 0, i64 2288// CHECK1-NEXT: [[TMP12:%.*]] = load double, ptr [[ARRAYIDX12]], align 8289// CHECK1-NEXT: [[ADD13:%.*]] = fadd double [[TMP12]], 1.000000e+00290// CHECK1-NEXT: store double [[ADD13]], ptr [[ARRAYIDX12]], align 8291// CHECK1-NEXT: [[TMP13:%.*]] = mul nsw i64 1, [[TMP5]]292// CHECK1-NEXT: [[ARRAYIDX14:%.*]] = getelementptr inbounds double, ptr [[TMP6]], i64 [[TMP13]]293// CHECK1-NEXT: [[ARRAYIDX15:%.*]] = getelementptr inbounds double, ptr [[ARRAYIDX14]], i64 3294// CHECK1-NEXT: [[TMP14:%.*]] = load double, ptr [[ARRAYIDX15]], align 8295// CHECK1-NEXT: [[ADD16:%.*]] = fadd double [[TMP14]], 1.000000e+00296// CHECK1-NEXT: store double [[ADD16]], ptr [[ARRAYIDX15]], align 8297// CHECK1-NEXT: [[X:%.*]] = getelementptr inbounds nuw [[STRUCT_TT:%.*]], ptr [[TMP7]], i32 0, i32 0298// CHECK1-NEXT: [[TMP15:%.*]] = load i64, ptr [[X]], align 8299// CHECK1-NEXT: [[ADD17:%.*]] = add nsw i64 [[TMP15]], 1300// CHECK1-NEXT: store i64 [[ADD17]], ptr [[X]], align 8301// CHECK1-NEXT: [[Y:%.*]] = getelementptr inbounds nuw [[STRUCT_TT]], ptr [[TMP7]], i32 0, i32 1302// CHECK1-NEXT: [[TMP16:%.*]] = load i8, ptr [[Y]], align 8303// CHECK1-NEXT: [[CONV18:%.*]] = sext i8 [[TMP16]] to i32304// CHECK1-NEXT: [[ADD19:%.*]] = add nsw i32 [[CONV18]], 1305// CHECK1-NEXT: [[CONV20:%.*]] = trunc i32 [[ADD19]] to i8306// CHECK1-NEXT: store i8 [[CONV20]], ptr [[Y]], align 8307// CHECK1-NEXT: [[CALL:%.*]] = call nonnull align 8 dereferenceable(8) ptr @_ZN2TTIxcEixEi(ptr nonnull align 8 dereferenceable(16) [[TMP7]], i32 0) #[[ATTR10:[0-9]+]]308// CHECK1-NEXT: [[TMP17:%.*]] = load i64, ptr [[CALL]], align 8309// CHECK1-NEXT: [[ADD21:%.*]] = add nsw i64 [[TMP17]], 1310// CHECK1-NEXT: store i64 [[ADD21]], ptr [[CALL]], align 8311// CHECK1-NEXT: call void @__kmpc_target_deinit()312// CHECK1-NEXT: ret void313// CHECK1: worker.exit:314// CHECK1-NEXT: ret void315//316//317// CHECK1-LABEL: define {{[^@]+}}@_ZN2TTIxcEixEi318// CHECK1-SAME: (ptr nonnull align 8 dereferenceable(16) [[THIS:%.*]], i32 [[I:%.*]]) #[[ATTR5:[0-9]+]] comdat align 2 {319// CHECK1-NEXT: entry:320// CHECK1-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 8321// CHECK1-NEXT: [[I_ADDR:%.*]] = alloca i32, align 4322// CHECK1-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 8323// CHECK1-NEXT: store i32 [[I]], ptr [[I_ADDR]], align 4324// CHECK1-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8325// CHECK1-NEXT: [[X:%.*]] = getelementptr inbounds nuw [[STRUCT_TT:%.*]], ptr [[THIS1]], i32 0, i32 0326// CHECK1-NEXT: ret ptr [[X]]327//328//329// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l90330// CHECK1-SAME: (ptr noalias [[DYN_PTR:%.*]], i64 [[A:%.*]], i64 [[AA:%.*]], i64 [[AAA:%.*]], ptr nonnull align 4 dereferenceable(40) [[B:%.*]]) #[[ATTR4]] {331// CHECK1-NEXT: entry:332// CHECK1-NEXT: [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 8333// CHECK1-NEXT: [[A_ADDR:%.*]] = alloca i64, align 8334// CHECK1-NEXT: [[AA_ADDR:%.*]] = alloca i64, align 8335// CHECK1-NEXT: [[AAA_ADDR:%.*]] = alloca i64, align 8336// CHECK1-NEXT: [[B_ADDR:%.*]] = alloca ptr, align 8337// CHECK1-NEXT: store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 8338// CHECK1-NEXT: store i64 [[A]], ptr [[A_ADDR]], align 8339// CHECK1-NEXT: store i64 [[AA]], ptr [[AA_ADDR]], align 8340// CHECK1-NEXT: store i64 [[AAA]], ptr [[AAA_ADDR]], align 8341// CHECK1-NEXT: store ptr [[B]], ptr [[B_ADDR]], align 8342// CHECK1-NEXT: [[TMP0:%.*]] = load ptr, ptr [[B_ADDR]], align 8343// CHECK1-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_target_init(ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l90_kernel_environment, ptr [[DYN_PTR]])344// CHECK1-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP1]], -1345// CHECK1-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]346// CHECK1: user_code.entry:347// CHECK1-NEXT: [[TMP2:%.*]] = load i32, ptr [[A_ADDR]], align 4348// CHECK1-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP2]], 1349// CHECK1-NEXT: store i32 [[ADD]], ptr [[A_ADDR]], align 4350// CHECK1-NEXT: [[TMP3:%.*]] = load i16, ptr [[AA_ADDR]], align 2351// CHECK1-NEXT: [[CONV:%.*]] = sext i16 [[TMP3]] to i32352// CHECK1-NEXT: [[ADD1:%.*]] = add nsw i32 [[CONV]], 1353// CHECK1-NEXT: [[CONV2:%.*]] = trunc i32 [[ADD1]] to i16354// CHECK1-NEXT: store i16 [[CONV2]], ptr [[AA_ADDR]], align 2355// CHECK1-NEXT: [[TMP4:%.*]] = load i8, ptr [[AAA_ADDR]], align 1356// CHECK1-NEXT: [[CONV3:%.*]] = sext i8 [[TMP4]] to i32357// CHECK1-NEXT: [[ADD4:%.*]] = add nsw i32 [[CONV3]], 1358// CHECK1-NEXT: [[CONV5:%.*]] = trunc i32 [[ADD4]] to i8359// CHECK1-NEXT: store i8 [[CONV5]], ptr [[AAA_ADDR]], align 1360// CHECK1-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x i32], ptr [[TMP0]], i64 0, i64 2361// CHECK1-NEXT: [[TMP5:%.*]] = load i32, ptr [[ARRAYIDX]], align 4362// CHECK1-NEXT: [[ADD6:%.*]] = add nsw i32 [[TMP5]], 1363// CHECK1-NEXT: store i32 [[ADD6]], ptr [[ARRAYIDX]], align 4364// CHECK1-NEXT: call void @__kmpc_target_deinit()365// CHECK1-NEXT: ret void366// CHECK1: worker.exit:367// CHECK1-NEXT: ret void368//369//370// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2S12r1Ei_l108371// CHECK1-SAME: (ptr noalias [[DYN_PTR:%.*]], ptr [[THIS:%.*]], i64 [[B:%.*]], i64 [[VLA:%.*]], i64 [[VLA1:%.*]], ptr nonnull align 2 dereferenceable(2) [[C:%.*]]) #[[ATTR4]] {372// CHECK1-NEXT: entry:373// CHECK1-NEXT: [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 8374// CHECK1-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 8375// CHECK1-NEXT: [[B_ADDR:%.*]] = alloca i64, align 8376// CHECK1-NEXT: [[VLA_ADDR:%.*]] = alloca i64, align 8377// CHECK1-NEXT: [[VLA_ADDR2:%.*]] = alloca i64, align 8378// CHECK1-NEXT: [[C_ADDR:%.*]] = alloca ptr, align 8379// CHECK1-NEXT: store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 8380// CHECK1-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 8381// CHECK1-NEXT: store i64 [[B]], ptr [[B_ADDR]], align 8382// CHECK1-NEXT: store i64 [[VLA]], ptr [[VLA_ADDR]], align 8383// CHECK1-NEXT: store i64 [[VLA1]], ptr [[VLA_ADDR2]], align 8384// CHECK1-NEXT: store ptr [[C]], ptr [[C_ADDR]], align 8385// CHECK1-NEXT: [[TMP0:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8386// CHECK1-NEXT: [[TMP1:%.*]] = load i64, ptr [[VLA_ADDR]], align 8387// CHECK1-NEXT: [[TMP2:%.*]] = load i64, ptr [[VLA_ADDR2]], align 8388// CHECK1-NEXT: [[TMP3:%.*]] = load ptr, ptr [[C_ADDR]], align 8389// CHECK1-NEXT: [[TMP4:%.*]] = call i32 @__kmpc_target_init(ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2S12r1Ei_l108_kernel_environment, ptr [[DYN_PTR]])390// CHECK1-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP4]], -1391// CHECK1-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]392// CHECK1: user_code.entry:393// CHECK1-NEXT: [[TMP5:%.*]] = load i32, ptr [[B_ADDR]], align 4394// CHECK1-NEXT: [[CONV:%.*]] = sitofp i32 [[TMP5]] to double395// CHECK1-NEXT: [[ADD:%.*]] = fadd double [[CONV]], 1.500000e+00396// CHECK1-NEXT: [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_S1:%.*]], ptr [[TMP0]], i32 0, i32 0397// CHECK1-NEXT: store double [[ADD]], ptr [[A]], align 8398// CHECK1-NEXT: [[A3:%.*]] = getelementptr inbounds nuw [[STRUCT_S1]], ptr [[TMP0]], i32 0, i32 0399// CHECK1-NEXT: [[TMP6:%.*]] = load double, ptr [[A3]], align 8400// CHECK1-NEXT: [[INC:%.*]] = fadd double [[TMP6]], 1.000000e+00401// CHECK1-NEXT: store double [[INC]], ptr [[A3]], align 8402// CHECK1-NEXT: [[CONV4:%.*]] = fptosi double [[INC]] to i16403// CHECK1-NEXT: [[TMP7:%.*]] = mul nsw i64 1, [[TMP2]]404// CHECK1-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds i16, ptr [[TMP3]], i64 [[TMP7]]405// CHECK1-NEXT: [[ARRAYIDX5:%.*]] = getelementptr inbounds i16, ptr [[ARRAYIDX]], i64 1406// CHECK1-NEXT: store i16 [[CONV4]], ptr [[ARRAYIDX5]], align 2407// CHECK1-NEXT: [[A6:%.*]] = getelementptr inbounds nuw [[STRUCT_S1]], ptr [[TMP0]], i32 0, i32 0408// CHECK1-NEXT: [[TMP8:%.*]] = load double, ptr [[A6]], align 8409// CHECK1-NEXT: [[CONV7:%.*]] = fptosi double [[TMP8]] to i32410// CHECK1-NEXT: [[A8:%.*]] = getelementptr inbounds nuw [[STRUCT_S1]], ptr [[TMP0]], i32 0, i32 0411// CHECK1-NEXT: [[CALL:%.*]] = call i32 @_Z3baziRd(i32 [[CONV7]], ptr nonnull align 8 dereferenceable(8) [[A8]]) #[[ATTR10]]412// CHECK1-NEXT: call void @__kmpc_target_deinit()413// CHECK1-NEXT: ret void414// CHECK1: worker.exit:415// CHECK1-NEXT: ret void416//417//418// CHECK1-LABEL: define {{[^@]+}}@_Z3baziRd419// CHECK1-SAME: (i32 [[F1:%.*]], ptr nonnull align 8 dereferenceable(8) [[A:%.*]]) #[[ATTR5]] {420// CHECK1-NEXT: entry:421// CHECK1-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 8422// CHECK1-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [2 x ptr], align 8423// CHECK1-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])424// CHECK1-NEXT: [[F:%.*]] = call align 8 ptr @__kmpc_alloc_shared(i64 4)425// CHECK1-NEXT: store i32 [[F1]], ptr [[F]], align 4426// CHECK1-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 8427// CHECK1-NEXT: [[TMP1:%.*]] = load ptr, ptr [[A_ADDR]], align 8428// CHECK1-NEXT: [[TMP2:%.*]] = getelementptr inbounds [2 x ptr], ptr [[CAPTURED_VARS_ADDRS]], i64 0, i64 0429// CHECK1-NEXT: store ptr [[F]], ptr [[TMP2]], align 8430// CHECK1-NEXT: [[TMP3:%.*]] = getelementptr inbounds [2 x ptr], ptr [[CAPTURED_VARS_ADDRS]], i64 0, i64 1431// CHECK1-NEXT: store ptr [[TMP1]], ptr [[TMP3]], align 8432// CHECK1-NEXT: call void @__kmpc_parallel_51(ptr @[[GLOB1]], i32 [[TMP0]], i32 1, i32 -1, i32 -1, ptr @_Z3baziRd_omp_outlined, ptr @_Z3baziRd_omp_outlined_wrapper, ptr [[CAPTURED_VARS_ADDRS]], i64 2)433// CHECK1-NEXT: [[TMP4:%.*]] = load i32, ptr [[F]], align 4434// CHECK1-NEXT: call void @__kmpc_free_shared(ptr [[F]], i64 4)435// CHECK1-NEXT: ret i32 [[TMP4]]436//437//438// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z16unreachable_callv_l142439// CHECK1-SAME: (ptr noalias [[DYN_PTR:%.*]]) #[[ATTR4]] {440// CHECK1-NEXT: entry:441// CHECK1-NEXT: [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 8442// CHECK1-NEXT: store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 8443// CHECK1-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_target_init(ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z16unreachable_callv_l142_kernel_environment, ptr [[DYN_PTR]])444// CHECK1-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP0]], -1445// CHECK1-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]446// CHECK1: user_code.entry:447// CHECK1-NEXT: call void @_Z6asserti(i32 0) #[[ATTR11:[0-9]+]]448// CHECK1-NEXT: unreachable449// CHECK1: worker.exit:450// CHECK1-NEXT: ret void451// CHECK1: 1:452// CHECK1-NEXT: call void @__kmpc_target_deinit()453// CHECK1-NEXT: ret void454//455//456// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l74457// CHECK1-SAME: (ptr noalias [[DYN_PTR:%.*]], i64 [[A:%.*]], i64 [[AA:%.*]], ptr nonnull align 4 dereferenceable(40) [[B:%.*]]) #[[ATTR4]] {458// CHECK1-NEXT: entry:459// CHECK1-NEXT: [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 8460// CHECK1-NEXT: [[A_ADDR:%.*]] = alloca i64, align 8461// CHECK1-NEXT: [[AA_ADDR:%.*]] = alloca i64, align 8462// CHECK1-NEXT: [[B_ADDR:%.*]] = alloca ptr, align 8463// CHECK1-NEXT: store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 8464// CHECK1-NEXT: store i64 [[A]], ptr [[A_ADDR]], align 8465// CHECK1-NEXT: store i64 [[AA]], ptr [[AA_ADDR]], align 8466// CHECK1-NEXT: store ptr [[B]], ptr [[B_ADDR]], align 8467// CHECK1-NEXT: [[TMP0:%.*]] = load ptr, ptr [[B_ADDR]], align 8468// CHECK1-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_target_init(ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l74_kernel_environment, ptr [[DYN_PTR]])469// CHECK1-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP1]], -1470// CHECK1-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]471// CHECK1: user_code.entry:472// CHECK1-NEXT: [[TMP2:%.*]] = load i32, ptr [[A_ADDR]], align 4473// CHECK1-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP2]], 1474// CHECK1-NEXT: store i32 [[ADD]], ptr [[A_ADDR]], align 4475// CHECK1-NEXT: [[TMP3:%.*]] = load i16, ptr [[AA_ADDR]], align 2476// CHECK1-NEXT: [[CONV:%.*]] = sext i16 [[TMP3]] to i32477// CHECK1-NEXT: [[ADD1:%.*]] = add nsw i32 [[CONV]], 1478// CHECK1-NEXT: [[CONV2:%.*]] = trunc i32 [[ADD1]] to i16479// CHECK1-NEXT: store i16 [[CONV2]], ptr [[AA_ADDR]], align 2480// CHECK1-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x i32], ptr [[TMP0]], i64 0, i64 2481// CHECK1-NEXT: [[TMP4:%.*]] = load i32, ptr [[ARRAYIDX]], align 4482// CHECK1-NEXT: [[ADD3:%.*]] = add nsw i32 [[TMP4]], 1483// CHECK1-NEXT: store i32 [[ADD3]], ptr [[ARRAYIDX]], align 4484// CHECK1-NEXT: call void @__kmpc_target_deinit()485// CHECK1-NEXT: ret void486// CHECK1: worker.exit:487// CHECK1-NEXT: ret void488//489//490// CHECK1-LABEL: define {{[^@]+}}@_Z3baziRd_omp_outlined491// CHECK1-SAME: (ptr noalias [[DOTGLOBAL_TID_:%.*]], ptr noalias [[DOTBOUND_TID_:%.*]], ptr nonnull align 4 dereferenceable(4) [[F:%.*]], ptr nonnull align 8 dereferenceable(8) [[A:%.*]]) #[[ATTR1]] {492// CHECK1-NEXT: entry:493// CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8494// CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8495// CHECK1-NEXT: [[F_ADDR:%.*]] = alloca ptr, align 8496// CHECK1-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 8497// CHECK1-NEXT: [[TMP:%.*]] = alloca ptr, align 8498// CHECK1-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8499// CHECK1-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8500// CHECK1-NEXT: store ptr [[F]], ptr [[F_ADDR]], align 8501// CHECK1-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 8502// CHECK1-NEXT: [[TMP0:%.*]] = load ptr, ptr [[F_ADDR]], align 8503// CHECK1-NEXT: [[TMP1:%.*]] = load ptr, ptr [[A_ADDR]], align 8504// CHECK1-NEXT: store ptr [[TMP1]], ptr [[TMP]], align 8505// CHECK1-NEXT: [[TMP2:%.*]] = load ptr, ptr [[TMP]], align 8506// CHECK1-NEXT: [[TMP3:%.*]] = load double, ptr [[TMP2]], align 8507// CHECK1-NEXT: [[ADD:%.*]] = fadd double 2.000000e+00, [[TMP3]]508// CHECK1-NEXT: [[CONV:%.*]] = fptosi double [[ADD]] to i32509// CHECK1-NEXT: store i32 [[CONV]], ptr [[TMP0]], align 4510// CHECK1-NEXT: ret void511//512//513// CHECK1-LABEL: define {{[^@]+}}@_Z3baziRd_omp_outlined_wrapper514// CHECK1-SAME: (i16 zeroext [[TMP0:%.*]], i32 [[TMP1:%.*]]) #[[ATTR8:[0-9]+]] {515// CHECK1-NEXT: entry:516// CHECK1-NEXT: [[DOTADDR:%.*]] = alloca i16, align 2517// CHECK1-NEXT: [[DOTADDR1:%.*]] = alloca i32, align 4518// CHECK1-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4519// CHECK1-NEXT: [[GLOBAL_ARGS:%.*]] = alloca ptr, align 8520// CHECK1-NEXT: store i16 [[TMP0]], ptr [[DOTADDR]], align 2521// CHECK1-NEXT: store i32 [[TMP1]], ptr [[DOTADDR1]], align 4522// CHECK1-NEXT: store i32 0, ptr [[DOTZERO_ADDR]], align 4523// CHECK1-NEXT: call void @__kmpc_get_shared_variables(ptr [[GLOBAL_ARGS]])524// CHECK1-NEXT: [[TMP2:%.*]] = load ptr, ptr [[GLOBAL_ARGS]], align 8525// CHECK1-NEXT: [[TMP3:%.*]] = getelementptr inbounds ptr, ptr [[TMP2]], i64 0526// CHECK1-NEXT: [[TMP4:%.*]] = load ptr, ptr [[TMP3]], align 8527// CHECK1-NEXT: [[TMP5:%.*]] = getelementptr inbounds ptr, ptr [[TMP2]], i64 1528// CHECK1-NEXT: [[TMP6:%.*]] = load ptr, ptr [[TMP5]], align 8529// CHECK1-NEXT: call void @_Z3baziRd_omp_outlined(ptr [[DOTADDR1]], ptr [[DOTZERO_ADDR]], ptr [[TMP4]], ptr [[TMP6]]) #[[ATTR2:[0-9]+]]530// CHECK1-NEXT: ret void531//532//533// CHECK2-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9targetBarPiS__l25534// CHECK2-SAME: (ptr noalias [[DYN_PTR:%.*]], ptr [[PTR1:%.*]], ptr nonnull align 4 dereferenceable(4) [[PTR2:%.*]]) #[[ATTR0:[0-9]+]] {535// CHECK2-NEXT: entry:536// CHECK2-NEXT: [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 4537// CHECK2-NEXT: [[PTR1_ADDR:%.*]] = alloca ptr, align 4538// CHECK2-NEXT: [[PTR2_ADDR:%.*]] = alloca ptr, align 4539// CHECK2-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [2 x ptr], align 4540// CHECK2-NEXT: store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 4541// CHECK2-NEXT: store ptr [[PTR1]], ptr [[PTR1_ADDR]], align 4542// CHECK2-NEXT: store ptr [[PTR2]], ptr [[PTR2_ADDR]], align 4543// CHECK2-NEXT: [[TMP0:%.*]] = load ptr, ptr [[PTR2_ADDR]], align 4544// CHECK2-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_target_init(ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9targetBarPiS__l25_kernel_environment, ptr [[DYN_PTR]])545// CHECK2-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP1]], -1546// CHECK2-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]547// CHECK2: user_code.entry:548// CHECK2-NEXT: [[TMP2:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1:[0-9]+]])549// CHECK2-NEXT: [[TMP3:%.*]] = getelementptr inbounds [2 x ptr], ptr [[CAPTURED_VARS_ADDRS]], i32 0, i32 0550// CHECK2-NEXT: store ptr [[PTR1_ADDR]], ptr [[TMP3]], align 4551// CHECK2-NEXT: [[TMP4:%.*]] = getelementptr inbounds [2 x ptr], ptr [[CAPTURED_VARS_ADDRS]], i32 0, i32 1552// CHECK2-NEXT: store ptr [[TMP0]], ptr [[TMP4]], align 4553// CHECK2-NEXT: call void @__kmpc_parallel_51(ptr @[[GLOB1]], i32 [[TMP2]], i32 1, i32 2, i32 -1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9targetBarPiS__l25_omp_outlined, ptr null, ptr [[CAPTURED_VARS_ADDRS]], i32 2)554// CHECK2-NEXT: call void @__kmpc_target_deinit()555// CHECK2-NEXT: ret void556// CHECK2: worker.exit:557// CHECK2-NEXT: ret void558//559//560// CHECK2-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9targetBarPiS__l25_omp_outlined561// CHECK2-SAME: (ptr noalias [[DOTGLOBAL_TID_:%.*]], ptr noalias [[DOTBOUND_TID_:%.*]], ptr nonnull align 4 dereferenceable(4) [[PTR1:%.*]], ptr nonnull align 4 dereferenceable(4) [[PTR2:%.*]]) #[[ATTR1:[0-9]+]] {562// CHECK2-NEXT: entry:563// CHECK2-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 4564// CHECK2-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 4565// CHECK2-NEXT: [[PTR1_ADDR:%.*]] = alloca ptr, align 4566// CHECK2-NEXT: [[PTR2_ADDR:%.*]] = alloca ptr, align 4567// CHECK2-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 4568// CHECK2-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 4569// CHECK2-NEXT: store ptr [[PTR1]], ptr [[PTR1_ADDR]], align 4570// CHECK2-NEXT: store ptr [[PTR2]], ptr [[PTR2_ADDR]], align 4571// CHECK2-NEXT: [[TMP0:%.*]] = load ptr, ptr [[PTR1_ADDR]], align 4572// CHECK2-NEXT: [[TMP1:%.*]] = load ptr, ptr [[PTR2_ADDR]], align 4573// CHECK2-NEXT: [[TMP2:%.*]] = load ptr, ptr [[TMP1]], align 4574// CHECK2-NEXT: [[TMP3:%.*]] = load i32, ptr [[TMP2]], align 4575// CHECK2-NEXT: [[TMP4:%.*]] = load ptr, ptr [[TMP0]], align 4576// CHECK2-NEXT: store i32 [[TMP3]], ptr [[TMP4]], align 4577// CHECK2-NEXT: ret void578//579//580// CHECK2-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3fooi_l39581// CHECK2-SAME: (ptr noalias [[DYN_PTR:%.*]]) #[[ATTR4:[0-9]+]] {582// CHECK2-NEXT: entry:583// CHECK2-NEXT: [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 4584// CHECK2-NEXT: store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 4585// CHECK2-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_target_init(ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3fooi_l39_kernel_environment, ptr [[DYN_PTR]])586// CHECK2-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP0]], -1587// CHECK2-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]588// CHECK2: user_code.entry:589// CHECK2-NEXT: call void @__kmpc_target_deinit()590// CHECK2-NEXT: ret void591// CHECK2: worker.exit:592// CHECK2-NEXT: ret void593//594//595// CHECK2-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3fooi_l47596// CHECK2-SAME: (ptr noalias [[DYN_PTR:%.*]], i32 [[AA:%.*]]) #[[ATTR4]] {597// CHECK2-NEXT: entry:598// CHECK2-NEXT: [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 4599// CHECK2-NEXT: [[AA_ADDR:%.*]] = alloca i32, align 4600// CHECK2-NEXT: store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 4601// CHECK2-NEXT: store i32 [[AA]], ptr [[AA_ADDR]], align 4602// CHECK2-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_target_init(ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3fooi_l47_kernel_environment, ptr [[DYN_PTR]])603// CHECK2-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP0]], -1604// CHECK2-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]605// CHECK2: user_code.entry:606// CHECK2-NEXT: [[TMP1:%.*]] = load i16, ptr [[AA_ADDR]], align 2607// CHECK2-NEXT: [[CONV:%.*]] = sext i16 [[TMP1]] to i32608// CHECK2-NEXT: [[ADD:%.*]] = add nsw i32 [[CONV]], 1609// CHECK2-NEXT: [[CONV1:%.*]] = trunc i32 [[ADD]] to i16610// CHECK2-NEXT: store i16 [[CONV1]], ptr [[AA_ADDR]], align 2611// CHECK2-NEXT: [[TMP2:%.*]] = load i16, ptr [[AA_ADDR]], align 2612// CHECK2-NEXT: [[CONV2:%.*]] = sext i16 [[TMP2]] to i32613// CHECK2-NEXT: [[ADD3:%.*]] = add nsw i32 [[CONV2]], 2614// CHECK2-NEXT: [[CONV4:%.*]] = trunc i32 [[ADD3]] to i16615// CHECK2-NEXT: store i16 [[CONV4]], ptr [[AA_ADDR]], align 2616// CHECK2-NEXT: call void @__kmpc_target_deinit()617// CHECK2-NEXT: ret void618// CHECK2: worker.exit:619// CHECK2-NEXT: ret void620//621//622// CHECK2-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3fooi_l53623// CHECK2-SAME: (ptr noalias [[DYN_PTR:%.*]], i32 [[A:%.*]], ptr nonnull align 4 dereferenceable(40) [[B:%.*]], i32 [[VLA:%.*]], ptr nonnull align 4 dereferenceable(4) [[BN:%.*]], ptr nonnull align 8 dereferenceable(400) [[C:%.*]], i32 [[VLA1:%.*]], i32 [[VLA3:%.*]], ptr nonnull align 8 dereferenceable(8) [[CN:%.*]], ptr nonnull align 8 dereferenceable(16) [[D:%.*]]) #[[ATTR4]] {624// CHECK2-NEXT: entry:625// CHECK2-NEXT: [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 4626// CHECK2-NEXT: [[A_ADDR:%.*]] = alloca i32, align 4627// CHECK2-NEXT: [[B_ADDR:%.*]] = alloca ptr, align 4628// CHECK2-NEXT: [[VLA_ADDR:%.*]] = alloca i32, align 4629// CHECK2-NEXT: [[BN_ADDR:%.*]] = alloca ptr, align 4630// CHECK2-NEXT: [[C_ADDR:%.*]] = alloca ptr, align 4631// CHECK2-NEXT: [[VLA_ADDR2:%.*]] = alloca i32, align 4632// CHECK2-NEXT: [[VLA_ADDR4:%.*]] = alloca i32, align 4633// CHECK2-NEXT: [[CN_ADDR:%.*]] = alloca ptr, align 4634// CHECK2-NEXT: [[D_ADDR:%.*]] = alloca ptr, align 4635// CHECK2-NEXT: store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 4636// CHECK2-NEXT: store i32 [[A]], ptr [[A_ADDR]], align 4637// CHECK2-NEXT: store ptr [[B]], ptr [[B_ADDR]], align 4638// CHECK2-NEXT: store i32 [[VLA]], ptr [[VLA_ADDR]], align 4639// CHECK2-NEXT: store ptr [[BN]], ptr [[BN_ADDR]], align 4640// CHECK2-NEXT: store ptr [[C]], ptr [[C_ADDR]], align 4641// CHECK2-NEXT: store i32 [[VLA1]], ptr [[VLA_ADDR2]], align 4642// CHECK2-NEXT: store i32 [[VLA3]], ptr [[VLA_ADDR4]], align 4643// CHECK2-NEXT: store ptr [[CN]], ptr [[CN_ADDR]], align 4644// CHECK2-NEXT: store ptr [[D]], ptr [[D_ADDR]], align 4645// CHECK2-NEXT: [[TMP0:%.*]] = load ptr, ptr [[B_ADDR]], align 4646// CHECK2-NEXT: [[TMP1:%.*]] = load i32, ptr [[VLA_ADDR]], align 4647// CHECK2-NEXT: [[TMP2:%.*]] = load ptr, ptr [[BN_ADDR]], align 4648// CHECK2-NEXT: [[TMP3:%.*]] = load ptr, ptr [[C_ADDR]], align 4649// CHECK2-NEXT: [[TMP4:%.*]] = load i32, ptr [[VLA_ADDR2]], align 4650// CHECK2-NEXT: [[TMP5:%.*]] = load i32, ptr [[VLA_ADDR4]], align 4651// CHECK2-NEXT: [[TMP6:%.*]] = load ptr, ptr [[CN_ADDR]], align 4652// CHECK2-NEXT: [[TMP7:%.*]] = load ptr, ptr [[D_ADDR]], align 4653// CHECK2-NEXT: [[TMP8:%.*]] = call i32 @__kmpc_target_init(ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3fooi_l53_kernel_environment, ptr [[DYN_PTR]])654// CHECK2-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP8]], -1655// CHECK2-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]656// CHECK2: user_code.entry:657// CHECK2-NEXT: [[TMP9:%.*]] = load i32, ptr [[A_ADDR]], align 4658// CHECK2-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP9]], 1659// CHECK2-NEXT: store i32 [[ADD]], ptr [[A_ADDR]], align 4660// CHECK2-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x float], ptr [[TMP0]], i32 0, i32 2661// CHECK2-NEXT: [[TMP10:%.*]] = load float, ptr [[ARRAYIDX]], align 4662// CHECK2-NEXT: [[CONV:%.*]] = fpext float [[TMP10]] to double663// CHECK2-NEXT: [[ADD5:%.*]] = fadd double [[CONV]], 1.000000e+00664// CHECK2-NEXT: [[CONV6:%.*]] = fptrunc double [[ADD5]] to float665// CHECK2-NEXT: store float [[CONV6]], ptr [[ARRAYIDX]], align 4666// CHECK2-NEXT: [[ARRAYIDX7:%.*]] = getelementptr inbounds float, ptr [[TMP2]], i32 3667// CHECK2-NEXT: [[TMP11:%.*]] = load float, ptr [[ARRAYIDX7]], align 4668// CHECK2-NEXT: [[CONV8:%.*]] = fpext float [[TMP11]] to double669// CHECK2-NEXT: [[ADD9:%.*]] = fadd double [[CONV8]], 1.000000e+00670// CHECK2-NEXT: [[CONV10:%.*]] = fptrunc double [[ADD9]] to float671// CHECK2-NEXT: store float [[CONV10]], ptr [[ARRAYIDX7]], align 4672// CHECK2-NEXT: [[ARRAYIDX11:%.*]] = getelementptr inbounds [5 x [10 x double]], ptr [[TMP3]], i32 0, i32 1673// CHECK2-NEXT: [[ARRAYIDX12:%.*]] = getelementptr inbounds [10 x double], ptr [[ARRAYIDX11]], i32 0, i32 2674// CHECK2-NEXT: [[TMP12:%.*]] = load double, ptr [[ARRAYIDX12]], align 8675// CHECK2-NEXT: [[ADD13:%.*]] = fadd double [[TMP12]], 1.000000e+00676// CHECK2-NEXT: store double [[ADD13]], ptr [[ARRAYIDX12]], align 8677// CHECK2-NEXT: [[TMP13:%.*]] = mul nsw i32 1, [[TMP5]]678// CHECK2-NEXT: [[ARRAYIDX14:%.*]] = getelementptr inbounds double, ptr [[TMP6]], i32 [[TMP13]]679// CHECK2-NEXT: [[ARRAYIDX15:%.*]] = getelementptr inbounds double, ptr [[ARRAYIDX14]], i32 3680// CHECK2-NEXT: [[TMP14:%.*]] = load double, ptr [[ARRAYIDX15]], align 8681// CHECK2-NEXT: [[ADD16:%.*]] = fadd double [[TMP14]], 1.000000e+00682// CHECK2-NEXT: store double [[ADD16]], ptr [[ARRAYIDX15]], align 8683// CHECK2-NEXT: [[X:%.*]] = getelementptr inbounds nuw [[STRUCT_TT:%.*]], ptr [[TMP7]], i32 0, i32 0684// CHECK2-NEXT: [[TMP15:%.*]] = load i64, ptr [[X]], align 8685// CHECK2-NEXT: [[ADD17:%.*]] = add nsw i64 [[TMP15]], 1686// CHECK2-NEXT: store i64 [[ADD17]], ptr [[X]], align 8687// CHECK2-NEXT: [[Y:%.*]] = getelementptr inbounds nuw [[STRUCT_TT]], ptr [[TMP7]], i32 0, i32 1688// CHECK2-NEXT: [[TMP16:%.*]] = load i8, ptr [[Y]], align 8689// CHECK2-NEXT: [[CONV18:%.*]] = sext i8 [[TMP16]] to i32690// CHECK2-NEXT: [[ADD19:%.*]] = add nsw i32 [[CONV18]], 1691// CHECK2-NEXT: [[CONV20:%.*]] = trunc i32 [[ADD19]] to i8692// CHECK2-NEXT: store i8 [[CONV20]], ptr [[Y]], align 8693// CHECK2-NEXT: [[CALL:%.*]] = call nonnull align 8 dereferenceable(8) ptr @_ZN2TTIxcEixEi(ptr nonnull align 8 dereferenceable(16) [[TMP7]], i32 0) #[[ATTR10:[0-9]+]]694// CHECK2-NEXT: [[TMP17:%.*]] = load i64, ptr [[CALL]], align 8695// CHECK2-NEXT: [[ADD21:%.*]] = add nsw i64 [[TMP17]], 1696// CHECK2-NEXT: store i64 [[ADD21]], ptr [[CALL]], align 8697// CHECK2-NEXT: call void @__kmpc_target_deinit()698// CHECK2-NEXT: ret void699// CHECK2: worker.exit:700// CHECK2-NEXT: ret void701//702//703// CHECK2-LABEL: define {{[^@]+}}@_ZN2TTIxcEixEi704// CHECK2-SAME: (ptr nonnull align 8 dereferenceable(16) [[THIS:%.*]], i32 [[I:%.*]]) #[[ATTR5:[0-9]+]] comdat align 2 {705// CHECK2-NEXT: entry:706// CHECK2-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 4707// CHECK2-NEXT: [[I_ADDR:%.*]] = alloca i32, align 4708// CHECK2-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 4709// CHECK2-NEXT: store i32 [[I]], ptr [[I_ADDR]], align 4710// CHECK2-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 4711// CHECK2-NEXT: [[X:%.*]] = getelementptr inbounds nuw [[STRUCT_TT:%.*]], ptr [[THIS1]], i32 0, i32 0712// CHECK2-NEXT: ret ptr [[X]]713//714//715// CHECK2-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l90716// CHECK2-SAME: (ptr noalias [[DYN_PTR:%.*]], i32 [[A:%.*]], i32 [[AA:%.*]], i32 [[AAA:%.*]], ptr nonnull align 4 dereferenceable(40) [[B:%.*]]) #[[ATTR4]] {717// CHECK2-NEXT: entry:718// CHECK2-NEXT: [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 4719// CHECK2-NEXT: [[A_ADDR:%.*]] = alloca i32, align 4720// CHECK2-NEXT: [[AA_ADDR:%.*]] = alloca i32, align 4721// CHECK2-NEXT: [[AAA_ADDR:%.*]] = alloca i32, align 4722// CHECK2-NEXT: [[B_ADDR:%.*]] = alloca ptr, align 4723// CHECK2-NEXT: store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 4724// CHECK2-NEXT: store i32 [[A]], ptr [[A_ADDR]], align 4725// CHECK2-NEXT: store i32 [[AA]], ptr [[AA_ADDR]], align 4726// CHECK2-NEXT: store i32 [[AAA]], ptr [[AAA_ADDR]], align 4727// CHECK2-NEXT: store ptr [[B]], ptr [[B_ADDR]], align 4728// CHECK2-NEXT: [[TMP0:%.*]] = load ptr, ptr [[B_ADDR]], align 4729// CHECK2-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_target_init(ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l90_kernel_environment, ptr [[DYN_PTR]])730// CHECK2-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP1]], -1731// CHECK2-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]732// CHECK2: user_code.entry:733// CHECK2-NEXT: [[TMP2:%.*]] = load i32, ptr [[A_ADDR]], align 4734// CHECK2-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP2]], 1735// CHECK2-NEXT: store i32 [[ADD]], ptr [[A_ADDR]], align 4736// CHECK2-NEXT: [[TMP3:%.*]] = load i16, ptr [[AA_ADDR]], align 2737// CHECK2-NEXT: [[CONV:%.*]] = sext i16 [[TMP3]] to i32738// CHECK2-NEXT: [[ADD1:%.*]] = add nsw i32 [[CONV]], 1739// CHECK2-NEXT: [[CONV2:%.*]] = trunc i32 [[ADD1]] to i16740// CHECK2-NEXT: store i16 [[CONV2]], ptr [[AA_ADDR]], align 2741// CHECK2-NEXT: [[TMP4:%.*]] = load i8, ptr [[AAA_ADDR]], align 1742// CHECK2-NEXT: [[CONV3:%.*]] = sext i8 [[TMP4]] to i32743// CHECK2-NEXT: [[ADD4:%.*]] = add nsw i32 [[CONV3]], 1744// CHECK2-NEXT: [[CONV5:%.*]] = trunc i32 [[ADD4]] to i8745// CHECK2-NEXT: store i8 [[CONV5]], ptr [[AAA_ADDR]], align 1746// CHECK2-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x i32], ptr [[TMP0]], i32 0, i32 2747// CHECK2-NEXT: [[TMP5:%.*]] = load i32, ptr [[ARRAYIDX]], align 4748// CHECK2-NEXT: [[ADD6:%.*]] = add nsw i32 [[TMP5]], 1749// CHECK2-NEXT: store i32 [[ADD6]], ptr [[ARRAYIDX]], align 4750// CHECK2-NEXT: call void @__kmpc_target_deinit()751// CHECK2-NEXT: ret void752// CHECK2: worker.exit:753// CHECK2-NEXT: ret void754//755//756// CHECK2-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2S12r1Ei_l108757// CHECK2-SAME: (ptr noalias [[DYN_PTR:%.*]], ptr [[THIS:%.*]], i32 [[B:%.*]], i32 [[VLA:%.*]], i32 [[VLA1:%.*]], ptr nonnull align 2 dereferenceable(2) [[C:%.*]]) #[[ATTR4]] {758// CHECK2-NEXT: entry:759// CHECK2-NEXT: [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 4760// CHECK2-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 4761// CHECK2-NEXT: [[B_ADDR:%.*]] = alloca i32, align 4762// CHECK2-NEXT: [[VLA_ADDR:%.*]] = alloca i32, align 4763// CHECK2-NEXT: [[VLA_ADDR2:%.*]] = alloca i32, align 4764// CHECK2-NEXT: [[C_ADDR:%.*]] = alloca ptr, align 4765// CHECK2-NEXT: store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 4766// CHECK2-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 4767// CHECK2-NEXT: store i32 [[B]], ptr [[B_ADDR]], align 4768// CHECK2-NEXT: store i32 [[VLA]], ptr [[VLA_ADDR]], align 4769// CHECK2-NEXT: store i32 [[VLA1]], ptr [[VLA_ADDR2]], align 4770// CHECK2-NEXT: store ptr [[C]], ptr [[C_ADDR]], align 4771// CHECK2-NEXT: [[TMP0:%.*]] = load ptr, ptr [[THIS_ADDR]], align 4772// CHECK2-NEXT: [[TMP1:%.*]] = load i32, ptr [[VLA_ADDR]], align 4773// CHECK2-NEXT: [[TMP2:%.*]] = load i32, ptr [[VLA_ADDR2]], align 4774// CHECK2-NEXT: [[TMP3:%.*]] = load ptr, ptr [[C_ADDR]], align 4775// CHECK2-NEXT: [[TMP4:%.*]] = call i32 @__kmpc_target_init(ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2S12r1Ei_l108_kernel_environment, ptr [[DYN_PTR]])776// CHECK2-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP4]], -1777// CHECK2-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]778// CHECK2: user_code.entry:779// CHECK2-NEXT: [[TMP5:%.*]] = load i32, ptr [[B_ADDR]], align 4780// CHECK2-NEXT: [[CONV:%.*]] = sitofp i32 [[TMP5]] to double781// CHECK2-NEXT: [[ADD:%.*]] = fadd double [[CONV]], 1.500000e+00782// CHECK2-NEXT: [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_S1:%.*]], ptr [[TMP0]], i32 0, i32 0783// CHECK2-NEXT: store double [[ADD]], ptr [[A]], align 8784// CHECK2-NEXT: [[A3:%.*]] = getelementptr inbounds nuw [[STRUCT_S1]], ptr [[TMP0]], i32 0, i32 0785// CHECK2-NEXT: [[TMP6:%.*]] = load double, ptr [[A3]], align 8786// CHECK2-NEXT: [[INC:%.*]] = fadd double [[TMP6]], 1.000000e+00787// CHECK2-NEXT: store double [[INC]], ptr [[A3]], align 8788// CHECK2-NEXT: [[CONV4:%.*]] = fptosi double [[INC]] to i16789// CHECK2-NEXT: [[TMP7:%.*]] = mul nsw i32 1, [[TMP2]]790// CHECK2-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds i16, ptr [[TMP3]], i32 [[TMP7]]791// CHECK2-NEXT: [[ARRAYIDX5:%.*]] = getelementptr inbounds i16, ptr [[ARRAYIDX]], i32 1792// CHECK2-NEXT: store i16 [[CONV4]], ptr [[ARRAYIDX5]], align 2793// CHECK2-NEXT: [[A6:%.*]] = getelementptr inbounds nuw [[STRUCT_S1]], ptr [[TMP0]], i32 0, i32 0794// CHECK2-NEXT: [[TMP8:%.*]] = load double, ptr [[A6]], align 8795// CHECK2-NEXT: [[CONV7:%.*]] = fptosi double [[TMP8]] to i32796// CHECK2-NEXT: [[A8:%.*]] = getelementptr inbounds nuw [[STRUCT_S1]], ptr [[TMP0]], i32 0, i32 0797// CHECK2-NEXT: [[CALL:%.*]] = call i32 @_Z3baziRd(i32 [[CONV7]], ptr nonnull align 8 dereferenceable(8) [[A8]]) #[[ATTR10]]798// CHECK2-NEXT: call void @__kmpc_target_deinit()799// CHECK2-NEXT: ret void800// CHECK2: worker.exit:801// CHECK2-NEXT: ret void802//803//804// CHECK2-LABEL: define {{[^@]+}}@_Z3baziRd805// CHECK2-SAME: (i32 [[F1:%.*]], ptr nonnull align 8 dereferenceable(8) [[A:%.*]]) #[[ATTR5]] {806// CHECK2-NEXT: entry:807// CHECK2-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 4808// CHECK2-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [2 x ptr], align 4809// CHECK2-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])810// CHECK2-NEXT: [[F:%.*]] = call align 8 ptr @__kmpc_alloc_shared(i32 4)811// CHECK2-NEXT: store i32 [[F1]], ptr [[F]], align 4812// CHECK2-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 4813// CHECK2-NEXT: [[TMP1:%.*]] = load ptr, ptr [[A_ADDR]], align 4814// CHECK2-NEXT: [[TMP2:%.*]] = getelementptr inbounds [2 x ptr], ptr [[CAPTURED_VARS_ADDRS]], i32 0, i32 0815// CHECK2-NEXT: store ptr [[F]], ptr [[TMP2]], align 4816// CHECK2-NEXT: [[TMP3:%.*]] = getelementptr inbounds [2 x ptr], ptr [[CAPTURED_VARS_ADDRS]], i32 0, i32 1817// CHECK2-NEXT: store ptr [[TMP1]], ptr [[TMP3]], align 4818// CHECK2-NEXT: call void @__kmpc_parallel_51(ptr @[[GLOB1]], i32 [[TMP0]], i32 1, i32 -1, i32 -1, ptr @_Z3baziRd_omp_outlined, ptr @_Z3baziRd_omp_outlined_wrapper, ptr [[CAPTURED_VARS_ADDRS]], i32 2)819// CHECK2-NEXT: [[TMP4:%.*]] = load i32, ptr [[F]], align 4820// CHECK2-NEXT: call void @__kmpc_free_shared(ptr [[F]], i32 4)821// CHECK2-NEXT: ret i32 [[TMP4]]822//823//824// CHECK2-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z16unreachable_callv_l142825// CHECK2-SAME: (ptr noalias [[DYN_PTR:%.*]]) #[[ATTR4]] {826// CHECK2-NEXT: entry:827// CHECK2-NEXT: [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 4828// CHECK2-NEXT: store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 4829// CHECK2-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_target_init(ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z16unreachable_callv_l142_kernel_environment, ptr [[DYN_PTR]])830// CHECK2-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP0]], -1831// CHECK2-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]832// CHECK2: user_code.entry:833// CHECK2-NEXT: call void @_Z6asserti(i32 0) #[[ATTR11:[0-9]+]]834// CHECK2-NEXT: unreachable835// CHECK2: worker.exit:836// CHECK2-NEXT: ret void837// CHECK2: 1:838// CHECK2-NEXT: call void @__kmpc_target_deinit()839// CHECK2-NEXT: ret void840//841//842// CHECK2-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l74843// CHECK2-SAME: (ptr noalias [[DYN_PTR:%.*]], i32 [[A:%.*]], i32 [[AA:%.*]], ptr nonnull align 4 dereferenceable(40) [[B:%.*]]) #[[ATTR4]] {844// CHECK2-NEXT: entry:845// CHECK2-NEXT: [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 4846// CHECK2-NEXT: [[A_ADDR:%.*]] = alloca i32, align 4847// CHECK2-NEXT: [[AA_ADDR:%.*]] = alloca i32, align 4848// CHECK2-NEXT: [[B_ADDR:%.*]] = alloca ptr, align 4849// CHECK2-NEXT: store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 4850// CHECK2-NEXT: store i32 [[A]], ptr [[A_ADDR]], align 4851// CHECK2-NEXT: store i32 [[AA]], ptr [[AA_ADDR]], align 4852// CHECK2-NEXT: store ptr [[B]], ptr [[B_ADDR]], align 4853// CHECK2-NEXT: [[TMP0:%.*]] = load ptr, ptr [[B_ADDR]], align 4854// CHECK2-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_target_init(ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l74_kernel_environment, ptr [[DYN_PTR]])855// CHECK2-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP1]], -1856// CHECK2-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]857// CHECK2: user_code.entry:858// CHECK2-NEXT: [[TMP2:%.*]] = load i32, ptr [[A_ADDR]], align 4859// CHECK2-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP2]], 1860// CHECK2-NEXT: store i32 [[ADD]], ptr [[A_ADDR]], align 4861// CHECK2-NEXT: [[TMP3:%.*]] = load i16, ptr [[AA_ADDR]], align 2862// CHECK2-NEXT: [[CONV:%.*]] = sext i16 [[TMP3]] to i32863// CHECK2-NEXT: [[ADD1:%.*]] = add nsw i32 [[CONV]], 1864// CHECK2-NEXT: [[CONV2:%.*]] = trunc i32 [[ADD1]] to i16865// CHECK2-NEXT: store i16 [[CONV2]], ptr [[AA_ADDR]], align 2866// CHECK2-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x i32], ptr [[TMP0]], i32 0, i32 2867// CHECK2-NEXT: [[TMP4:%.*]] = load i32, ptr [[ARRAYIDX]], align 4868// CHECK2-NEXT: [[ADD3:%.*]] = add nsw i32 [[TMP4]], 1869// CHECK2-NEXT: store i32 [[ADD3]], ptr [[ARRAYIDX]], align 4870// CHECK2-NEXT: call void @__kmpc_target_deinit()871// CHECK2-NEXT: ret void872// CHECK2: worker.exit:873// CHECK2-NEXT: ret void874//875//876// CHECK2-LABEL: define {{[^@]+}}@_Z3baziRd_omp_outlined877// CHECK2-SAME: (ptr noalias [[DOTGLOBAL_TID_:%.*]], ptr noalias [[DOTBOUND_TID_:%.*]], ptr nonnull align 4 dereferenceable(4) [[F:%.*]], ptr nonnull align 8 dereferenceable(8) [[A:%.*]]) #[[ATTR1]] {878// CHECK2-NEXT: entry:879// CHECK2-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 4880// CHECK2-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 4881// CHECK2-NEXT: [[F_ADDR:%.*]] = alloca ptr, align 4882// CHECK2-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 4883// CHECK2-NEXT: [[TMP:%.*]] = alloca ptr, align 4884// CHECK2-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 4885// CHECK2-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 4886// CHECK2-NEXT: store ptr [[F]], ptr [[F_ADDR]], align 4887// CHECK2-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 4888// CHECK2-NEXT: [[TMP0:%.*]] = load ptr, ptr [[F_ADDR]], align 4889// CHECK2-NEXT: [[TMP1:%.*]] = load ptr, ptr [[A_ADDR]], align 4890// CHECK2-NEXT: store ptr [[TMP1]], ptr [[TMP]], align 4891// CHECK2-NEXT: [[TMP2:%.*]] = load ptr, ptr [[TMP]], align 4892// CHECK2-NEXT: [[TMP3:%.*]] = load double, ptr [[TMP2]], align 8893// CHECK2-NEXT: [[ADD:%.*]] = fadd double 2.000000e+00, [[TMP3]]894// CHECK2-NEXT: [[CONV:%.*]] = fptosi double [[ADD]] to i32895// CHECK2-NEXT: store i32 [[CONV]], ptr [[TMP0]], align 4896// CHECK2-NEXT: ret void897//898//899// CHECK2-LABEL: define {{[^@]+}}@_Z3baziRd_omp_outlined_wrapper900// CHECK2-SAME: (i16 zeroext [[TMP0:%.*]], i32 [[TMP1:%.*]]) #[[ATTR8:[0-9]+]] {901// CHECK2-NEXT: entry:902// CHECK2-NEXT: [[DOTADDR:%.*]] = alloca i16, align 2903// CHECK2-NEXT: [[DOTADDR1:%.*]] = alloca i32, align 4904// CHECK2-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4905// CHECK2-NEXT: [[GLOBAL_ARGS:%.*]] = alloca ptr, align 4906// CHECK2-NEXT: store i16 [[TMP0]], ptr [[DOTADDR]], align 2907// CHECK2-NEXT: store i32 [[TMP1]], ptr [[DOTADDR1]], align 4908// CHECK2-NEXT: store i32 0, ptr [[DOTZERO_ADDR]], align 4909// CHECK2-NEXT: call void @__kmpc_get_shared_variables(ptr [[GLOBAL_ARGS]])910// CHECK2-NEXT: [[TMP2:%.*]] = load ptr, ptr [[GLOBAL_ARGS]], align 4911// CHECK2-NEXT: [[TMP3:%.*]] = getelementptr inbounds ptr, ptr [[TMP2]], i32 0912// CHECK2-NEXT: [[TMP4:%.*]] = load ptr, ptr [[TMP3]], align 4913// CHECK2-NEXT: [[TMP5:%.*]] = getelementptr inbounds ptr, ptr [[TMP2]], i32 1914// CHECK2-NEXT: [[TMP6:%.*]] = load ptr, ptr [[TMP5]], align 4915// CHECK2-NEXT: call void @_Z3baziRd_omp_outlined(ptr [[DOTADDR1]], ptr [[DOTZERO_ADDR]], ptr [[TMP4]], ptr [[TMP6]]) #[[ATTR2:[0-9]+]]916// CHECK2-NEXT: ret void917//918