1984 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// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix=CHECK13// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s4// RUN: %clang_cc1 -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix=CHECK15// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix=CHECK36// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s7// RUN: %clang_cc1 -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix=CHECK38 9// RUN: %clang_cc1 -verify -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"10// RUN: %clang_cc1 -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s11// RUN: %clang_cc1 -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"12// RUN: %clang_cc1 -verify -fopenmp-simd -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"13// RUN: %clang_cc1 -fopenmp-simd -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s14// RUN: %clang_cc1 -fopenmp-simd -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"15// expected-no-diagnostics16#ifndef HEADER17#define HEADER18 19 20 21double Ga = 1.0;22double Gb = 2.0;23double Gc = 3.0;24double Gd = 4.0;25 26int foo(short a, short b, short c, short d){27 static float Sa = 5.0;28 static float Sb = 6.0;29 static float Sc = 7.0;30 static float Sd = 8.0;31 32 33 // 3 local vars being captured.34 35 36 37 38 // 3 static vars being captured.39 40 41 42 43 // 3 static global vars being captured.44 45 46 47 48 // Capture b, Gb, Sb, Gc, c, Sc, d, Gd, Sd49 #pragma omp target if(Ga>0.0 && a>0 && Sa>0.0)50 {51 b += 1;52 Gb += 1.0;53 Sb += 1.0;54 55 // The parallel region only uses 3 captures.56 // Capture d, Gd, Sd,57 58 #pragma omp parallel if(Gc>0.0 && c>0 && Sc>0.0)59 {60 d += 1;61 Gd += 1.0;62 Sd += 1.0;63 }64 }65 return a + b + c + d + (int)Sa + (int)Sb + (int)Sc + (int)Sd;66}67 68int bar(short a, short b, short c, short d){69 static float Sa = 9.0;70 static float Sb = 10.0;71 static float Sc = 11.0;72 static float Sd = 12.0;73 74 // Capture a, b, c, d75 #pragma omp parallel76 {77 78 // 3 local vars being captured.79 80 81 82 83 // 3 static vars being captured.84 85 86 87 88 // 3 static global vars being captured.89 90 91 92 93 // Capture b, Gb, Sb, Gc, c, Sc, d, Gd, Sd94 #pragma omp target if(Ga>0.0 && a>0 && Sa>0.0)95 {96 b += 1;97 Gb += 1.0;98 Sb += 1.0;99 100 101 // Capture d, Gd, Sd102 #pragma omp parallel if(Gc>0.0 && c>0 && Sc>0.0)103 {104 d += 1;105 Gd += 1.0;106 Sd += 1.0;107 }108 }109 }110 return a + b + c + d + (int)Sa + (int)Sb + (int)Sc + (int)Sd;111}112 113///114/// Tests with template functions.115///116 117 118template<typename T>119int tbar(T a, T b, T c, T d){120 static float Sa = 17.0;121 static float Sb = 18.0;122 static float Sc = 19.0;123 static float Sd = 20.0;124 125 // Capture a, b, c, d126 #pragma omp parallel127 {128 129 // 3 local vars being captured.130 131 132 133 134 // 3 static vars being captured.135 136 137 138 139 // 3 static global vars being captured.140 141 142 143 144 // Capture b, Gb, Sb, Gc, c, Sc, d, Gd, Sd145 #pragma omp target if(Ga>0.0 && a>0 && Sa>0.0)146 {147 b += 1;148 Gb += 1.0;149 Sb += 1.0;150 151 152 // Capture d, Gd, Sd153 #pragma omp parallel if(Gc>0.0 && c>0 && Sc>0.0)154 {155 d += 1;156 Gd += 1.0;157 Sd += 1.0;158 }159 }160 }161 return a + b + c + d + (int)Sa + (int)Sb + (int)Sc + (int)Sd;162}163 164int tbar2(short a, short b, short c, short d){165 return tbar(a, b, c, d);166}167 168#endif169// CHECK1-LABEL: define {{[^@]+}}@_Z3foossss170// CHECK1-SAME: (i16 noundef signext [[A:%.*]], i16 noundef signext [[B:%.*]], i16 noundef signext [[C:%.*]], i16 noundef signext [[D:%.*]]) #[[ATTR0:[0-9]+]] {171// CHECK1-NEXT: entry:172// CHECK1-NEXT: [[A_ADDR:%.*]] = alloca i16, align 2173// CHECK1-NEXT: [[B_ADDR:%.*]] = alloca i16, align 2174// CHECK1-NEXT: [[C_ADDR:%.*]] = alloca i16, align 2175// CHECK1-NEXT: [[D_ADDR:%.*]] = alloca i16, align 2176// CHECK1-NEXT: [[B_CASTED:%.*]] = alloca i64, align 8177// CHECK1-NEXT: [[GB_CASTED:%.*]] = alloca i64, align 8178// CHECK1-NEXT: [[SB_CASTED:%.*]] = alloca i64, align 8179// CHECK1-NEXT: [[GC_CASTED:%.*]] = alloca i64, align 8180// CHECK1-NEXT: [[C_CASTED:%.*]] = alloca i64, align 8181// CHECK1-NEXT: [[SC_CASTED:%.*]] = alloca i64, align 8182// CHECK1-NEXT: [[D_CASTED:%.*]] = alloca i64, align 8183// CHECK1-NEXT: [[GD_CASTED:%.*]] = alloca i64, align 8184// CHECK1-NEXT: [[SD_CASTED:%.*]] = alloca i64, align 8185// CHECK1-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [9 x ptr], align 8186// CHECK1-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [9 x ptr], align 8187// CHECK1-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [9 x ptr], align 8188// CHECK1-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8189// CHECK1-NEXT: store i16 [[A]], ptr [[A_ADDR]], align 2190// CHECK1-NEXT: store i16 [[B]], ptr [[B_ADDR]], align 2191// CHECK1-NEXT: store i16 [[C]], ptr [[C_ADDR]], align 2192// CHECK1-NEXT: store i16 [[D]], ptr [[D_ADDR]], align 2193// CHECK1-NEXT: [[TMP0:%.*]] = load i16, ptr [[B_ADDR]], align 2194// CHECK1-NEXT: store i16 [[TMP0]], ptr [[B_CASTED]], align 2195// CHECK1-NEXT: [[TMP1:%.*]] = load i64, ptr [[B_CASTED]], align 8196// CHECK1-NEXT: [[TMP2:%.*]] = load double, ptr @Gb, align 8197// CHECK1-NEXT: store double [[TMP2]], ptr [[GB_CASTED]], align 8198// CHECK1-NEXT: [[TMP3:%.*]] = load i64, ptr [[GB_CASTED]], align 8199// CHECK1-NEXT: [[TMP4:%.*]] = load float, ptr @_ZZ3foossssE2Sb, align 4200// CHECK1-NEXT: store float [[TMP4]], ptr [[SB_CASTED]], align 4201// CHECK1-NEXT: [[TMP5:%.*]] = load i64, ptr [[SB_CASTED]], align 8202// CHECK1-NEXT: [[TMP6:%.*]] = load double, ptr @Gc, align 8203// CHECK1-NEXT: store double [[TMP6]], ptr [[GC_CASTED]], align 8204// CHECK1-NEXT: [[TMP7:%.*]] = load i64, ptr [[GC_CASTED]], align 8205// CHECK1-NEXT: [[TMP8:%.*]] = load i16, ptr [[C_ADDR]], align 2206// CHECK1-NEXT: store i16 [[TMP8]], ptr [[C_CASTED]], align 2207// CHECK1-NEXT: [[TMP9:%.*]] = load i64, ptr [[C_CASTED]], align 8208// CHECK1-NEXT: [[TMP10:%.*]] = load float, ptr @_ZZ3foossssE2Sc, align 4209// CHECK1-NEXT: store float [[TMP10]], ptr [[SC_CASTED]], align 4210// CHECK1-NEXT: [[TMP11:%.*]] = load i64, ptr [[SC_CASTED]], align 8211// CHECK1-NEXT: [[TMP12:%.*]] = load i16, ptr [[D_ADDR]], align 2212// CHECK1-NEXT: store i16 [[TMP12]], ptr [[D_CASTED]], align 2213// CHECK1-NEXT: [[TMP13:%.*]] = load i64, ptr [[D_CASTED]], align 8214// CHECK1-NEXT: [[TMP14:%.*]] = load double, ptr @Gd, align 8215// CHECK1-NEXT: store double [[TMP14]], ptr [[GD_CASTED]], align 8216// CHECK1-NEXT: [[TMP15:%.*]] = load i64, ptr [[GD_CASTED]], align 8217// CHECK1-NEXT: [[TMP16:%.*]] = load float, ptr @_ZZ3foossssE2Sd, align 4218// CHECK1-NEXT: store float [[TMP16]], ptr [[SD_CASTED]], align 4219// CHECK1-NEXT: [[TMP17:%.*]] = load i64, ptr [[SD_CASTED]], align 8220// CHECK1-NEXT: [[TMP18:%.*]] = load double, ptr @Ga, align 8221// CHECK1-NEXT: [[CMP:%.*]] = fcmp ogt double [[TMP18]], 0.000000e+00222// CHECK1-NEXT: br i1 [[CMP]], label [[LAND_LHS_TRUE:%.*]], label [[OMP_IF_ELSE:%.*]]223// CHECK1: land.lhs.true:224// CHECK1-NEXT: [[TMP19:%.*]] = load i16, ptr [[A_ADDR]], align 2225// CHECK1-NEXT: [[CONV:%.*]] = sext i16 [[TMP19]] to i32226// CHECK1-NEXT: [[CMP1:%.*]] = icmp sgt i32 [[CONV]], 0227// CHECK1-NEXT: br i1 [[CMP1]], label [[LAND_LHS_TRUE2:%.*]], label [[OMP_IF_ELSE]]228// CHECK1: land.lhs.true2:229// CHECK1-NEXT: [[TMP20:%.*]] = load float, ptr @_ZZ3foossssE2Sa, align 4230// CHECK1-NEXT: [[CONV3:%.*]] = fpext float [[TMP20]] to double231// CHECK1-NEXT: [[CMP4:%.*]] = fcmp ogt double [[CONV3]], 0.000000e+00232// CHECK1-NEXT: br i1 [[CMP4]], label [[OMP_IF_THEN:%.*]], label [[OMP_IF_ELSE]]233// CHECK1: omp_if.then:234// CHECK1-NEXT: [[TMP21:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0235// CHECK1-NEXT: store i64 [[TMP1]], ptr [[TMP21]], align 8236// CHECK1-NEXT: [[TMP22:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0237// CHECK1-NEXT: store i64 [[TMP1]], ptr [[TMP22]], align 8238// CHECK1-NEXT: [[TMP23:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 0239// CHECK1-NEXT: store ptr null, ptr [[TMP23]], align 8240// CHECK1-NEXT: [[TMP24:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 1241// CHECK1-NEXT: store i64 [[TMP3]], ptr [[TMP24]], align 8242// CHECK1-NEXT: [[TMP25:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 1243// CHECK1-NEXT: store i64 [[TMP3]], ptr [[TMP25]], align 8244// CHECK1-NEXT: [[TMP26:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 1245// CHECK1-NEXT: store ptr null, ptr [[TMP26]], align 8246// CHECK1-NEXT: [[TMP27:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 2247// CHECK1-NEXT: store i64 [[TMP5]], ptr [[TMP27]], align 8248// CHECK1-NEXT: [[TMP28:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 2249// CHECK1-NEXT: store i64 [[TMP5]], ptr [[TMP28]], align 8250// CHECK1-NEXT: [[TMP29:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 2251// CHECK1-NEXT: store ptr null, ptr [[TMP29]], align 8252// CHECK1-NEXT: [[TMP30:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 3253// CHECK1-NEXT: store i64 [[TMP7]], ptr [[TMP30]], align 8254// CHECK1-NEXT: [[TMP31:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 3255// CHECK1-NEXT: store i64 [[TMP7]], ptr [[TMP31]], align 8256// CHECK1-NEXT: [[TMP32:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 3257// CHECK1-NEXT: store ptr null, ptr [[TMP32]], align 8258// CHECK1-NEXT: [[TMP33:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 4259// CHECK1-NEXT: store i64 [[TMP9]], ptr [[TMP33]], align 8260// CHECK1-NEXT: [[TMP34:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 4261// CHECK1-NEXT: store i64 [[TMP9]], ptr [[TMP34]], align 8262// CHECK1-NEXT: [[TMP35:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 4263// CHECK1-NEXT: store ptr null, ptr [[TMP35]], align 8264// CHECK1-NEXT: [[TMP36:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 5265// CHECK1-NEXT: store i64 [[TMP11]], ptr [[TMP36]], align 8266// CHECK1-NEXT: [[TMP37:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 5267// CHECK1-NEXT: store i64 [[TMP11]], ptr [[TMP37]], align 8268// CHECK1-NEXT: [[TMP38:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 5269// CHECK1-NEXT: store ptr null, ptr [[TMP38]], align 8270// CHECK1-NEXT: [[TMP39:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 6271// CHECK1-NEXT: store i64 [[TMP13]], ptr [[TMP39]], align 8272// CHECK1-NEXT: [[TMP40:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 6273// CHECK1-NEXT: store i64 [[TMP13]], ptr [[TMP40]], align 8274// CHECK1-NEXT: [[TMP41:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 6275// CHECK1-NEXT: store ptr null, ptr [[TMP41]], align 8276// CHECK1-NEXT: [[TMP42:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 7277// CHECK1-NEXT: store i64 [[TMP15]], ptr [[TMP42]], align 8278// CHECK1-NEXT: [[TMP43:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 7279// CHECK1-NEXT: store i64 [[TMP15]], ptr [[TMP43]], align 8280// CHECK1-NEXT: [[TMP44:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 7281// CHECK1-NEXT: store ptr null, ptr [[TMP44]], align 8282// CHECK1-NEXT: [[TMP45:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 8283// CHECK1-NEXT: store i64 [[TMP17]], ptr [[TMP45]], align 8284// CHECK1-NEXT: [[TMP46:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 8285// CHECK1-NEXT: store i64 [[TMP17]], ptr [[TMP46]], align 8286// CHECK1-NEXT: [[TMP47:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 8287// CHECK1-NEXT: store ptr null, ptr [[TMP47]], align 8288// CHECK1-NEXT: [[TMP48:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0289// CHECK1-NEXT: [[TMP49:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0290// CHECK1-NEXT: [[TMP50:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0291// CHECK1-NEXT: store i32 3, ptr [[TMP50]], align 4292// CHECK1-NEXT: [[TMP51:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1293// CHECK1-NEXT: store i32 9, ptr [[TMP51]], align 4294// CHECK1-NEXT: [[TMP52:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2295// CHECK1-NEXT: store ptr [[TMP48]], ptr [[TMP52]], align 8296// CHECK1-NEXT: [[TMP53:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3297// CHECK1-NEXT: store ptr [[TMP49]], ptr [[TMP53]], align 8298// CHECK1-NEXT: [[TMP54:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4299// CHECK1-NEXT: store ptr @.offload_sizes, ptr [[TMP54]], align 8300// CHECK1-NEXT: [[TMP55:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5301// CHECK1-NEXT: store ptr @.offload_maptypes, ptr [[TMP55]], align 8302// CHECK1-NEXT: [[TMP56:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6303// CHECK1-NEXT: store ptr null, ptr [[TMP56]], align 8304// CHECK1-NEXT: [[TMP57:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7305// CHECK1-NEXT: store ptr null, ptr [[TMP57]], align 8306// CHECK1-NEXT: [[TMP58:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8307// CHECK1-NEXT: store i64 0, ptr [[TMP58]], align 8308// CHECK1-NEXT: [[TMP59:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9309// CHECK1-NEXT: store i64 0, ptr [[TMP59]], align 8310// CHECK1-NEXT: [[TMP60:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10311// CHECK1-NEXT: store [3 x i32] [i32 -1, i32 0, i32 0], ptr [[TMP60]], align 4312// CHECK1-NEXT: [[TMP61:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11313// CHECK1-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP61]], align 4314// CHECK1-NEXT: [[TMP62:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12315// CHECK1-NEXT: store i32 0, ptr [[TMP62]], align 4316// CHECK1-NEXT: [[TMP63:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1:[0-9]+]], i64 -1, i32 -1, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3foossss_l49.region_id, ptr [[KERNEL_ARGS]])317// CHECK1-NEXT: [[TMP64:%.*]] = icmp ne i32 [[TMP63]], 0318// CHECK1-NEXT: br i1 [[TMP64]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]319// CHECK1: omp_offload.failed:320// CHECK1-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3foossss_l49(i64 [[TMP1]], i64 [[TMP3]], i64 [[TMP5]], i64 [[TMP7]], i64 [[TMP9]], i64 [[TMP11]], i64 [[TMP13]], i64 [[TMP15]], i64 [[TMP17]]) #[[ATTR2:[0-9]+]]321// CHECK1-NEXT: br label [[OMP_OFFLOAD_CONT]]322// CHECK1: omp_offload.cont:323// CHECK1-NEXT: br label [[OMP_IF_END:%.*]]324// CHECK1: omp_if.else:325// CHECK1-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3foossss_l49(i64 [[TMP1]], i64 [[TMP3]], i64 [[TMP5]], i64 [[TMP7]], i64 [[TMP9]], i64 [[TMP11]], i64 [[TMP13]], i64 [[TMP15]], i64 [[TMP17]]) #[[ATTR2]]326// CHECK1-NEXT: br label [[OMP_IF_END]]327// CHECK1: omp_if.end:328// CHECK1-NEXT: [[TMP65:%.*]] = load i16, ptr [[A_ADDR]], align 2329// CHECK1-NEXT: [[CONV5:%.*]] = sext i16 [[TMP65]] to i32330// CHECK1-NEXT: [[TMP66:%.*]] = load i16, ptr [[B_ADDR]], align 2331// CHECK1-NEXT: [[CONV6:%.*]] = sext i16 [[TMP66]] to i32332// CHECK1-NEXT: [[ADD:%.*]] = add nsw i32 [[CONV5]], [[CONV6]]333// CHECK1-NEXT: [[TMP67:%.*]] = load i16, ptr [[C_ADDR]], align 2334// CHECK1-NEXT: [[CONV7:%.*]] = sext i16 [[TMP67]] to i32335// CHECK1-NEXT: [[ADD8:%.*]] = add nsw i32 [[ADD]], [[CONV7]]336// CHECK1-NEXT: [[TMP68:%.*]] = load i16, ptr [[D_ADDR]], align 2337// CHECK1-NEXT: [[CONV9:%.*]] = sext i16 [[TMP68]] to i32338// CHECK1-NEXT: [[ADD10:%.*]] = add nsw i32 [[ADD8]], [[CONV9]]339// CHECK1-NEXT: [[TMP69:%.*]] = load float, ptr @_ZZ3foossssE2Sa, align 4340// CHECK1-NEXT: [[CONV11:%.*]] = fptosi float [[TMP69]] to i32341// CHECK1-NEXT: [[ADD12:%.*]] = add nsw i32 [[ADD10]], [[CONV11]]342// CHECK1-NEXT: [[TMP70:%.*]] = load float, ptr @_ZZ3foossssE2Sb, align 4343// CHECK1-NEXT: [[CONV13:%.*]] = fptosi float [[TMP70]] to i32344// CHECK1-NEXT: [[ADD14:%.*]] = add nsw i32 [[ADD12]], [[CONV13]]345// CHECK1-NEXT: [[TMP71:%.*]] = load float, ptr @_ZZ3foossssE2Sc, align 4346// CHECK1-NEXT: [[CONV15:%.*]] = fptosi float [[TMP71]] to i32347// CHECK1-NEXT: [[ADD16:%.*]] = add nsw i32 [[ADD14]], [[CONV15]]348// CHECK1-NEXT: [[TMP72:%.*]] = load float, ptr @_ZZ3foossssE2Sd, align 4349// CHECK1-NEXT: [[CONV17:%.*]] = fptosi float [[TMP72]] to i32350// CHECK1-NEXT: [[ADD18:%.*]] = add nsw i32 [[ADD16]], [[CONV17]]351// CHECK1-NEXT: ret i32 [[ADD18]]352//353//354// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3foossss_l49355// CHECK1-SAME: (i64 noundef [[B:%.*]], i64 noundef [[GB:%.*]], i64 noundef [[SB:%.*]], i64 noundef [[GC:%.*]], i64 noundef [[C:%.*]], i64 noundef [[SC:%.*]], i64 noundef [[D:%.*]], i64 noundef [[GD:%.*]], i64 noundef [[SD:%.*]]) #[[ATTR1:[0-9]+]] {356// CHECK1-NEXT: entry:357// CHECK1-NEXT: [[B_ADDR:%.*]] = alloca i64, align 8358// CHECK1-NEXT: [[GB_ADDR:%.*]] = alloca i64, align 8359// CHECK1-NEXT: [[SB_ADDR:%.*]] = alloca i64, align 8360// CHECK1-NEXT: [[GC_ADDR:%.*]] = alloca i64, align 8361// CHECK1-NEXT: [[C_ADDR:%.*]] = alloca i64, align 8362// CHECK1-NEXT: [[SC_ADDR:%.*]] = alloca i64, align 8363// CHECK1-NEXT: [[D_ADDR:%.*]] = alloca i64, align 8364// CHECK1-NEXT: [[GD_ADDR:%.*]] = alloca i64, align 8365// CHECK1-NEXT: [[SD_ADDR:%.*]] = alloca i64, align 8366// CHECK1-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 4367// CHECK1-NEXT: [[DOTBOUND_ZERO_ADDR:%.*]] = alloca i32, align 4368// CHECK1-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])369// CHECK1-NEXT: store i64 [[B]], ptr [[B_ADDR]], align 8370// CHECK1-NEXT: store i64 [[GB]], ptr [[GB_ADDR]], align 8371// CHECK1-NEXT: store i64 [[SB]], ptr [[SB_ADDR]], align 8372// CHECK1-NEXT: store i64 [[GC]], ptr [[GC_ADDR]], align 8373// CHECK1-NEXT: store i64 [[C]], ptr [[C_ADDR]], align 8374// CHECK1-NEXT: store i64 [[SC]], ptr [[SC_ADDR]], align 8375// CHECK1-NEXT: store i64 [[D]], ptr [[D_ADDR]], align 8376// CHECK1-NEXT: store i64 [[GD]], ptr [[GD_ADDR]], align 8377// CHECK1-NEXT: store i64 [[SD]], ptr [[SD_ADDR]], align 8378// CHECK1-NEXT: [[TMP1:%.*]] = load i16, ptr [[B_ADDR]], align 2379// CHECK1-NEXT: [[CONV:%.*]] = sext i16 [[TMP1]] to i32380// CHECK1-NEXT: [[ADD:%.*]] = add nsw i32 [[CONV]], 1381// CHECK1-NEXT: [[CONV1:%.*]] = trunc i32 [[ADD]] to i16382// CHECK1-NEXT: store i16 [[CONV1]], ptr [[B_ADDR]], align 2383// CHECK1-NEXT: [[TMP2:%.*]] = load double, ptr [[GB_ADDR]], align 8384// CHECK1-NEXT: [[ADD2:%.*]] = fadd double [[TMP2]], 1.000000e+00385// CHECK1-NEXT: store double [[ADD2]], ptr [[GB_ADDR]], align 8386// CHECK1-NEXT: [[TMP3:%.*]] = load float, ptr [[SB_ADDR]], align 4387// CHECK1-NEXT: [[CONV3:%.*]] = fpext float [[TMP3]] to double388// CHECK1-NEXT: [[ADD4:%.*]] = fadd double [[CONV3]], 1.000000e+00389// CHECK1-NEXT: [[CONV5:%.*]] = fptrunc double [[ADD4]] to float390// CHECK1-NEXT: store float [[CONV5]], ptr [[SB_ADDR]], align 4391// CHECK1-NEXT: [[TMP4:%.*]] = load double, ptr [[GC_ADDR]], align 8392// CHECK1-NEXT: [[CMP:%.*]] = fcmp ogt double [[TMP4]], 0.000000e+00393// CHECK1-NEXT: br i1 [[CMP]], label [[LAND_LHS_TRUE:%.*]], label [[OMP_IF_ELSE:%.*]]394// CHECK1: land.lhs.true:395// CHECK1-NEXT: [[TMP5:%.*]] = load i16, ptr [[C_ADDR]], align 2396// CHECK1-NEXT: [[CONV6:%.*]] = sext i16 [[TMP5]] to i32397// CHECK1-NEXT: [[CMP7:%.*]] = icmp sgt i32 [[CONV6]], 0398// CHECK1-NEXT: br i1 [[CMP7]], label [[LAND_LHS_TRUE8:%.*]], label [[OMP_IF_ELSE]]399// CHECK1: land.lhs.true8:400// CHECK1-NEXT: [[TMP6:%.*]] = load float, ptr [[SC_ADDR]], align 4401// CHECK1-NEXT: [[CONV9:%.*]] = fpext float [[TMP6]] to double402// CHECK1-NEXT: [[CMP10:%.*]] = fcmp ogt double [[CONV9]], 0.000000e+00403// CHECK1-NEXT: br i1 [[CMP10]], label [[OMP_IF_THEN:%.*]], label [[OMP_IF_ELSE]]404// CHECK1: omp_if.then:405// CHECK1-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB1]], i32 3, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3foossss_l49.omp_outlined, ptr [[D_ADDR]], ptr [[GD_ADDR]], ptr [[SD_ADDR]])406// CHECK1-NEXT: br label [[OMP_IF_END:%.*]]407// CHECK1: omp_if.else:408// CHECK1-NEXT: call void @__kmpc_serialized_parallel(ptr @[[GLOB1]], i32 [[TMP0]])409// CHECK1-NEXT: store i32 [[TMP0]], ptr [[DOTTHREADID_TEMP_]], align 4410// CHECK1-NEXT: store i32 0, ptr [[DOTBOUND_ZERO_ADDR]], align 4411// CHECK1-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3foossss_l49.omp_outlined(ptr [[DOTTHREADID_TEMP_]], ptr [[DOTBOUND_ZERO_ADDR]], ptr [[D_ADDR]], ptr [[GD_ADDR]], ptr [[SD_ADDR]]) #[[ATTR2]]412// CHECK1-NEXT: call void @__kmpc_end_serialized_parallel(ptr @[[GLOB1]], i32 [[TMP0]])413// CHECK1-NEXT: br label [[OMP_IF_END]]414// CHECK1: omp_if.end:415// CHECK1-NEXT: ret void416//417//418// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3foossss_l49.omp_outlined419// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 2 dereferenceable(2) [[D:%.*]], ptr noundef nonnull align 8 dereferenceable(8) [[GD:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[SD:%.*]]) #[[ATTR1]] {420// CHECK1-NEXT: entry:421// CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8422// CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8423// CHECK1-NEXT: [[D_ADDR:%.*]] = alloca ptr, align 8424// CHECK1-NEXT: [[GD_ADDR:%.*]] = alloca ptr, align 8425// CHECK1-NEXT: [[SD_ADDR:%.*]] = alloca ptr, align 8426// CHECK1-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8427// CHECK1-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8428// CHECK1-NEXT: store ptr [[D]], ptr [[D_ADDR]], align 8429// CHECK1-NEXT: store ptr [[GD]], ptr [[GD_ADDR]], align 8430// CHECK1-NEXT: store ptr [[SD]], ptr [[SD_ADDR]], align 8431// CHECK1-NEXT: [[TMP0:%.*]] = load ptr, ptr [[D_ADDR]], align 8432// CHECK1-NEXT: [[TMP1:%.*]] = load ptr, ptr [[GD_ADDR]], align 8433// CHECK1-NEXT: [[TMP2:%.*]] = load ptr, ptr [[SD_ADDR]], align 8434// CHECK1-NEXT: [[TMP3:%.*]] = load i16, ptr [[TMP0]], align 2435// CHECK1-NEXT: [[CONV:%.*]] = sext i16 [[TMP3]] to i32436// CHECK1-NEXT: [[ADD:%.*]] = add nsw i32 [[CONV]], 1437// CHECK1-NEXT: [[CONV1:%.*]] = trunc i32 [[ADD]] to i16438// CHECK1-NEXT: store i16 [[CONV1]], ptr [[TMP0]], align 2439// CHECK1-NEXT: [[TMP4:%.*]] = load double, ptr [[TMP1]], align 8440// CHECK1-NEXT: [[ADD2:%.*]] = fadd double [[TMP4]], 1.000000e+00441// CHECK1-NEXT: store double [[ADD2]], ptr [[TMP1]], align 8442// CHECK1-NEXT: [[TMP5:%.*]] = load float, ptr [[TMP2]], align 4443// CHECK1-NEXT: [[CONV3:%.*]] = fpext float [[TMP5]] to double444// CHECK1-NEXT: [[ADD4:%.*]] = fadd double [[CONV3]], 1.000000e+00445// CHECK1-NEXT: [[CONV5:%.*]] = fptrunc double [[ADD4]] to float446// CHECK1-NEXT: store float [[CONV5]], ptr [[TMP2]], align 4447// CHECK1-NEXT: ret void448//449//450// CHECK1-LABEL: define {{[^@]+}}@_Z3barssss451// CHECK1-SAME: (i16 noundef signext [[A:%.*]], i16 noundef signext [[B:%.*]], i16 noundef signext [[C:%.*]], i16 noundef signext [[D:%.*]]) #[[ATTR0]] {452// CHECK1-NEXT: entry:453// CHECK1-NEXT: [[A_ADDR:%.*]] = alloca i16, align 2454// CHECK1-NEXT: [[B_ADDR:%.*]] = alloca i16, align 2455// CHECK1-NEXT: [[C_ADDR:%.*]] = alloca i16, align 2456// CHECK1-NEXT: [[D_ADDR:%.*]] = alloca i16, align 2457// CHECK1-NEXT: store i16 [[A]], ptr [[A_ADDR]], align 2458// CHECK1-NEXT: store i16 [[B]], ptr [[B_ADDR]], align 2459// CHECK1-NEXT: store i16 [[C]], ptr [[C_ADDR]], align 2460// CHECK1-NEXT: store i16 [[D]], ptr [[D_ADDR]], align 2461// CHECK1-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB1]], i32 4, ptr @_Z3barssss.omp_outlined, ptr [[A_ADDR]], ptr [[B_ADDR]], ptr [[C_ADDR]], ptr [[D_ADDR]])462// CHECK1-NEXT: [[TMP0:%.*]] = load i16, ptr [[A_ADDR]], align 2463// CHECK1-NEXT: [[CONV:%.*]] = sext i16 [[TMP0]] to i32464// CHECK1-NEXT: [[TMP1:%.*]] = load i16, ptr [[B_ADDR]], align 2465// CHECK1-NEXT: [[CONV1:%.*]] = sext i16 [[TMP1]] to i32466// CHECK1-NEXT: [[ADD:%.*]] = add nsw i32 [[CONV]], [[CONV1]]467// CHECK1-NEXT: [[TMP2:%.*]] = load i16, ptr [[C_ADDR]], align 2468// CHECK1-NEXT: [[CONV2:%.*]] = sext i16 [[TMP2]] to i32469// CHECK1-NEXT: [[ADD3:%.*]] = add nsw i32 [[ADD]], [[CONV2]]470// CHECK1-NEXT: [[TMP3:%.*]] = load i16, ptr [[D_ADDR]], align 2471// CHECK1-NEXT: [[CONV4:%.*]] = sext i16 [[TMP3]] to i32472// CHECK1-NEXT: [[ADD5:%.*]] = add nsw i32 [[ADD3]], [[CONV4]]473// CHECK1-NEXT: [[TMP4:%.*]] = load float, ptr @_ZZ3barssssE2Sa, align 4474// CHECK1-NEXT: [[CONV6:%.*]] = fptosi float [[TMP4]] to i32475// CHECK1-NEXT: [[ADD7:%.*]] = add nsw i32 [[ADD5]], [[CONV6]]476// CHECK1-NEXT: [[TMP5:%.*]] = load float, ptr @_ZZ3barssssE2Sb, align 4477// CHECK1-NEXT: [[CONV8:%.*]] = fptosi float [[TMP5]] to i32478// CHECK1-NEXT: [[ADD9:%.*]] = add nsw i32 [[ADD7]], [[CONV8]]479// CHECK1-NEXT: [[TMP6:%.*]] = load float, ptr @_ZZ3barssssE2Sc, align 4480// CHECK1-NEXT: [[CONV10:%.*]] = fptosi float [[TMP6]] to i32481// CHECK1-NEXT: [[ADD11:%.*]] = add nsw i32 [[ADD9]], [[CONV10]]482// CHECK1-NEXT: [[TMP7:%.*]] = load float, ptr @_ZZ3barssssE2Sd, align 4483// CHECK1-NEXT: [[CONV12:%.*]] = fptosi float [[TMP7]] to i32484// CHECK1-NEXT: [[ADD13:%.*]] = add nsw i32 [[ADD11]], [[CONV12]]485// CHECK1-NEXT: ret i32 [[ADD13]]486//487//488// CHECK1-LABEL: define {{[^@]+}}@_Z3barssss.omp_outlined489// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 2 dereferenceable(2) [[A:%.*]], ptr noundef nonnull align 2 dereferenceable(2) [[B:%.*]], ptr noundef nonnull align 2 dereferenceable(2) [[C:%.*]], ptr noundef nonnull align 2 dereferenceable(2) [[D:%.*]]) #[[ATTR1]] {490// CHECK1-NEXT: entry:491// CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8492// CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8493// CHECK1-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 8494// CHECK1-NEXT: [[B_ADDR:%.*]] = alloca ptr, align 8495// CHECK1-NEXT: [[C_ADDR:%.*]] = alloca ptr, align 8496// CHECK1-NEXT: [[D_ADDR:%.*]] = alloca ptr, align 8497// CHECK1-NEXT: [[B_CASTED:%.*]] = alloca i64, align 8498// CHECK1-NEXT: [[GB_CASTED:%.*]] = alloca i64, align 8499// CHECK1-NEXT: [[SB_CASTED:%.*]] = alloca i64, align 8500// CHECK1-NEXT: [[GC_CASTED:%.*]] = alloca i64, align 8501// CHECK1-NEXT: [[C_CASTED:%.*]] = alloca i64, align 8502// CHECK1-NEXT: [[SC_CASTED:%.*]] = alloca i64, align 8503// CHECK1-NEXT: [[D_CASTED:%.*]] = alloca i64, align 8504// CHECK1-NEXT: [[GD_CASTED:%.*]] = alloca i64, align 8505// CHECK1-NEXT: [[SD_CASTED:%.*]] = alloca i64, align 8506// CHECK1-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [9 x ptr], align 8507// CHECK1-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [9 x ptr], align 8508// CHECK1-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [9 x ptr], align 8509// CHECK1-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8510// CHECK1-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8511// CHECK1-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8512// CHECK1-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 8513// CHECK1-NEXT: store ptr [[B]], ptr [[B_ADDR]], align 8514// CHECK1-NEXT: store ptr [[C]], ptr [[C_ADDR]], align 8515// CHECK1-NEXT: store ptr [[D]], ptr [[D_ADDR]], align 8516// CHECK1-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 8517// CHECK1-NEXT: [[TMP1:%.*]] = load ptr, ptr [[B_ADDR]], align 8518// CHECK1-NEXT: [[TMP2:%.*]] = load ptr, ptr [[C_ADDR]], align 8519// CHECK1-NEXT: [[TMP3:%.*]] = load ptr, ptr [[D_ADDR]], align 8520// CHECK1-NEXT: [[TMP4:%.*]] = load i16, ptr [[TMP1]], align 2521// CHECK1-NEXT: store i16 [[TMP4]], ptr [[B_CASTED]], align 2522// CHECK1-NEXT: [[TMP5:%.*]] = load i64, ptr [[B_CASTED]], align 8523// CHECK1-NEXT: [[TMP6:%.*]] = load double, ptr @Gb, align 8524// CHECK1-NEXT: store double [[TMP6]], ptr [[GB_CASTED]], align 8525// CHECK1-NEXT: [[TMP7:%.*]] = load i64, ptr [[GB_CASTED]], align 8526// CHECK1-NEXT: [[TMP8:%.*]] = load float, ptr @_ZZ3barssssE2Sb, align 4527// CHECK1-NEXT: store float [[TMP8]], ptr [[SB_CASTED]], align 4528// CHECK1-NEXT: [[TMP9:%.*]] = load i64, ptr [[SB_CASTED]], align 8529// CHECK1-NEXT: [[TMP10:%.*]] = load double, ptr @Gc, align 8530// CHECK1-NEXT: store double [[TMP10]], ptr [[GC_CASTED]], align 8531// CHECK1-NEXT: [[TMP11:%.*]] = load i64, ptr [[GC_CASTED]], align 8532// CHECK1-NEXT: [[TMP12:%.*]] = load i16, ptr [[TMP2]], align 2533// CHECK1-NEXT: store i16 [[TMP12]], ptr [[C_CASTED]], align 2534// CHECK1-NEXT: [[TMP13:%.*]] = load i64, ptr [[C_CASTED]], align 8535// CHECK1-NEXT: [[TMP14:%.*]] = load float, ptr @_ZZ3barssssE2Sc, align 4536// CHECK1-NEXT: store float [[TMP14]], ptr [[SC_CASTED]], align 4537// CHECK1-NEXT: [[TMP15:%.*]] = load i64, ptr [[SC_CASTED]], align 8538// CHECK1-NEXT: [[TMP16:%.*]] = load i16, ptr [[TMP3]], align 2539// CHECK1-NEXT: store i16 [[TMP16]], ptr [[D_CASTED]], align 2540// CHECK1-NEXT: [[TMP17:%.*]] = load i64, ptr [[D_CASTED]], align 8541// CHECK1-NEXT: [[TMP18:%.*]] = load double, ptr @Gd, align 8542// CHECK1-NEXT: store double [[TMP18]], ptr [[GD_CASTED]], align 8543// CHECK1-NEXT: [[TMP19:%.*]] = load i64, ptr [[GD_CASTED]], align 8544// CHECK1-NEXT: [[TMP20:%.*]] = load float, ptr @_ZZ3barssssE2Sd, align 4545// CHECK1-NEXT: store float [[TMP20]], ptr [[SD_CASTED]], align 4546// CHECK1-NEXT: [[TMP21:%.*]] = load i64, ptr [[SD_CASTED]], align 8547// CHECK1-NEXT: [[TMP22:%.*]] = load double, ptr @Ga, align 8548// CHECK1-NEXT: [[CMP:%.*]] = fcmp ogt double [[TMP22]], 0.000000e+00549// CHECK1-NEXT: br i1 [[CMP]], label [[LAND_LHS_TRUE:%.*]], label [[OMP_IF_ELSE:%.*]]550// CHECK1: land.lhs.true:551// CHECK1-NEXT: [[TMP23:%.*]] = load i16, ptr [[TMP0]], align 2552// CHECK1-NEXT: [[CONV:%.*]] = sext i16 [[TMP23]] to i32553// CHECK1-NEXT: [[CMP1:%.*]] = icmp sgt i32 [[CONV]], 0554// CHECK1-NEXT: br i1 [[CMP1]], label [[LAND_LHS_TRUE2:%.*]], label [[OMP_IF_ELSE]]555// CHECK1: land.lhs.true2:556// CHECK1-NEXT: [[TMP24:%.*]] = load float, ptr @_ZZ3barssssE2Sa, align 4557// CHECK1-NEXT: [[CONV3:%.*]] = fpext float [[TMP24]] to double558// CHECK1-NEXT: [[CMP4:%.*]] = fcmp ogt double [[CONV3]], 0.000000e+00559// CHECK1-NEXT: br i1 [[CMP4]], label [[OMP_IF_THEN:%.*]], label [[OMP_IF_ELSE]]560// CHECK1: omp_if.then:561// CHECK1-NEXT: [[TMP25:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0562// CHECK1-NEXT: store i64 [[TMP5]], ptr [[TMP25]], align 8563// CHECK1-NEXT: [[TMP26:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0564// CHECK1-NEXT: store i64 [[TMP5]], ptr [[TMP26]], align 8565// CHECK1-NEXT: [[TMP27:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 0566// CHECK1-NEXT: store ptr null, ptr [[TMP27]], align 8567// CHECK1-NEXT: [[TMP28:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 1568// CHECK1-NEXT: store i64 [[TMP7]], ptr [[TMP28]], align 8569// CHECK1-NEXT: [[TMP29:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 1570// CHECK1-NEXT: store i64 [[TMP7]], ptr [[TMP29]], align 8571// CHECK1-NEXT: [[TMP30:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 1572// CHECK1-NEXT: store ptr null, ptr [[TMP30]], align 8573// CHECK1-NEXT: [[TMP31:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 2574// CHECK1-NEXT: store i64 [[TMP9]], ptr [[TMP31]], align 8575// CHECK1-NEXT: [[TMP32:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 2576// CHECK1-NEXT: store i64 [[TMP9]], ptr [[TMP32]], align 8577// CHECK1-NEXT: [[TMP33:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 2578// CHECK1-NEXT: store ptr null, ptr [[TMP33]], align 8579// CHECK1-NEXT: [[TMP34:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 3580// CHECK1-NEXT: store i64 [[TMP11]], ptr [[TMP34]], align 8581// CHECK1-NEXT: [[TMP35:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 3582// CHECK1-NEXT: store i64 [[TMP11]], ptr [[TMP35]], align 8583// CHECK1-NEXT: [[TMP36:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 3584// CHECK1-NEXT: store ptr null, ptr [[TMP36]], align 8585// CHECK1-NEXT: [[TMP37:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 4586// CHECK1-NEXT: store i64 [[TMP13]], ptr [[TMP37]], align 8587// CHECK1-NEXT: [[TMP38:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 4588// CHECK1-NEXT: store i64 [[TMP13]], ptr [[TMP38]], align 8589// CHECK1-NEXT: [[TMP39:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 4590// CHECK1-NEXT: store ptr null, ptr [[TMP39]], align 8591// CHECK1-NEXT: [[TMP40:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 5592// CHECK1-NEXT: store i64 [[TMP15]], ptr [[TMP40]], align 8593// CHECK1-NEXT: [[TMP41:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 5594// CHECK1-NEXT: store i64 [[TMP15]], ptr [[TMP41]], align 8595// CHECK1-NEXT: [[TMP42:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 5596// CHECK1-NEXT: store ptr null, ptr [[TMP42]], align 8597// CHECK1-NEXT: [[TMP43:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 6598// CHECK1-NEXT: store i64 [[TMP17]], ptr [[TMP43]], align 8599// CHECK1-NEXT: [[TMP44:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 6600// CHECK1-NEXT: store i64 [[TMP17]], ptr [[TMP44]], align 8601// CHECK1-NEXT: [[TMP45:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 6602// CHECK1-NEXT: store ptr null, ptr [[TMP45]], align 8603// CHECK1-NEXT: [[TMP46:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 7604// CHECK1-NEXT: store i64 [[TMP19]], ptr [[TMP46]], align 8605// CHECK1-NEXT: [[TMP47:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 7606// CHECK1-NEXT: store i64 [[TMP19]], ptr [[TMP47]], align 8607// CHECK1-NEXT: [[TMP48:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 7608// CHECK1-NEXT: store ptr null, ptr [[TMP48]], align 8609// CHECK1-NEXT: [[TMP49:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 8610// CHECK1-NEXT: store i64 [[TMP21]], ptr [[TMP49]], align 8611// CHECK1-NEXT: [[TMP50:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 8612// CHECK1-NEXT: store i64 [[TMP21]], ptr [[TMP50]], align 8613// CHECK1-NEXT: [[TMP51:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 8614// CHECK1-NEXT: store ptr null, ptr [[TMP51]], align 8615// CHECK1-NEXT: [[TMP52:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0616// CHECK1-NEXT: [[TMP53:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0617// CHECK1-NEXT: [[TMP54:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0618// CHECK1-NEXT: store i32 3, ptr [[TMP54]], align 4619// CHECK1-NEXT: [[TMP55:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1620// CHECK1-NEXT: store i32 9, ptr [[TMP55]], align 4621// CHECK1-NEXT: [[TMP56:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2622// CHECK1-NEXT: store ptr [[TMP52]], ptr [[TMP56]], align 8623// CHECK1-NEXT: [[TMP57:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3624// CHECK1-NEXT: store ptr [[TMP53]], ptr [[TMP57]], align 8625// CHECK1-NEXT: [[TMP58:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4626// CHECK1-NEXT: store ptr @.offload_sizes.1, ptr [[TMP58]], align 8627// CHECK1-NEXT: [[TMP59:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5628// CHECK1-NEXT: store ptr @.offload_maptypes.2, ptr [[TMP59]], align 8629// CHECK1-NEXT: [[TMP60:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6630// CHECK1-NEXT: store ptr null, ptr [[TMP60]], align 8631// CHECK1-NEXT: [[TMP61:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7632// CHECK1-NEXT: store ptr null, ptr [[TMP61]], align 8633// CHECK1-NEXT: [[TMP62:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8634// CHECK1-NEXT: store i64 0, ptr [[TMP62]], align 8635// CHECK1-NEXT: [[TMP63:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9636// CHECK1-NEXT: store i64 0, ptr [[TMP63]], align 8637// CHECK1-NEXT: [[TMP64:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10638// CHECK1-NEXT: store [3 x i32] [i32 -1, i32 0, i32 0], ptr [[TMP64]], align 4639// CHECK1-NEXT: [[TMP65:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11640// CHECK1-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP65]], align 4641// CHECK1-NEXT: [[TMP66:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12642// CHECK1-NEXT: store i32 0, ptr [[TMP66]], align 4643// CHECK1-NEXT: [[TMP67:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 -1, i32 -1, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3barssss_l94.region_id, ptr [[KERNEL_ARGS]])644// CHECK1-NEXT: [[TMP68:%.*]] = icmp ne i32 [[TMP67]], 0645// CHECK1-NEXT: br i1 [[TMP68]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]646// CHECK1: omp_offload.failed:647// CHECK1-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3barssss_l94(i64 [[TMP5]], i64 [[TMP7]], i64 [[TMP9]], i64 [[TMP11]], i64 [[TMP13]], i64 [[TMP15]], i64 [[TMP17]], i64 [[TMP19]], i64 [[TMP21]]) #[[ATTR2]]648// CHECK1-NEXT: br label [[OMP_OFFLOAD_CONT]]649// CHECK1: omp_offload.cont:650// CHECK1-NEXT: br label [[OMP_IF_END:%.*]]651// CHECK1: omp_if.else:652// CHECK1-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3barssss_l94(i64 [[TMP5]], i64 [[TMP7]], i64 [[TMP9]], i64 [[TMP11]], i64 [[TMP13]], i64 [[TMP15]], i64 [[TMP17]], i64 [[TMP19]], i64 [[TMP21]]) #[[ATTR2]]653// CHECK1-NEXT: br label [[OMP_IF_END]]654// CHECK1: omp_if.end:655// CHECK1-NEXT: ret void656//657//658// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3barssss_l94659// CHECK1-SAME: (i64 noundef [[B:%.*]], i64 noundef [[GB:%.*]], i64 noundef [[SB:%.*]], i64 noundef [[GC:%.*]], i64 noundef [[C:%.*]], i64 noundef [[SC:%.*]], i64 noundef [[D:%.*]], i64 noundef [[GD:%.*]], i64 noundef [[SD:%.*]]) #[[ATTR1]] {660// CHECK1-NEXT: entry:661// CHECK1-NEXT: [[B_ADDR:%.*]] = alloca i64, align 8662// CHECK1-NEXT: [[GB_ADDR:%.*]] = alloca i64, align 8663// CHECK1-NEXT: [[SB_ADDR:%.*]] = alloca i64, align 8664// CHECK1-NEXT: [[GC_ADDR:%.*]] = alloca i64, align 8665// CHECK1-NEXT: [[C_ADDR:%.*]] = alloca i64, align 8666// CHECK1-NEXT: [[SC_ADDR:%.*]] = alloca i64, align 8667// CHECK1-NEXT: [[D_ADDR:%.*]] = alloca i64, align 8668// CHECK1-NEXT: [[GD_ADDR:%.*]] = alloca i64, align 8669// CHECK1-NEXT: [[SD_ADDR:%.*]] = alloca i64, align 8670// CHECK1-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 4671// CHECK1-NEXT: [[DOTBOUND_ZERO_ADDR:%.*]] = alloca i32, align 4672// CHECK1-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])673// CHECK1-NEXT: store i64 [[B]], ptr [[B_ADDR]], align 8674// CHECK1-NEXT: store i64 [[GB]], ptr [[GB_ADDR]], align 8675// CHECK1-NEXT: store i64 [[SB]], ptr [[SB_ADDR]], align 8676// CHECK1-NEXT: store i64 [[GC]], ptr [[GC_ADDR]], align 8677// CHECK1-NEXT: store i64 [[C]], ptr [[C_ADDR]], align 8678// CHECK1-NEXT: store i64 [[SC]], ptr [[SC_ADDR]], align 8679// CHECK1-NEXT: store i64 [[D]], ptr [[D_ADDR]], align 8680// CHECK1-NEXT: store i64 [[GD]], ptr [[GD_ADDR]], align 8681// CHECK1-NEXT: store i64 [[SD]], ptr [[SD_ADDR]], align 8682// CHECK1-NEXT: [[TMP1:%.*]] = load i16, ptr [[B_ADDR]], align 2683// CHECK1-NEXT: [[CONV:%.*]] = sext i16 [[TMP1]] to i32684// CHECK1-NEXT: [[ADD:%.*]] = add nsw i32 [[CONV]], 1685// CHECK1-NEXT: [[CONV1:%.*]] = trunc i32 [[ADD]] to i16686// CHECK1-NEXT: store i16 [[CONV1]], ptr [[B_ADDR]], align 2687// CHECK1-NEXT: [[TMP2:%.*]] = load double, ptr [[GB_ADDR]], align 8688// CHECK1-NEXT: [[ADD2:%.*]] = fadd double [[TMP2]], 1.000000e+00689// CHECK1-NEXT: store double [[ADD2]], ptr [[GB_ADDR]], align 8690// CHECK1-NEXT: [[TMP3:%.*]] = load float, ptr [[SB_ADDR]], align 4691// CHECK1-NEXT: [[CONV3:%.*]] = fpext float [[TMP3]] to double692// CHECK1-NEXT: [[ADD4:%.*]] = fadd double [[CONV3]], 1.000000e+00693// CHECK1-NEXT: [[CONV5:%.*]] = fptrunc double [[ADD4]] to float694// CHECK1-NEXT: store float [[CONV5]], ptr [[SB_ADDR]], align 4695// CHECK1-NEXT: [[TMP4:%.*]] = load double, ptr [[GC_ADDR]], align 8696// CHECK1-NEXT: [[CMP:%.*]] = fcmp ogt double [[TMP4]], 0.000000e+00697// CHECK1-NEXT: br i1 [[CMP]], label [[LAND_LHS_TRUE:%.*]], label [[OMP_IF_ELSE:%.*]]698// CHECK1: land.lhs.true:699// CHECK1-NEXT: [[TMP5:%.*]] = load i16, ptr [[C_ADDR]], align 2700// CHECK1-NEXT: [[CONV6:%.*]] = sext i16 [[TMP5]] to i32701// CHECK1-NEXT: [[CMP7:%.*]] = icmp sgt i32 [[CONV6]], 0702// CHECK1-NEXT: br i1 [[CMP7]], label [[LAND_LHS_TRUE8:%.*]], label [[OMP_IF_ELSE]]703// CHECK1: land.lhs.true8:704// CHECK1-NEXT: [[TMP6:%.*]] = load float, ptr [[SC_ADDR]], align 4705// CHECK1-NEXT: [[CONV9:%.*]] = fpext float [[TMP6]] to double706// CHECK1-NEXT: [[CMP10:%.*]] = fcmp ogt double [[CONV9]], 0.000000e+00707// CHECK1-NEXT: br i1 [[CMP10]], label [[OMP_IF_THEN:%.*]], label [[OMP_IF_ELSE]]708// CHECK1: omp_if.then:709// CHECK1-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB1]], i32 3, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3barssss_l94.omp_outlined, ptr [[D_ADDR]], ptr [[GD_ADDR]], ptr [[SD_ADDR]])710// CHECK1-NEXT: br label [[OMP_IF_END:%.*]]711// CHECK1: omp_if.else:712// CHECK1-NEXT: call void @__kmpc_serialized_parallel(ptr @[[GLOB1]], i32 [[TMP0]])713// CHECK1-NEXT: store i32 [[TMP0]], ptr [[DOTTHREADID_TEMP_]], align 4714// CHECK1-NEXT: store i32 0, ptr [[DOTBOUND_ZERO_ADDR]], align 4715// CHECK1-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3barssss_l94.omp_outlined(ptr [[DOTTHREADID_TEMP_]], ptr [[DOTBOUND_ZERO_ADDR]], ptr [[D_ADDR]], ptr [[GD_ADDR]], ptr [[SD_ADDR]]) #[[ATTR2]]716// CHECK1-NEXT: call void @__kmpc_end_serialized_parallel(ptr @[[GLOB1]], i32 [[TMP0]])717// CHECK1-NEXT: br label [[OMP_IF_END]]718// CHECK1: omp_if.end:719// CHECK1-NEXT: ret void720//721//722// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3barssss_l94.omp_outlined723// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 2 dereferenceable(2) [[D:%.*]], ptr noundef nonnull align 8 dereferenceable(8) [[GD:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[SD:%.*]]) #[[ATTR1]] {724// CHECK1-NEXT: entry:725// CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8726// CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8727// CHECK1-NEXT: [[D_ADDR:%.*]] = alloca ptr, align 8728// CHECK1-NEXT: [[GD_ADDR:%.*]] = alloca ptr, align 8729// CHECK1-NEXT: [[SD_ADDR:%.*]] = alloca ptr, align 8730// CHECK1-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8731// CHECK1-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8732// CHECK1-NEXT: store ptr [[D]], ptr [[D_ADDR]], align 8733// CHECK1-NEXT: store ptr [[GD]], ptr [[GD_ADDR]], align 8734// CHECK1-NEXT: store ptr [[SD]], ptr [[SD_ADDR]], align 8735// CHECK1-NEXT: [[TMP0:%.*]] = load ptr, ptr [[D_ADDR]], align 8736// CHECK1-NEXT: [[TMP1:%.*]] = load ptr, ptr [[GD_ADDR]], align 8737// CHECK1-NEXT: [[TMP2:%.*]] = load ptr, ptr [[SD_ADDR]], align 8738// CHECK1-NEXT: [[TMP3:%.*]] = load i16, ptr [[TMP0]], align 2739// CHECK1-NEXT: [[CONV:%.*]] = sext i16 [[TMP3]] to i32740// CHECK1-NEXT: [[ADD:%.*]] = add nsw i32 [[CONV]], 1741// CHECK1-NEXT: [[CONV1:%.*]] = trunc i32 [[ADD]] to i16742// CHECK1-NEXT: store i16 [[CONV1]], ptr [[TMP0]], align 2743// CHECK1-NEXT: [[TMP4:%.*]] = load double, ptr [[TMP1]], align 8744// CHECK1-NEXT: [[ADD2:%.*]] = fadd double [[TMP4]], 1.000000e+00745// CHECK1-NEXT: store double [[ADD2]], ptr [[TMP1]], align 8746// CHECK1-NEXT: [[TMP5:%.*]] = load float, ptr [[TMP2]], align 4747// CHECK1-NEXT: [[CONV3:%.*]] = fpext float [[TMP5]] to double748// CHECK1-NEXT: [[ADD4:%.*]] = fadd double [[CONV3]], 1.000000e+00749// CHECK1-NEXT: [[CONV5:%.*]] = fptrunc double [[ADD4]] to float750// CHECK1-NEXT: store float [[CONV5]], ptr [[TMP2]], align 4751// CHECK1-NEXT: ret void752//753//754// CHECK1-LABEL: define {{[^@]+}}@_Z5tbar2ssss755// CHECK1-SAME: (i16 noundef signext [[A:%.*]], i16 noundef signext [[B:%.*]], i16 noundef signext [[C:%.*]], i16 noundef signext [[D:%.*]]) #[[ATTR0]] {756// CHECK1-NEXT: entry:757// CHECK1-NEXT: [[A_ADDR:%.*]] = alloca i16, align 2758// CHECK1-NEXT: [[B_ADDR:%.*]] = alloca i16, align 2759// CHECK1-NEXT: [[C_ADDR:%.*]] = alloca i16, align 2760// CHECK1-NEXT: [[D_ADDR:%.*]] = alloca i16, align 2761// CHECK1-NEXT: store i16 [[A]], ptr [[A_ADDR]], align 2762// CHECK1-NEXT: store i16 [[B]], ptr [[B_ADDR]], align 2763// CHECK1-NEXT: store i16 [[C]], ptr [[C_ADDR]], align 2764// CHECK1-NEXT: store i16 [[D]], ptr [[D_ADDR]], align 2765// CHECK1-NEXT: [[TMP0:%.*]] = load i16, ptr [[A_ADDR]], align 2766// CHECK1-NEXT: [[TMP1:%.*]] = load i16, ptr [[B_ADDR]], align 2767// CHECK1-NEXT: [[TMP2:%.*]] = load i16, ptr [[C_ADDR]], align 2768// CHECK1-NEXT: [[TMP3:%.*]] = load i16, ptr [[D_ADDR]], align 2769// CHECK1-NEXT: [[CALL:%.*]] = call noundef signext i32 @_Z4tbarIsEiT_S0_S0_S0_(i16 noundef signext [[TMP0]], i16 noundef signext [[TMP1]], i16 noundef signext [[TMP2]], i16 noundef signext [[TMP3]])770// CHECK1-NEXT: ret i32 [[CALL]]771//772//773// CHECK1-LABEL: define {{[^@]+}}@_Z4tbarIsEiT_S0_S0_S0_774// CHECK1-SAME: (i16 noundef signext [[A:%.*]], i16 noundef signext [[B:%.*]], i16 noundef signext [[C:%.*]], i16 noundef signext [[D:%.*]]) #[[ATTR0]] comdat {775// CHECK1-NEXT: entry:776// CHECK1-NEXT: [[A_ADDR:%.*]] = alloca i16, align 2777// CHECK1-NEXT: [[B_ADDR:%.*]] = alloca i16, align 2778// CHECK1-NEXT: [[C_ADDR:%.*]] = alloca i16, align 2779// CHECK1-NEXT: [[D_ADDR:%.*]] = alloca i16, align 2780// CHECK1-NEXT: store i16 [[A]], ptr [[A_ADDR]], align 2781// CHECK1-NEXT: store i16 [[B]], ptr [[B_ADDR]], align 2782// CHECK1-NEXT: store i16 [[C]], ptr [[C_ADDR]], align 2783// CHECK1-NEXT: store i16 [[D]], ptr [[D_ADDR]], align 2784// CHECK1-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB1]], i32 4, ptr @_Z4tbarIsEiT_S0_S0_S0_.omp_outlined, ptr [[A_ADDR]], ptr [[B_ADDR]], ptr [[C_ADDR]], ptr [[D_ADDR]])785// CHECK1-NEXT: [[TMP0:%.*]] = load i16, ptr [[A_ADDR]], align 2786// CHECK1-NEXT: [[CONV:%.*]] = sext i16 [[TMP0]] to i32787// CHECK1-NEXT: [[TMP1:%.*]] = load i16, ptr [[B_ADDR]], align 2788// CHECK1-NEXT: [[CONV1:%.*]] = sext i16 [[TMP1]] to i32789// CHECK1-NEXT: [[ADD:%.*]] = add nsw i32 [[CONV]], [[CONV1]]790// CHECK1-NEXT: [[TMP2:%.*]] = load i16, ptr [[C_ADDR]], align 2791// CHECK1-NEXT: [[CONV2:%.*]] = sext i16 [[TMP2]] to i32792// CHECK1-NEXT: [[ADD3:%.*]] = add nsw i32 [[ADD]], [[CONV2]]793// CHECK1-NEXT: [[TMP3:%.*]] = load i16, ptr [[D_ADDR]], align 2794// CHECK1-NEXT: [[CONV4:%.*]] = sext i16 [[TMP3]] to i32795// CHECK1-NEXT: [[ADD5:%.*]] = add nsw i32 [[ADD3]], [[CONV4]]796// CHECK1-NEXT: [[TMP4:%.*]] = load float, ptr @_ZZ4tbarIsEiT_S0_S0_S0_E2Sa, align 4797// CHECK1-NEXT: [[CONV6:%.*]] = fptosi float [[TMP4]] to i32798// CHECK1-NEXT: [[ADD7:%.*]] = add nsw i32 [[ADD5]], [[CONV6]]799// CHECK1-NEXT: [[TMP5:%.*]] = load float, ptr @_ZZ4tbarIsEiT_S0_S0_S0_E2Sb, align 4800// CHECK1-NEXT: [[CONV8:%.*]] = fptosi float [[TMP5]] to i32801// CHECK1-NEXT: [[ADD9:%.*]] = add nsw i32 [[ADD7]], [[CONV8]]802// CHECK1-NEXT: [[TMP6:%.*]] = load float, ptr @_ZZ4tbarIsEiT_S0_S0_S0_E2Sc, align 4803// CHECK1-NEXT: [[CONV10:%.*]] = fptosi float [[TMP6]] to i32804// CHECK1-NEXT: [[ADD11:%.*]] = add nsw i32 [[ADD9]], [[CONV10]]805// CHECK1-NEXT: [[TMP7:%.*]] = load float, ptr @_ZZ4tbarIsEiT_S0_S0_S0_E2Sd, align 4806// CHECK1-NEXT: [[CONV12:%.*]] = fptosi float [[TMP7]] to i32807// CHECK1-NEXT: [[ADD13:%.*]] = add nsw i32 [[ADD11]], [[CONV12]]808// CHECK1-NEXT: ret i32 [[ADD13]]809//810//811// CHECK1-LABEL: define {{[^@]+}}@_Z4tbarIsEiT_S0_S0_S0_.omp_outlined812// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 2 dereferenceable(2) [[A:%.*]], ptr noundef nonnull align 2 dereferenceable(2) [[B:%.*]], ptr noundef nonnull align 2 dereferenceable(2) [[C:%.*]], ptr noundef nonnull align 2 dereferenceable(2) [[D:%.*]]) #[[ATTR1]] {813// CHECK1-NEXT: entry:814// CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8815// CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8816// CHECK1-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 8817// CHECK1-NEXT: [[B_ADDR:%.*]] = alloca ptr, align 8818// CHECK1-NEXT: [[C_ADDR:%.*]] = alloca ptr, align 8819// CHECK1-NEXT: [[D_ADDR:%.*]] = alloca ptr, align 8820// CHECK1-NEXT: [[B_CASTED:%.*]] = alloca i64, align 8821// CHECK1-NEXT: [[GB_CASTED:%.*]] = alloca i64, align 8822// CHECK1-NEXT: [[SB_CASTED:%.*]] = alloca i64, align 8823// CHECK1-NEXT: [[GC_CASTED:%.*]] = alloca i64, align 8824// CHECK1-NEXT: [[C_CASTED:%.*]] = alloca i64, align 8825// CHECK1-NEXT: [[SC_CASTED:%.*]] = alloca i64, align 8826// CHECK1-NEXT: [[D_CASTED:%.*]] = alloca i64, align 8827// CHECK1-NEXT: [[GD_CASTED:%.*]] = alloca i64, align 8828// CHECK1-NEXT: [[SD_CASTED:%.*]] = alloca i64, align 8829// CHECK1-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [9 x ptr], align 8830// CHECK1-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [9 x ptr], align 8831// CHECK1-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [9 x ptr], align 8832// CHECK1-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8833// CHECK1-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8834// CHECK1-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8835// CHECK1-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 8836// CHECK1-NEXT: store ptr [[B]], ptr [[B_ADDR]], align 8837// CHECK1-NEXT: store ptr [[C]], ptr [[C_ADDR]], align 8838// CHECK1-NEXT: store ptr [[D]], ptr [[D_ADDR]], align 8839// CHECK1-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 8840// CHECK1-NEXT: [[TMP1:%.*]] = load ptr, ptr [[B_ADDR]], align 8841// CHECK1-NEXT: [[TMP2:%.*]] = load ptr, ptr [[C_ADDR]], align 8842// CHECK1-NEXT: [[TMP3:%.*]] = load ptr, ptr [[D_ADDR]], align 8843// CHECK1-NEXT: [[TMP4:%.*]] = load i16, ptr [[TMP1]], align 2844// CHECK1-NEXT: store i16 [[TMP4]], ptr [[B_CASTED]], align 2845// CHECK1-NEXT: [[TMP5:%.*]] = load i64, ptr [[B_CASTED]], align 8846// CHECK1-NEXT: [[TMP6:%.*]] = load double, ptr @Gb, align 8847// CHECK1-NEXT: store double [[TMP6]], ptr [[GB_CASTED]], align 8848// CHECK1-NEXT: [[TMP7:%.*]] = load i64, ptr [[GB_CASTED]], align 8849// CHECK1-NEXT: [[TMP8:%.*]] = load float, ptr @_ZZ4tbarIsEiT_S0_S0_S0_E2Sb, align 4850// CHECK1-NEXT: store float [[TMP8]], ptr [[SB_CASTED]], align 4851// CHECK1-NEXT: [[TMP9:%.*]] = load i64, ptr [[SB_CASTED]], align 8852// CHECK1-NEXT: [[TMP10:%.*]] = load double, ptr @Gc, align 8853// CHECK1-NEXT: store double [[TMP10]], ptr [[GC_CASTED]], align 8854// CHECK1-NEXT: [[TMP11:%.*]] = load i64, ptr [[GC_CASTED]], align 8855// CHECK1-NEXT: [[TMP12:%.*]] = load i16, ptr [[TMP2]], align 2856// CHECK1-NEXT: store i16 [[TMP12]], ptr [[C_CASTED]], align 2857// CHECK1-NEXT: [[TMP13:%.*]] = load i64, ptr [[C_CASTED]], align 8858// CHECK1-NEXT: [[TMP14:%.*]] = load float, ptr @_ZZ4tbarIsEiT_S0_S0_S0_E2Sc, align 4859// CHECK1-NEXT: store float [[TMP14]], ptr [[SC_CASTED]], align 4860// CHECK1-NEXT: [[TMP15:%.*]] = load i64, ptr [[SC_CASTED]], align 8861// CHECK1-NEXT: [[TMP16:%.*]] = load i16, ptr [[TMP3]], align 2862// CHECK1-NEXT: store i16 [[TMP16]], ptr [[D_CASTED]], align 2863// CHECK1-NEXT: [[TMP17:%.*]] = load i64, ptr [[D_CASTED]], align 8864// CHECK1-NEXT: [[TMP18:%.*]] = load double, ptr @Gd, align 8865// CHECK1-NEXT: store double [[TMP18]], ptr [[GD_CASTED]], align 8866// CHECK1-NEXT: [[TMP19:%.*]] = load i64, ptr [[GD_CASTED]], align 8867// CHECK1-NEXT: [[TMP20:%.*]] = load float, ptr @_ZZ4tbarIsEiT_S0_S0_S0_E2Sd, align 4868// CHECK1-NEXT: store float [[TMP20]], ptr [[SD_CASTED]], align 4869// CHECK1-NEXT: [[TMP21:%.*]] = load i64, ptr [[SD_CASTED]], align 8870// CHECK1-NEXT: [[TMP22:%.*]] = load double, ptr @Ga, align 8871// CHECK1-NEXT: [[CMP:%.*]] = fcmp ogt double [[TMP22]], 0.000000e+00872// CHECK1-NEXT: br i1 [[CMP]], label [[LAND_LHS_TRUE:%.*]], label [[OMP_IF_ELSE:%.*]]873// CHECK1: land.lhs.true:874// CHECK1-NEXT: [[TMP23:%.*]] = load i16, ptr [[TMP0]], align 2875// CHECK1-NEXT: [[CONV:%.*]] = sext i16 [[TMP23]] to i32876// CHECK1-NEXT: [[CMP1:%.*]] = icmp sgt i32 [[CONV]], 0877// CHECK1-NEXT: br i1 [[CMP1]], label [[LAND_LHS_TRUE2:%.*]], label [[OMP_IF_ELSE]]878// CHECK1: land.lhs.true2:879// CHECK1-NEXT: [[TMP24:%.*]] = load float, ptr @_ZZ4tbarIsEiT_S0_S0_S0_E2Sa, align 4880// CHECK1-NEXT: [[CONV3:%.*]] = fpext float [[TMP24]] to double881// CHECK1-NEXT: [[CMP4:%.*]] = fcmp ogt double [[CONV3]], 0.000000e+00882// CHECK1-NEXT: br i1 [[CMP4]], label [[OMP_IF_THEN:%.*]], label [[OMP_IF_ELSE]]883// CHECK1: omp_if.then:884// CHECK1-NEXT: [[TMP25:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0885// CHECK1-NEXT: store i64 [[TMP5]], ptr [[TMP25]], align 8886// CHECK1-NEXT: [[TMP26:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0887// CHECK1-NEXT: store i64 [[TMP5]], ptr [[TMP26]], align 8888// CHECK1-NEXT: [[TMP27:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 0889// CHECK1-NEXT: store ptr null, ptr [[TMP27]], align 8890// CHECK1-NEXT: [[TMP28:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 1891// CHECK1-NEXT: store i64 [[TMP7]], ptr [[TMP28]], align 8892// CHECK1-NEXT: [[TMP29:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 1893// CHECK1-NEXT: store i64 [[TMP7]], ptr [[TMP29]], align 8894// CHECK1-NEXT: [[TMP30:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 1895// CHECK1-NEXT: store ptr null, ptr [[TMP30]], align 8896// CHECK1-NEXT: [[TMP31:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 2897// CHECK1-NEXT: store i64 [[TMP9]], ptr [[TMP31]], align 8898// CHECK1-NEXT: [[TMP32:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 2899// CHECK1-NEXT: store i64 [[TMP9]], ptr [[TMP32]], align 8900// CHECK1-NEXT: [[TMP33:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 2901// CHECK1-NEXT: store ptr null, ptr [[TMP33]], align 8902// CHECK1-NEXT: [[TMP34:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 3903// CHECK1-NEXT: store i64 [[TMP11]], ptr [[TMP34]], align 8904// CHECK1-NEXT: [[TMP35:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 3905// CHECK1-NEXT: store i64 [[TMP11]], ptr [[TMP35]], align 8906// CHECK1-NEXT: [[TMP36:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 3907// CHECK1-NEXT: store ptr null, ptr [[TMP36]], align 8908// CHECK1-NEXT: [[TMP37:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 4909// CHECK1-NEXT: store i64 [[TMP13]], ptr [[TMP37]], align 8910// CHECK1-NEXT: [[TMP38:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 4911// CHECK1-NEXT: store i64 [[TMP13]], ptr [[TMP38]], align 8912// CHECK1-NEXT: [[TMP39:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 4913// CHECK1-NEXT: store ptr null, ptr [[TMP39]], align 8914// CHECK1-NEXT: [[TMP40:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 5915// CHECK1-NEXT: store i64 [[TMP15]], ptr [[TMP40]], align 8916// CHECK1-NEXT: [[TMP41:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 5917// CHECK1-NEXT: store i64 [[TMP15]], ptr [[TMP41]], align 8918// CHECK1-NEXT: [[TMP42:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 5919// CHECK1-NEXT: store ptr null, ptr [[TMP42]], align 8920// CHECK1-NEXT: [[TMP43:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 6921// CHECK1-NEXT: store i64 [[TMP17]], ptr [[TMP43]], align 8922// CHECK1-NEXT: [[TMP44:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 6923// CHECK1-NEXT: store i64 [[TMP17]], ptr [[TMP44]], align 8924// CHECK1-NEXT: [[TMP45:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 6925// CHECK1-NEXT: store ptr null, ptr [[TMP45]], align 8926// CHECK1-NEXT: [[TMP46:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 7927// CHECK1-NEXT: store i64 [[TMP19]], ptr [[TMP46]], align 8928// CHECK1-NEXT: [[TMP47:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 7929// CHECK1-NEXT: store i64 [[TMP19]], ptr [[TMP47]], align 8930// CHECK1-NEXT: [[TMP48:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 7931// CHECK1-NEXT: store ptr null, ptr [[TMP48]], align 8932// CHECK1-NEXT: [[TMP49:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 8933// CHECK1-NEXT: store i64 [[TMP21]], ptr [[TMP49]], align 8934// CHECK1-NEXT: [[TMP50:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 8935// CHECK1-NEXT: store i64 [[TMP21]], ptr [[TMP50]], align 8936// CHECK1-NEXT: [[TMP51:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 8937// CHECK1-NEXT: store ptr null, ptr [[TMP51]], align 8938// CHECK1-NEXT: [[TMP52:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0939// CHECK1-NEXT: [[TMP53:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0940// CHECK1-NEXT: [[TMP54:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0941// CHECK1-NEXT: store i32 3, ptr [[TMP54]], align 4942// CHECK1-NEXT: [[TMP55:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1943// CHECK1-NEXT: store i32 9, ptr [[TMP55]], align 4944// CHECK1-NEXT: [[TMP56:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2945// CHECK1-NEXT: store ptr [[TMP52]], ptr [[TMP56]], align 8946// CHECK1-NEXT: [[TMP57:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3947// CHECK1-NEXT: store ptr [[TMP53]], ptr [[TMP57]], align 8948// CHECK1-NEXT: [[TMP58:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4949// CHECK1-NEXT: store ptr @.offload_sizes.3, ptr [[TMP58]], align 8950// CHECK1-NEXT: [[TMP59:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5951// CHECK1-NEXT: store ptr @.offload_maptypes.4, ptr [[TMP59]], align 8952// CHECK1-NEXT: [[TMP60:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6953// CHECK1-NEXT: store ptr null, ptr [[TMP60]], align 8954// CHECK1-NEXT: [[TMP61:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7955// CHECK1-NEXT: store ptr null, ptr [[TMP61]], align 8956// CHECK1-NEXT: [[TMP62:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8957// CHECK1-NEXT: store i64 0, ptr [[TMP62]], align 8958// CHECK1-NEXT: [[TMP63:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9959// CHECK1-NEXT: store i64 0, ptr [[TMP63]], align 8960// CHECK1-NEXT: [[TMP64:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10961// CHECK1-NEXT: store [3 x i32] [i32 -1, i32 0, i32 0], ptr [[TMP64]], align 4962// CHECK1-NEXT: [[TMP65:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11963// CHECK1-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP65]], align 4964// CHECK1-NEXT: [[TMP66:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12965// CHECK1-NEXT: store i32 0, ptr [[TMP66]], align 4966// CHECK1-NEXT: [[TMP67:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 -1, i32 -1, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z4tbarIsEiT_S0_S0_S0__l145.region_id, ptr [[KERNEL_ARGS]])967// CHECK1-NEXT: [[TMP68:%.*]] = icmp ne i32 [[TMP67]], 0968// CHECK1-NEXT: br i1 [[TMP68]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]969// CHECK1: omp_offload.failed:970// CHECK1-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z4tbarIsEiT_S0_S0_S0__l145(i64 [[TMP5]], i64 [[TMP7]], i64 [[TMP9]], i64 [[TMP11]], i64 [[TMP13]], i64 [[TMP15]], i64 [[TMP17]], i64 [[TMP19]], i64 [[TMP21]]) #[[ATTR2]]971// CHECK1-NEXT: br label [[OMP_OFFLOAD_CONT]]972// CHECK1: omp_offload.cont:973// CHECK1-NEXT: br label [[OMP_IF_END:%.*]]974// CHECK1: omp_if.else:975// CHECK1-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z4tbarIsEiT_S0_S0_S0__l145(i64 [[TMP5]], i64 [[TMP7]], i64 [[TMP9]], i64 [[TMP11]], i64 [[TMP13]], i64 [[TMP15]], i64 [[TMP17]], i64 [[TMP19]], i64 [[TMP21]]) #[[ATTR2]]976// CHECK1-NEXT: br label [[OMP_IF_END]]977// CHECK1: omp_if.end:978// CHECK1-NEXT: ret void979//980//981// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z4tbarIsEiT_S0_S0_S0__l145982// CHECK1-SAME: (i64 noundef [[B:%.*]], i64 noundef [[GB:%.*]], i64 noundef [[SB:%.*]], i64 noundef [[GC:%.*]], i64 noundef [[C:%.*]], i64 noundef [[SC:%.*]], i64 noundef [[D:%.*]], i64 noundef [[GD:%.*]], i64 noundef [[SD:%.*]]) #[[ATTR1]] {983// CHECK1-NEXT: entry:984// CHECK1-NEXT: [[B_ADDR:%.*]] = alloca i64, align 8985// CHECK1-NEXT: [[GB_ADDR:%.*]] = alloca i64, align 8986// CHECK1-NEXT: [[SB_ADDR:%.*]] = alloca i64, align 8987// CHECK1-NEXT: [[GC_ADDR:%.*]] = alloca i64, align 8988// CHECK1-NEXT: [[C_ADDR:%.*]] = alloca i64, align 8989// CHECK1-NEXT: [[SC_ADDR:%.*]] = alloca i64, align 8990// CHECK1-NEXT: [[D_ADDR:%.*]] = alloca i64, align 8991// CHECK1-NEXT: [[GD_ADDR:%.*]] = alloca i64, align 8992// CHECK1-NEXT: [[SD_ADDR:%.*]] = alloca i64, align 8993// CHECK1-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 4994// CHECK1-NEXT: [[DOTBOUND_ZERO_ADDR:%.*]] = alloca i32, align 4995// CHECK1-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])996// CHECK1-NEXT: store i64 [[B]], ptr [[B_ADDR]], align 8997// CHECK1-NEXT: store i64 [[GB]], ptr [[GB_ADDR]], align 8998// CHECK1-NEXT: store i64 [[SB]], ptr [[SB_ADDR]], align 8999// CHECK1-NEXT: store i64 [[GC]], ptr [[GC_ADDR]], align 81000// CHECK1-NEXT: store i64 [[C]], ptr [[C_ADDR]], align 81001// CHECK1-NEXT: store i64 [[SC]], ptr [[SC_ADDR]], align 81002// CHECK1-NEXT: store i64 [[D]], ptr [[D_ADDR]], align 81003// CHECK1-NEXT: store i64 [[GD]], ptr [[GD_ADDR]], align 81004// CHECK1-NEXT: store i64 [[SD]], ptr [[SD_ADDR]], align 81005// CHECK1-NEXT: [[TMP1:%.*]] = load i16, ptr [[B_ADDR]], align 21006// CHECK1-NEXT: [[CONV:%.*]] = sext i16 [[TMP1]] to i321007// CHECK1-NEXT: [[ADD:%.*]] = add nsw i32 [[CONV]], 11008// CHECK1-NEXT: [[CONV1:%.*]] = trunc i32 [[ADD]] to i161009// CHECK1-NEXT: store i16 [[CONV1]], ptr [[B_ADDR]], align 21010// CHECK1-NEXT: [[TMP2:%.*]] = load double, ptr [[GB_ADDR]], align 81011// CHECK1-NEXT: [[ADD2:%.*]] = fadd double [[TMP2]], 1.000000e+001012// CHECK1-NEXT: store double [[ADD2]], ptr [[GB_ADDR]], align 81013// CHECK1-NEXT: [[TMP3:%.*]] = load float, ptr [[SB_ADDR]], align 41014// CHECK1-NEXT: [[CONV3:%.*]] = fpext float [[TMP3]] to double1015// CHECK1-NEXT: [[ADD4:%.*]] = fadd double [[CONV3]], 1.000000e+001016// CHECK1-NEXT: [[CONV5:%.*]] = fptrunc double [[ADD4]] to float1017// CHECK1-NEXT: store float [[CONV5]], ptr [[SB_ADDR]], align 41018// CHECK1-NEXT: [[TMP4:%.*]] = load double, ptr [[GC_ADDR]], align 81019// CHECK1-NEXT: [[CMP:%.*]] = fcmp ogt double [[TMP4]], 0.000000e+001020// CHECK1-NEXT: br i1 [[CMP]], label [[LAND_LHS_TRUE:%.*]], label [[OMP_IF_ELSE:%.*]]1021// CHECK1: land.lhs.true:1022// CHECK1-NEXT: [[TMP5:%.*]] = load i16, ptr [[C_ADDR]], align 21023// CHECK1-NEXT: [[CONV6:%.*]] = sext i16 [[TMP5]] to i321024// CHECK1-NEXT: [[CMP7:%.*]] = icmp sgt i32 [[CONV6]], 01025// CHECK1-NEXT: br i1 [[CMP7]], label [[LAND_LHS_TRUE8:%.*]], label [[OMP_IF_ELSE]]1026// CHECK1: land.lhs.true8:1027// CHECK1-NEXT: [[TMP6:%.*]] = load float, ptr [[SC_ADDR]], align 41028// CHECK1-NEXT: [[CONV9:%.*]] = fpext float [[TMP6]] to double1029// CHECK1-NEXT: [[CMP10:%.*]] = fcmp ogt double [[CONV9]], 0.000000e+001030// CHECK1-NEXT: br i1 [[CMP10]], label [[OMP_IF_THEN:%.*]], label [[OMP_IF_ELSE]]1031// CHECK1: omp_if.then:1032// CHECK1-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB1]], i32 3, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z4tbarIsEiT_S0_S0_S0__l145.omp_outlined, ptr [[D_ADDR]], ptr [[GD_ADDR]], ptr [[SD_ADDR]])1033// CHECK1-NEXT: br label [[OMP_IF_END:%.*]]1034// CHECK1: omp_if.else:1035// CHECK1-NEXT: call void @__kmpc_serialized_parallel(ptr @[[GLOB1]], i32 [[TMP0]])1036// CHECK1-NEXT: store i32 [[TMP0]], ptr [[DOTTHREADID_TEMP_]], align 41037// CHECK1-NEXT: store i32 0, ptr [[DOTBOUND_ZERO_ADDR]], align 41038// CHECK1-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z4tbarIsEiT_S0_S0_S0__l145.omp_outlined(ptr [[DOTTHREADID_TEMP_]], ptr [[DOTBOUND_ZERO_ADDR]], ptr [[D_ADDR]], ptr [[GD_ADDR]], ptr [[SD_ADDR]]) #[[ATTR2]]1039// CHECK1-NEXT: call void @__kmpc_end_serialized_parallel(ptr @[[GLOB1]], i32 [[TMP0]])1040// CHECK1-NEXT: br label [[OMP_IF_END]]1041// CHECK1: omp_if.end:1042// CHECK1-NEXT: ret void1043//1044//1045// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z4tbarIsEiT_S0_S0_S0__l145.omp_outlined1046// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 2 dereferenceable(2) [[D:%.*]], ptr noundef nonnull align 8 dereferenceable(8) [[GD:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[SD:%.*]]) #[[ATTR1]] {1047// CHECK1-NEXT: entry:1048// CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 81049// CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 81050// CHECK1-NEXT: [[D_ADDR:%.*]] = alloca ptr, align 81051// CHECK1-NEXT: [[GD_ADDR:%.*]] = alloca ptr, align 81052// CHECK1-NEXT: [[SD_ADDR:%.*]] = alloca ptr, align 81053// CHECK1-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 81054// CHECK1-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 81055// CHECK1-NEXT: store ptr [[D]], ptr [[D_ADDR]], align 81056// CHECK1-NEXT: store ptr [[GD]], ptr [[GD_ADDR]], align 81057// CHECK1-NEXT: store ptr [[SD]], ptr [[SD_ADDR]], align 81058// CHECK1-NEXT: [[TMP0:%.*]] = load ptr, ptr [[D_ADDR]], align 81059// CHECK1-NEXT: [[TMP1:%.*]] = load ptr, ptr [[GD_ADDR]], align 81060// CHECK1-NEXT: [[TMP2:%.*]] = load ptr, ptr [[SD_ADDR]], align 81061// CHECK1-NEXT: [[TMP3:%.*]] = load i16, ptr [[TMP0]], align 21062// CHECK1-NEXT: [[CONV:%.*]] = sext i16 [[TMP3]] to i321063// CHECK1-NEXT: [[ADD:%.*]] = add nsw i32 [[CONV]], 11064// CHECK1-NEXT: [[CONV1:%.*]] = trunc i32 [[ADD]] to i161065// CHECK1-NEXT: store i16 [[CONV1]], ptr [[TMP0]], align 21066// CHECK1-NEXT: [[TMP4:%.*]] = load double, ptr [[TMP1]], align 81067// CHECK1-NEXT: [[ADD2:%.*]] = fadd double [[TMP4]], 1.000000e+001068// CHECK1-NEXT: store double [[ADD2]], ptr [[TMP1]], align 81069// CHECK1-NEXT: [[TMP5:%.*]] = load float, ptr [[TMP2]], align 41070// CHECK1-NEXT: [[CONV3:%.*]] = fpext float [[TMP5]] to double1071// CHECK1-NEXT: [[ADD4:%.*]] = fadd double [[CONV3]], 1.000000e+001072// CHECK1-NEXT: [[CONV5:%.*]] = fptrunc double [[ADD4]] to float1073// CHECK1-NEXT: store float [[CONV5]], ptr [[TMP2]], align 41074// CHECK1-NEXT: ret void1075//1076//1077// CHECK3-LABEL: define {{[^@]+}}@_Z3foossss1078// CHECK3-SAME: (i16 noundef signext [[A:%.*]], i16 noundef signext [[B:%.*]], i16 noundef signext [[C:%.*]], i16 noundef signext [[D:%.*]]) #[[ATTR0:[0-9]+]] {1079// CHECK3-NEXT: entry:1080// CHECK3-NEXT: [[A_ADDR:%.*]] = alloca i16, align 21081// CHECK3-NEXT: [[B_ADDR:%.*]] = alloca i16, align 21082// CHECK3-NEXT: [[C_ADDR:%.*]] = alloca i16, align 21083// CHECK3-NEXT: [[D_ADDR:%.*]] = alloca i16, align 21084// CHECK3-NEXT: [[B_CASTED:%.*]] = alloca i32, align 41085// CHECK3-NEXT: [[SB_CASTED:%.*]] = alloca i32, align 41086// CHECK3-NEXT: [[C_CASTED:%.*]] = alloca i32, align 41087// CHECK3-NEXT: [[SC_CASTED:%.*]] = alloca i32, align 41088// CHECK3-NEXT: [[D_CASTED:%.*]] = alloca i32, align 41089// CHECK3-NEXT: [[SD_CASTED:%.*]] = alloca i32, align 41090// CHECK3-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [9 x ptr], align 41091// CHECK3-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [9 x ptr], align 41092// CHECK3-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [9 x ptr], align 41093// CHECK3-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 81094// CHECK3-NEXT: store i16 [[A]], ptr [[A_ADDR]], align 21095// CHECK3-NEXT: store i16 [[B]], ptr [[B_ADDR]], align 21096// CHECK3-NEXT: store i16 [[C]], ptr [[C_ADDR]], align 21097// CHECK3-NEXT: store i16 [[D]], ptr [[D_ADDR]], align 21098// CHECK3-NEXT: [[TMP0:%.*]] = load i16, ptr [[B_ADDR]], align 21099// CHECK3-NEXT: store i16 [[TMP0]], ptr [[B_CASTED]], align 21100// CHECK3-NEXT: [[TMP1:%.*]] = load i32, ptr [[B_CASTED]], align 41101// CHECK3-NEXT: [[TMP2:%.*]] = load float, ptr @_ZZ3foossssE2Sb, align 41102// CHECK3-NEXT: store float [[TMP2]], ptr [[SB_CASTED]], align 41103// CHECK3-NEXT: [[TMP3:%.*]] = load i32, ptr [[SB_CASTED]], align 41104// CHECK3-NEXT: [[TMP4:%.*]] = load i16, ptr [[C_ADDR]], align 21105// CHECK3-NEXT: store i16 [[TMP4]], ptr [[C_CASTED]], align 21106// CHECK3-NEXT: [[TMP5:%.*]] = load i32, ptr [[C_CASTED]], align 41107// CHECK3-NEXT: [[TMP6:%.*]] = load float, ptr @_ZZ3foossssE2Sc, align 41108// CHECK3-NEXT: store float [[TMP6]], ptr [[SC_CASTED]], align 41109// CHECK3-NEXT: [[TMP7:%.*]] = load i32, ptr [[SC_CASTED]], align 41110// CHECK3-NEXT: [[TMP8:%.*]] = load i16, ptr [[D_ADDR]], align 21111// CHECK3-NEXT: store i16 [[TMP8]], ptr [[D_CASTED]], align 21112// CHECK3-NEXT: [[TMP9:%.*]] = load i32, ptr [[D_CASTED]], align 41113// CHECK3-NEXT: [[TMP10:%.*]] = load float, ptr @_ZZ3foossssE2Sd, align 41114// CHECK3-NEXT: store float [[TMP10]], ptr [[SD_CASTED]], align 41115// CHECK3-NEXT: [[TMP11:%.*]] = load i32, ptr [[SD_CASTED]], align 41116// CHECK3-NEXT: [[TMP12:%.*]] = load double, ptr @Ga, align 81117// CHECK3-NEXT: [[CMP:%.*]] = fcmp ogt double [[TMP12]], 0.000000e+001118// CHECK3-NEXT: br i1 [[CMP]], label [[LAND_LHS_TRUE:%.*]], label [[OMP_IF_ELSE:%.*]]1119// CHECK3: land.lhs.true:1120// CHECK3-NEXT: [[TMP13:%.*]] = load i16, ptr [[A_ADDR]], align 21121// CHECK3-NEXT: [[CONV:%.*]] = sext i16 [[TMP13]] to i321122// CHECK3-NEXT: [[CMP1:%.*]] = icmp sgt i32 [[CONV]], 01123// CHECK3-NEXT: br i1 [[CMP1]], label [[LAND_LHS_TRUE2:%.*]], label [[OMP_IF_ELSE]]1124// CHECK3: land.lhs.true2:1125// CHECK3-NEXT: [[TMP14:%.*]] = load float, ptr @_ZZ3foossssE2Sa, align 41126// CHECK3-NEXT: [[CONV3:%.*]] = fpext float [[TMP14]] to double1127// CHECK3-NEXT: [[CMP4:%.*]] = fcmp ogt double [[CONV3]], 0.000000e+001128// CHECK3-NEXT: br i1 [[CMP4]], label [[OMP_IF_THEN:%.*]], label [[OMP_IF_ELSE]]1129// CHECK3: omp_if.then:1130// CHECK3-NEXT: [[TMP15:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 01131// CHECK3-NEXT: store i32 [[TMP1]], ptr [[TMP15]], align 41132// CHECK3-NEXT: [[TMP16:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 01133// CHECK3-NEXT: store i32 [[TMP1]], ptr [[TMP16]], align 41134// CHECK3-NEXT: [[TMP17:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 01135// CHECK3-NEXT: store ptr null, ptr [[TMP17]], align 41136// CHECK3-NEXT: [[TMP18:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 11137// CHECK3-NEXT: store ptr @Gb, ptr [[TMP18]], align 41138// CHECK3-NEXT: [[TMP19:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 11139// CHECK3-NEXT: store ptr @Gb, ptr [[TMP19]], align 41140// CHECK3-NEXT: [[TMP20:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 11141// CHECK3-NEXT: store ptr null, ptr [[TMP20]], align 41142// CHECK3-NEXT: [[TMP21:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 21143// CHECK3-NEXT: store i32 [[TMP3]], ptr [[TMP21]], align 41144// CHECK3-NEXT: [[TMP22:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 21145// CHECK3-NEXT: store i32 [[TMP3]], ptr [[TMP22]], align 41146// CHECK3-NEXT: [[TMP23:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 21147// CHECK3-NEXT: store ptr null, ptr [[TMP23]], align 41148// CHECK3-NEXT: [[TMP24:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 31149// CHECK3-NEXT: store ptr @Gc, ptr [[TMP24]], align 41150// CHECK3-NEXT: [[TMP25:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 31151// CHECK3-NEXT: store ptr @Gc, ptr [[TMP25]], align 41152// CHECK3-NEXT: [[TMP26:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 31153// CHECK3-NEXT: store ptr null, ptr [[TMP26]], align 41154// CHECK3-NEXT: [[TMP27:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 41155// CHECK3-NEXT: store i32 [[TMP5]], ptr [[TMP27]], align 41156// CHECK3-NEXT: [[TMP28:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 41157// CHECK3-NEXT: store i32 [[TMP5]], ptr [[TMP28]], align 41158// CHECK3-NEXT: [[TMP29:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 41159// CHECK3-NEXT: store ptr null, ptr [[TMP29]], align 41160// CHECK3-NEXT: [[TMP30:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 51161// CHECK3-NEXT: store i32 [[TMP7]], ptr [[TMP30]], align 41162// CHECK3-NEXT: [[TMP31:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 51163// CHECK3-NEXT: store i32 [[TMP7]], ptr [[TMP31]], align 41164// CHECK3-NEXT: [[TMP32:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 51165// CHECK3-NEXT: store ptr null, ptr [[TMP32]], align 41166// CHECK3-NEXT: [[TMP33:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 61167// CHECK3-NEXT: store i32 [[TMP9]], ptr [[TMP33]], align 41168// CHECK3-NEXT: [[TMP34:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 61169// CHECK3-NEXT: store i32 [[TMP9]], ptr [[TMP34]], align 41170// CHECK3-NEXT: [[TMP35:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 61171// CHECK3-NEXT: store ptr null, ptr [[TMP35]], align 41172// CHECK3-NEXT: [[TMP36:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 71173// CHECK3-NEXT: store ptr @Gd, ptr [[TMP36]], align 41174// CHECK3-NEXT: [[TMP37:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 71175// CHECK3-NEXT: store ptr @Gd, ptr [[TMP37]], align 41176// CHECK3-NEXT: [[TMP38:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 71177// CHECK3-NEXT: store ptr null, ptr [[TMP38]], align 41178// CHECK3-NEXT: [[TMP39:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 81179// CHECK3-NEXT: store i32 [[TMP11]], ptr [[TMP39]], align 41180// CHECK3-NEXT: [[TMP40:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 81181// CHECK3-NEXT: store i32 [[TMP11]], ptr [[TMP40]], align 41182// CHECK3-NEXT: [[TMP41:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 81183// CHECK3-NEXT: store ptr null, ptr [[TMP41]], align 41184// CHECK3-NEXT: [[TMP42:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 01185// CHECK3-NEXT: [[TMP43:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 01186// CHECK3-NEXT: [[TMP44:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 01187// CHECK3-NEXT: store i32 3, ptr [[TMP44]], align 41188// CHECK3-NEXT: [[TMP45:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11189// CHECK3-NEXT: store i32 9, ptr [[TMP45]], align 41190// CHECK3-NEXT: [[TMP46:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 21191// CHECK3-NEXT: store ptr [[TMP42]], ptr [[TMP46]], align 41192// CHECK3-NEXT: [[TMP47:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 31193// CHECK3-NEXT: store ptr [[TMP43]], ptr [[TMP47]], align 41194// CHECK3-NEXT: [[TMP48:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 41195// CHECK3-NEXT: store ptr @.offload_sizes, ptr [[TMP48]], align 41196// CHECK3-NEXT: [[TMP49:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 51197// CHECK3-NEXT: store ptr @.offload_maptypes, ptr [[TMP49]], align 41198// CHECK3-NEXT: [[TMP50:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 61199// CHECK3-NEXT: store ptr null, ptr [[TMP50]], align 41200// CHECK3-NEXT: [[TMP51:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 71201// CHECK3-NEXT: store ptr null, ptr [[TMP51]], align 41202// CHECK3-NEXT: [[TMP52:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 81203// CHECK3-NEXT: store i64 0, ptr [[TMP52]], align 81204// CHECK3-NEXT: [[TMP53:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 91205// CHECK3-NEXT: store i64 0, ptr [[TMP53]], align 81206// CHECK3-NEXT: [[TMP54:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 101207// CHECK3-NEXT: store [3 x i32] [i32 -1, i32 0, i32 0], ptr [[TMP54]], align 41208// CHECK3-NEXT: [[TMP55:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 111209// CHECK3-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP55]], align 41210// CHECK3-NEXT: [[TMP56:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 121211// CHECK3-NEXT: store i32 0, ptr [[TMP56]], align 41212// CHECK3-NEXT: [[TMP57:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1:[0-9]+]], i64 -1, i32 -1, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3foossss_l49.region_id, ptr [[KERNEL_ARGS]])1213// CHECK3-NEXT: [[TMP58:%.*]] = icmp ne i32 [[TMP57]], 01214// CHECK3-NEXT: br i1 [[TMP58]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]1215// CHECK3: omp_offload.failed:1216// CHECK3-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3foossss_l49(i32 [[TMP1]], ptr @Gb, i32 [[TMP3]], ptr @Gc, i32 [[TMP5]], i32 [[TMP7]], i32 [[TMP9]], ptr @Gd, i32 [[TMP11]]) #[[ATTR2:[0-9]+]]1217// CHECK3-NEXT: br label [[OMP_OFFLOAD_CONT]]1218// CHECK3: omp_offload.cont:1219// CHECK3-NEXT: br label [[OMP_IF_END:%.*]]1220// CHECK3: omp_if.else:1221// CHECK3-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3foossss_l49(i32 [[TMP1]], ptr @Gb, i32 [[TMP3]], ptr @Gc, i32 [[TMP5]], i32 [[TMP7]], i32 [[TMP9]], ptr @Gd, i32 [[TMP11]]) #[[ATTR2]]1222// CHECK3-NEXT: br label [[OMP_IF_END]]1223// CHECK3: omp_if.end:1224// CHECK3-NEXT: [[TMP59:%.*]] = load i16, ptr [[A_ADDR]], align 21225// CHECK3-NEXT: [[CONV5:%.*]] = sext i16 [[TMP59]] to i321226// CHECK3-NEXT: [[TMP60:%.*]] = load i16, ptr [[B_ADDR]], align 21227// CHECK3-NEXT: [[CONV6:%.*]] = sext i16 [[TMP60]] to i321228// CHECK3-NEXT: [[ADD:%.*]] = add nsw i32 [[CONV5]], [[CONV6]]1229// CHECK3-NEXT: [[TMP61:%.*]] = load i16, ptr [[C_ADDR]], align 21230// CHECK3-NEXT: [[CONV7:%.*]] = sext i16 [[TMP61]] to i321231// CHECK3-NEXT: [[ADD8:%.*]] = add nsw i32 [[ADD]], [[CONV7]]1232// CHECK3-NEXT: [[TMP62:%.*]] = load i16, ptr [[D_ADDR]], align 21233// CHECK3-NEXT: [[CONV9:%.*]] = sext i16 [[TMP62]] to i321234// CHECK3-NEXT: [[ADD10:%.*]] = add nsw i32 [[ADD8]], [[CONV9]]1235// CHECK3-NEXT: [[TMP63:%.*]] = load float, ptr @_ZZ3foossssE2Sa, align 41236// CHECK3-NEXT: [[CONV11:%.*]] = fptosi float [[TMP63]] to i321237// CHECK3-NEXT: [[ADD12:%.*]] = add nsw i32 [[ADD10]], [[CONV11]]1238// CHECK3-NEXT: [[TMP64:%.*]] = load float, ptr @_ZZ3foossssE2Sb, align 41239// CHECK3-NEXT: [[CONV13:%.*]] = fptosi float [[TMP64]] to i321240// CHECK3-NEXT: [[ADD14:%.*]] = add nsw i32 [[ADD12]], [[CONV13]]1241// CHECK3-NEXT: [[TMP65:%.*]] = load float, ptr @_ZZ3foossssE2Sc, align 41242// CHECK3-NEXT: [[CONV15:%.*]] = fptosi float [[TMP65]] to i321243// CHECK3-NEXT: [[ADD16:%.*]] = add nsw i32 [[ADD14]], [[CONV15]]1244// CHECK3-NEXT: [[TMP66:%.*]] = load float, ptr @_ZZ3foossssE2Sd, align 41245// CHECK3-NEXT: [[CONV17:%.*]] = fptosi float [[TMP66]] to i321246// CHECK3-NEXT: [[ADD18:%.*]] = add nsw i32 [[ADD16]], [[CONV17]]1247// CHECK3-NEXT: ret i32 [[ADD18]]1248//1249//1250// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3foossss_l491251// CHECK3-SAME: (i32 noundef [[B:%.*]], ptr noundef nonnull align 4 dereferenceable(8) [[GB:%.*]], i32 noundef [[SB:%.*]], ptr noundef nonnull align 4 dereferenceable(8) [[GC:%.*]], i32 noundef [[C:%.*]], i32 noundef [[SC:%.*]], i32 noundef [[D:%.*]], ptr noundef nonnull align 4 dereferenceable(8) [[GD:%.*]], i32 noundef [[SD:%.*]]) #[[ATTR1:[0-9]+]] {1252// CHECK3-NEXT: entry:1253// CHECK3-NEXT: [[B_ADDR:%.*]] = alloca i32, align 41254// CHECK3-NEXT: [[GB_ADDR:%.*]] = alloca ptr, align 41255// CHECK3-NEXT: [[SB_ADDR:%.*]] = alloca i32, align 41256// CHECK3-NEXT: [[GC_ADDR:%.*]] = alloca ptr, align 41257// CHECK3-NEXT: [[C_ADDR:%.*]] = alloca i32, align 41258// CHECK3-NEXT: [[SC_ADDR:%.*]] = alloca i32, align 41259// CHECK3-NEXT: [[D_ADDR:%.*]] = alloca i32, align 41260// CHECK3-NEXT: [[GD_ADDR:%.*]] = alloca ptr, align 41261// CHECK3-NEXT: [[SD_ADDR:%.*]] = alloca i32, align 41262// CHECK3-NEXT: [[GB1:%.*]] = alloca double, align 81263// CHECK3-NEXT: [[GC2:%.*]] = alloca double, align 81264// CHECK3-NEXT: [[GD3:%.*]] = alloca double, align 81265// CHECK3-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 41266// CHECK3-NEXT: [[DOTBOUND_ZERO_ADDR:%.*]] = alloca i32, align 41267// CHECK3-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])1268// CHECK3-NEXT: store i32 [[B]], ptr [[B_ADDR]], align 41269// CHECK3-NEXT: store ptr [[GB]], ptr [[GB_ADDR]], align 41270// CHECK3-NEXT: store i32 [[SB]], ptr [[SB_ADDR]], align 41271// CHECK3-NEXT: store ptr [[GC]], ptr [[GC_ADDR]], align 41272// CHECK3-NEXT: store i32 [[C]], ptr [[C_ADDR]], align 41273// CHECK3-NEXT: store i32 [[SC]], ptr [[SC_ADDR]], align 41274// CHECK3-NEXT: store i32 [[D]], ptr [[D_ADDR]], align 41275// CHECK3-NEXT: store ptr [[GD]], ptr [[GD_ADDR]], align 41276// CHECK3-NEXT: store i32 [[SD]], ptr [[SD_ADDR]], align 41277// CHECK3-NEXT: [[TMP1:%.*]] = load ptr, ptr [[GB_ADDR]], align 41278// CHECK3-NEXT: [[TMP2:%.*]] = load ptr, ptr [[GC_ADDR]], align 41279// CHECK3-NEXT: [[TMP3:%.*]] = load ptr, ptr [[GD_ADDR]], align 41280// CHECK3-NEXT: [[TMP4:%.*]] = load double, ptr [[TMP1]], align 81281// CHECK3-NEXT: store double [[TMP4]], ptr [[GB1]], align 81282// CHECK3-NEXT: [[TMP5:%.*]] = load double, ptr [[TMP2]], align 81283// CHECK3-NEXT: store double [[TMP5]], ptr [[GC2]], align 81284// CHECK3-NEXT: [[TMP6:%.*]] = load double, ptr [[TMP3]], align 81285// CHECK3-NEXT: store double [[TMP6]], ptr [[GD3]], align 81286// CHECK3-NEXT: [[TMP7:%.*]] = load i16, ptr [[B_ADDR]], align 21287// CHECK3-NEXT: [[CONV:%.*]] = sext i16 [[TMP7]] to i321288// CHECK3-NEXT: [[ADD:%.*]] = add nsw i32 [[CONV]], 11289// CHECK3-NEXT: [[CONV4:%.*]] = trunc i32 [[ADD]] to i161290// CHECK3-NEXT: store i16 [[CONV4]], ptr [[B_ADDR]], align 21291// CHECK3-NEXT: [[TMP8:%.*]] = load double, ptr [[GB1]], align 81292// CHECK3-NEXT: [[ADD5:%.*]] = fadd double [[TMP8]], 1.000000e+001293// CHECK3-NEXT: store double [[ADD5]], ptr [[GB1]], align 81294// CHECK3-NEXT: [[TMP9:%.*]] = load float, ptr [[SB_ADDR]], align 41295// CHECK3-NEXT: [[CONV6:%.*]] = fpext float [[TMP9]] to double1296// CHECK3-NEXT: [[ADD7:%.*]] = fadd double [[CONV6]], 1.000000e+001297// CHECK3-NEXT: [[CONV8:%.*]] = fptrunc double [[ADD7]] to float1298// CHECK3-NEXT: store float [[CONV8]], ptr [[SB_ADDR]], align 41299// CHECK3-NEXT: [[TMP10:%.*]] = load double, ptr [[GC2]], align 81300// CHECK3-NEXT: [[CMP:%.*]] = fcmp ogt double [[TMP10]], 0.000000e+001301// CHECK3-NEXT: br i1 [[CMP]], label [[LAND_LHS_TRUE:%.*]], label [[OMP_IF_ELSE:%.*]]1302// CHECK3: land.lhs.true:1303// CHECK3-NEXT: [[TMP11:%.*]] = load i16, ptr [[C_ADDR]], align 21304// CHECK3-NEXT: [[CONV9:%.*]] = sext i16 [[TMP11]] to i321305// CHECK3-NEXT: [[CMP10:%.*]] = icmp sgt i32 [[CONV9]], 01306// CHECK3-NEXT: br i1 [[CMP10]], label [[LAND_LHS_TRUE11:%.*]], label [[OMP_IF_ELSE]]1307// CHECK3: land.lhs.true11:1308// CHECK3-NEXT: [[TMP12:%.*]] = load float, ptr [[SC_ADDR]], align 41309// CHECK3-NEXT: [[CONV12:%.*]] = fpext float [[TMP12]] to double1310// CHECK3-NEXT: [[CMP13:%.*]] = fcmp ogt double [[CONV12]], 0.000000e+001311// CHECK3-NEXT: br i1 [[CMP13]], label [[OMP_IF_THEN:%.*]], label [[OMP_IF_ELSE]]1312// CHECK3: omp_if.then:1313// CHECK3-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB1]], i32 3, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3foossss_l49.omp_outlined, ptr [[D_ADDR]], ptr [[GD3]], ptr [[SD_ADDR]])1314// CHECK3-NEXT: br label [[OMP_IF_END:%.*]]1315// CHECK3: omp_if.else:1316// CHECK3-NEXT: call void @__kmpc_serialized_parallel(ptr @[[GLOB1]], i32 [[TMP0]])1317// CHECK3-NEXT: store i32 [[TMP0]], ptr [[DOTTHREADID_TEMP_]], align 41318// CHECK3-NEXT: store i32 0, ptr [[DOTBOUND_ZERO_ADDR]], align 41319// CHECK3-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3foossss_l49.omp_outlined(ptr [[DOTTHREADID_TEMP_]], ptr [[DOTBOUND_ZERO_ADDR]], ptr [[D_ADDR]], ptr [[GD3]], ptr [[SD_ADDR]]) #[[ATTR2]]1320// CHECK3-NEXT: call void @__kmpc_end_serialized_parallel(ptr @[[GLOB1]], i32 [[TMP0]])1321// CHECK3-NEXT: br label [[OMP_IF_END]]1322// CHECK3: omp_if.end:1323// CHECK3-NEXT: ret void1324//1325//1326// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3foossss_l49.omp_outlined1327// CHECK3-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 2 dereferenceable(2) [[D:%.*]], ptr noundef nonnull align 4 dereferenceable(8) [[GD:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[SD:%.*]]) #[[ATTR1]] {1328// CHECK3-NEXT: entry:1329// CHECK3-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 41330// CHECK3-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 41331// CHECK3-NEXT: [[D_ADDR:%.*]] = alloca ptr, align 41332// CHECK3-NEXT: [[GD_ADDR:%.*]] = alloca ptr, align 41333// CHECK3-NEXT: [[SD_ADDR:%.*]] = alloca ptr, align 41334// CHECK3-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 41335// CHECK3-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 41336// CHECK3-NEXT: store ptr [[D]], ptr [[D_ADDR]], align 41337// CHECK3-NEXT: store ptr [[GD]], ptr [[GD_ADDR]], align 41338// CHECK3-NEXT: store ptr [[SD]], ptr [[SD_ADDR]], align 41339// CHECK3-NEXT: [[TMP0:%.*]] = load ptr, ptr [[D_ADDR]], align 41340// CHECK3-NEXT: [[TMP1:%.*]] = load ptr, ptr [[GD_ADDR]], align 41341// CHECK3-NEXT: [[TMP2:%.*]] = load ptr, ptr [[SD_ADDR]], align 41342// CHECK3-NEXT: [[TMP3:%.*]] = load i16, ptr [[TMP0]], align 21343// CHECK3-NEXT: [[CONV:%.*]] = sext i16 [[TMP3]] to i321344// CHECK3-NEXT: [[ADD:%.*]] = add nsw i32 [[CONV]], 11345// CHECK3-NEXT: [[CONV1:%.*]] = trunc i32 [[ADD]] to i161346// CHECK3-NEXT: store i16 [[CONV1]], ptr [[TMP0]], align 21347// CHECK3-NEXT: [[TMP4:%.*]] = load double, ptr [[TMP1]], align 81348// CHECK3-NEXT: [[ADD2:%.*]] = fadd double [[TMP4]], 1.000000e+001349// CHECK3-NEXT: store double [[ADD2]], ptr [[TMP1]], align 81350// CHECK3-NEXT: [[TMP5:%.*]] = load float, ptr [[TMP2]], align 41351// CHECK3-NEXT: [[CONV3:%.*]] = fpext float [[TMP5]] to double1352// CHECK3-NEXT: [[ADD4:%.*]] = fadd double [[CONV3]], 1.000000e+001353// CHECK3-NEXT: [[CONV5:%.*]] = fptrunc double [[ADD4]] to float1354// CHECK3-NEXT: store float [[CONV5]], ptr [[TMP2]], align 41355// CHECK3-NEXT: ret void1356//1357//1358// CHECK3-LABEL: define {{[^@]+}}@_Z3barssss1359// CHECK3-SAME: (i16 noundef signext [[A:%.*]], i16 noundef signext [[B:%.*]], i16 noundef signext [[C:%.*]], i16 noundef signext [[D:%.*]]) #[[ATTR0]] {1360// CHECK3-NEXT: entry:1361// CHECK3-NEXT: [[A_ADDR:%.*]] = alloca i16, align 21362// CHECK3-NEXT: [[B_ADDR:%.*]] = alloca i16, align 21363// CHECK3-NEXT: [[C_ADDR:%.*]] = alloca i16, align 21364// CHECK3-NEXT: [[D_ADDR:%.*]] = alloca i16, align 21365// CHECK3-NEXT: store i16 [[A]], ptr [[A_ADDR]], align 21366// CHECK3-NEXT: store i16 [[B]], ptr [[B_ADDR]], align 21367// CHECK3-NEXT: store i16 [[C]], ptr [[C_ADDR]], align 21368// CHECK3-NEXT: store i16 [[D]], ptr [[D_ADDR]], align 21369// CHECK3-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB1]], i32 4, ptr @_Z3barssss.omp_outlined, ptr [[A_ADDR]], ptr [[B_ADDR]], ptr [[C_ADDR]], ptr [[D_ADDR]])1370// CHECK3-NEXT: [[TMP0:%.*]] = load i16, ptr [[A_ADDR]], align 21371// CHECK3-NEXT: [[CONV:%.*]] = sext i16 [[TMP0]] to i321372// CHECK3-NEXT: [[TMP1:%.*]] = load i16, ptr [[B_ADDR]], align 21373// CHECK3-NEXT: [[CONV1:%.*]] = sext i16 [[TMP1]] to i321374// CHECK3-NEXT: [[ADD:%.*]] = add nsw i32 [[CONV]], [[CONV1]]1375// CHECK3-NEXT: [[TMP2:%.*]] = load i16, ptr [[C_ADDR]], align 21376// CHECK3-NEXT: [[CONV2:%.*]] = sext i16 [[TMP2]] to i321377// CHECK3-NEXT: [[ADD3:%.*]] = add nsw i32 [[ADD]], [[CONV2]]1378// CHECK3-NEXT: [[TMP3:%.*]] = load i16, ptr [[D_ADDR]], align 21379// CHECK3-NEXT: [[CONV4:%.*]] = sext i16 [[TMP3]] to i321380// CHECK3-NEXT: [[ADD5:%.*]] = add nsw i32 [[ADD3]], [[CONV4]]1381// CHECK3-NEXT: [[TMP4:%.*]] = load float, ptr @_ZZ3barssssE2Sa, align 41382// CHECK3-NEXT: [[CONV6:%.*]] = fptosi float [[TMP4]] to i321383// CHECK3-NEXT: [[ADD7:%.*]] = add nsw i32 [[ADD5]], [[CONV6]]1384// CHECK3-NEXT: [[TMP5:%.*]] = load float, ptr @_ZZ3barssssE2Sb, align 41385// CHECK3-NEXT: [[CONV8:%.*]] = fptosi float [[TMP5]] to i321386// CHECK3-NEXT: [[ADD9:%.*]] = add nsw i32 [[ADD7]], [[CONV8]]1387// CHECK3-NEXT: [[TMP6:%.*]] = load float, ptr @_ZZ3barssssE2Sc, align 41388// CHECK3-NEXT: [[CONV10:%.*]] = fptosi float [[TMP6]] to i321389// CHECK3-NEXT: [[ADD11:%.*]] = add nsw i32 [[ADD9]], [[CONV10]]1390// CHECK3-NEXT: [[TMP7:%.*]] = load float, ptr @_ZZ3barssssE2Sd, align 41391// CHECK3-NEXT: [[CONV12:%.*]] = fptosi float [[TMP7]] to i321392// CHECK3-NEXT: [[ADD13:%.*]] = add nsw i32 [[ADD11]], [[CONV12]]1393// CHECK3-NEXT: ret i32 [[ADD13]]1394//1395//1396// CHECK3-LABEL: define {{[^@]+}}@_Z3barssss.omp_outlined1397// CHECK3-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 2 dereferenceable(2) [[A:%.*]], ptr noundef nonnull align 2 dereferenceable(2) [[B:%.*]], ptr noundef nonnull align 2 dereferenceable(2) [[C:%.*]], ptr noundef nonnull align 2 dereferenceable(2) [[D:%.*]]) #[[ATTR1]] {1398// CHECK3-NEXT: entry:1399// CHECK3-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 41400// CHECK3-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 41401// CHECK3-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 41402// CHECK3-NEXT: [[B_ADDR:%.*]] = alloca ptr, align 41403// CHECK3-NEXT: [[C_ADDR:%.*]] = alloca ptr, align 41404// CHECK3-NEXT: [[D_ADDR:%.*]] = alloca ptr, align 41405// CHECK3-NEXT: [[B_CASTED:%.*]] = alloca i32, align 41406// CHECK3-NEXT: [[SB_CASTED:%.*]] = alloca i32, align 41407// CHECK3-NEXT: [[C_CASTED:%.*]] = alloca i32, align 41408// CHECK3-NEXT: [[SC_CASTED:%.*]] = alloca i32, align 41409// CHECK3-NEXT: [[D_CASTED:%.*]] = alloca i32, align 41410// CHECK3-NEXT: [[SD_CASTED:%.*]] = alloca i32, align 41411// CHECK3-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [9 x ptr], align 41412// CHECK3-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [9 x ptr], align 41413// CHECK3-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [9 x ptr], align 41414// CHECK3-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 81415// CHECK3-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 41416// CHECK3-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 41417// CHECK3-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 41418// CHECK3-NEXT: store ptr [[B]], ptr [[B_ADDR]], align 41419// CHECK3-NEXT: store ptr [[C]], ptr [[C_ADDR]], align 41420// CHECK3-NEXT: store ptr [[D]], ptr [[D_ADDR]], align 41421// CHECK3-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 41422// CHECK3-NEXT: [[TMP1:%.*]] = load ptr, ptr [[B_ADDR]], align 41423// CHECK3-NEXT: [[TMP2:%.*]] = load ptr, ptr [[C_ADDR]], align 41424// CHECK3-NEXT: [[TMP3:%.*]] = load ptr, ptr [[D_ADDR]], align 41425// CHECK3-NEXT: [[TMP4:%.*]] = load i16, ptr [[TMP1]], align 21426// CHECK3-NEXT: store i16 [[TMP4]], ptr [[B_CASTED]], align 21427// CHECK3-NEXT: [[TMP5:%.*]] = load i32, ptr [[B_CASTED]], align 41428// CHECK3-NEXT: [[TMP6:%.*]] = load float, ptr @_ZZ3barssssE2Sb, align 41429// CHECK3-NEXT: store float [[TMP6]], ptr [[SB_CASTED]], align 41430// CHECK3-NEXT: [[TMP7:%.*]] = load i32, ptr [[SB_CASTED]], align 41431// CHECK3-NEXT: [[TMP8:%.*]] = load i16, ptr [[TMP2]], align 21432// CHECK3-NEXT: store i16 [[TMP8]], ptr [[C_CASTED]], align 21433// CHECK3-NEXT: [[TMP9:%.*]] = load i32, ptr [[C_CASTED]], align 41434// CHECK3-NEXT: [[TMP10:%.*]] = load float, ptr @_ZZ3barssssE2Sc, align 41435// CHECK3-NEXT: store float [[TMP10]], ptr [[SC_CASTED]], align 41436// CHECK3-NEXT: [[TMP11:%.*]] = load i32, ptr [[SC_CASTED]], align 41437// CHECK3-NEXT: [[TMP12:%.*]] = load i16, ptr [[TMP3]], align 21438// CHECK3-NEXT: store i16 [[TMP12]], ptr [[D_CASTED]], align 21439// CHECK3-NEXT: [[TMP13:%.*]] = load i32, ptr [[D_CASTED]], align 41440// CHECK3-NEXT: [[TMP14:%.*]] = load float, ptr @_ZZ3barssssE2Sd, align 41441// CHECK3-NEXT: store float [[TMP14]], ptr [[SD_CASTED]], align 41442// CHECK3-NEXT: [[TMP15:%.*]] = load i32, ptr [[SD_CASTED]], align 41443// CHECK3-NEXT: [[TMP16:%.*]] = load double, ptr @Ga, align 81444// CHECK3-NEXT: [[CMP:%.*]] = fcmp ogt double [[TMP16]], 0.000000e+001445// CHECK3-NEXT: br i1 [[CMP]], label [[LAND_LHS_TRUE:%.*]], label [[OMP_IF_ELSE:%.*]]1446// CHECK3: land.lhs.true:1447// CHECK3-NEXT: [[TMP17:%.*]] = load i16, ptr [[TMP0]], align 21448// CHECK3-NEXT: [[CONV:%.*]] = sext i16 [[TMP17]] to i321449// CHECK3-NEXT: [[CMP1:%.*]] = icmp sgt i32 [[CONV]], 01450// CHECK3-NEXT: br i1 [[CMP1]], label [[LAND_LHS_TRUE2:%.*]], label [[OMP_IF_ELSE]]1451// CHECK3: land.lhs.true2:1452// CHECK3-NEXT: [[TMP18:%.*]] = load float, ptr @_ZZ3barssssE2Sa, align 41453// CHECK3-NEXT: [[CONV3:%.*]] = fpext float [[TMP18]] to double1454// CHECK3-NEXT: [[CMP4:%.*]] = fcmp ogt double [[CONV3]], 0.000000e+001455// CHECK3-NEXT: br i1 [[CMP4]], label [[OMP_IF_THEN:%.*]], label [[OMP_IF_ELSE]]1456// CHECK3: omp_if.then:1457// CHECK3-NEXT: [[TMP19:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 01458// CHECK3-NEXT: store i32 [[TMP5]], ptr [[TMP19]], align 41459// CHECK3-NEXT: [[TMP20:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 01460// CHECK3-NEXT: store i32 [[TMP5]], ptr [[TMP20]], align 41461// CHECK3-NEXT: [[TMP21:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 01462// CHECK3-NEXT: store ptr null, ptr [[TMP21]], align 41463// CHECK3-NEXT: [[TMP22:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 11464// CHECK3-NEXT: store ptr @Gb, ptr [[TMP22]], align 41465// CHECK3-NEXT: [[TMP23:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 11466// CHECK3-NEXT: store ptr @Gb, ptr [[TMP23]], align 41467// CHECK3-NEXT: [[TMP24:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 11468// CHECK3-NEXT: store ptr null, ptr [[TMP24]], align 41469// CHECK3-NEXT: [[TMP25:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 21470// CHECK3-NEXT: store i32 [[TMP7]], ptr [[TMP25]], align 41471// CHECK3-NEXT: [[TMP26:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 21472// CHECK3-NEXT: store i32 [[TMP7]], ptr [[TMP26]], align 41473// CHECK3-NEXT: [[TMP27:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 21474// CHECK3-NEXT: store ptr null, ptr [[TMP27]], align 41475// CHECK3-NEXT: [[TMP28:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 31476// CHECK3-NEXT: store ptr @Gc, ptr [[TMP28]], align 41477// CHECK3-NEXT: [[TMP29:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 31478// CHECK3-NEXT: store ptr @Gc, ptr [[TMP29]], align 41479// CHECK3-NEXT: [[TMP30:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 31480// CHECK3-NEXT: store ptr null, ptr [[TMP30]], align 41481// CHECK3-NEXT: [[TMP31:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 41482// CHECK3-NEXT: store i32 [[TMP9]], ptr [[TMP31]], align 41483// CHECK3-NEXT: [[TMP32:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 41484// CHECK3-NEXT: store i32 [[TMP9]], ptr [[TMP32]], align 41485// CHECK3-NEXT: [[TMP33:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 41486// CHECK3-NEXT: store ptr null, ptr [[TMP33]], align 41487// CHECK3-NEXT: [[TMP34:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 51488// CHECK3-NEXT: store i32 [[TMP11]], ptr [[TMP34]], align 41489// CHECK3-NEXT: [[TMP35:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 51490// CHECK3-NEXT: store i32 [[TMP11]], ptr [[TMP35]], align 41491// CHECK3-NEXT: [[TMP36:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 51492// CHECK3-NEXT: store ptr null, ptr [[TMP36]], align 41493// CHECK3-NEXT: [[TMP37:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 61494// CHECK3-NEXT: store i32 [[TMP13]], ptr [[TMP37]], align 41495// CHECK3-NEXT: [[TMP38:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 61496// CHECK3-NEXT: store i32 [[TMP13]], ptr [[TMP38]], align 41497// CHECK3-NEXT: [[TMP39:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 61498// CHECK3-NEXT: store ptr null, ptr [[TMP39]], align 41499// CHECK3-NEXT: [[TMP40:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 71500// CHECK3-NEXT: store ptr @Gd, ptr [[TMP40]], align 41501// CHECK3-NEXT: [[TMP41:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 71502// CHECK3-NEXT: store ptr @Gd, ptr [[TMP41]], align 41503// CHECK3-NEXT: [[TMP42:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 71504// CHECK3-NEXT: store ptr null, ptr [[TMP42]], align 41505// CHECK3-NEXT: [[TMP43:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 81506// CHECK3-NEXT: store i32 [[TMP15]], ptr [[TMP43]], align 41507// CHECK3-NEXT: [[TMP44:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 81508// CHECK3-NEXT: store i32 [[TMP15]], ptr [[TMP44]], align 41509// CHECK3-NEXT: [[TMP45:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 81510// CHECK3-NEXT: store ptr null, ptr [[TMP45]], align 41511// CHECK3-NEXT: [[TMP46:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 01512// CHECK3-NEXT: [[TMP47:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 01513// CHECK3-NEXT: [[TMP48:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 01514// CHECK3-NEXT: store i32 3, ptr [[TMP48]], align 41515// CHECK3-NEXT: [[TMP49:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11516// CHECK3-NEXT: store i32 9, ptr [[TMP49]], align 41517// CHECK3-NEXT: [[TMP50:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 21518// CHECK3-NEXT: store ptr [[TMP46]], ptr [[TMP50]], align 41519// CHECK3-NEXT: [[TMP51:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 31520// CHECK3-NEXT: store ptr [[TMP47]], ptr [[TMP51]], align 41521// CHECK3-NEXT: [[TMP52:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 41522// CHECK3-NEXT: store ptr @.offload_sizes.1, ptr [[TMP52]], align 41523// CHECK3-NEXT: [[TMP53:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 51524// CHECK3-NEXT: store ptr @.offload_maptypes.2, ptr [[TMP53]], align 41525// CHECK3-NEXT: [[TMP54:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 61526// CHECK3-NEXT: store ptr null, ptr [[TMP54]], align 41527// CHECK3-NEXT: [[TMP55:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 71528// CHECK3-NEXT: store ptr null, ptr [[TMP55]], align 41529// CHECK3-NEXT: [[TMP56:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 81530// CHECK3-NEXT: store i64 0, ptr [[TMP56]], align 81531// CHECK3-NEXT: [[TMP57:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 91532// CHECK3-NEXT: store i64 0, ptr [[TMP57]], align 81533// CHECK3-NEXT: [[TMP58:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 101534// CHECK3-NEXT: store [3 x i32] [i32 -1, i32 0, i32 0], ptr [[TMP58]], align 41535// CHECK3-NEXT: [[TMP59:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 111536// CHECK3-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP59]], align 41537// CHECK3-NEXT: [[TMP60:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 121538// CHECK3-NEXT: store i32 0, ptr [[TMP60]], align 41539// CHECK3-NEXT: [[TMP61:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 -1, i32 -1, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3barssss_l94.region_id, ptr [[KERNEL_ARGS]])1540// CHECK3-NEXT: [[TMP62:%.*]] = icmp ne i32 [[TMP61]], 01541// CHECK3-NEXT: br i1 [[TMP62]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]1542// CHECK3: omp_offload.failed:1543// CHECK3-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3barssss_l94(i32 [[TMP5]], ptr @Gb, i32 [[TMP7]], ptr @Gc, i32 [[TMP9]], i32 [[TMP11]], i32 [[TMP13]], ptr @Gd, i32 [[TMP15]]) #[[ATTR2]]1544// CHECK3-NEXT: br label [[OMP_OFFLOAD_CONT]]1545// CHECK3: omp_offload.cont:1546// CHECK3-NEXT: br label [[OMP_IF_END:%.*]]1547// CHECK3: omp_if.else:1548// CHECK3-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3barssss_l94(i32 [[TMP5]], ptr @Gb, i32 [[TMP7]], ptr @Gc, i32 [[TMP9]], i32 [[TMP11]], i32 [[TMP13]], ptr @Gd, i32 [[TMP15]]) #[[ATTR2]]1549// CHECK3-NEXT: br label [[OMP_IF_END]]1550// CHECK3: omp_if.end:1551// CHECK3-NEXT: ret void1552//1553//1554// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3barssss_l941555// CHECK3-SAME: (i32 noundef [[B:%.*]], ptr noundef nonnull align 4 dereferenceable(8) [[GB:%.*]], i32 noundef [[SB:%.*]], ptr noundef nonnull align 4 dereferenceable(8) [[GC:%.*]], i32 noundef [[C:%.*]], i32 noundef [[SC:%.*]], i32 noundef [[D:%.*]], ptr noundef nonnull align 4 dereferenceable(8) [[GD:%.*]], i32 noundef [[SD:%.*]]) #[[ATTR1]] {1556// CHECK3-NEXT: entry:1557// CHECK3-NEXT: [[B_ADDR:%.*]] = alloca i32, align 41558// CHECK3-NEXT: [[GB_ADDR:%.*]] = alloca ptr, align 41559// CHECK3-NEXT: [[SB_ADDR:%.*]] = alloca i32, align 41560// CHECK3-NEXT: [[GC_ADDR:%.*]] = alloca ptr, align 41561// CHECK3-NEXT: [[C_ADDR:%.*]] = alloca i32, align 41562// CHECK3-NEXT: [[SC_ADDR:%.*]] = alloca i32, align 41563// CHECK3-NEXT: [[D_ADDR:%.*]] = alloca i32, align 41564// CHECK3-NEXT: [[GD_ADDR:%.*]] = alloca ptr, align 41565// CHECK3-NEXT: [[SD_ADDR:%.*]] = alloca i32, align 41566// CHECK3-NEXT: [[GB1:%.*]] = alloca double, align 81567// CHECK3-NEXT: [[GC2:%.*]] = alloca double, align 81568// CHECK3-NEXT: [[GD3:%.*]] = alloca double, align 81569// CHECK3-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 41570// CHECK3-NEXT: [[DOTBOUND_ZERO_ADDR:%.*]] = alloca i32, align 41571// CHECK3-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])1572// CHECK3-NEXT: store i32 [[B]], ptr [[B_ADDR]], align 41573// CHECK3-NEXT: store ptr [[GB]], ptr [[GB_ADDR]], align 41574// CHECK3-NEXT: store i32 [[SB]], ptr [[SB_ADDR]], align 41575// CHECK3-NEXT: store ptr [[GC]], ptr [[GC_ADDR]], align 41576// CHECK3-NEXT: store i32 [[C]], ptr [[C_ADDR]], align 41577// CHECK3-NEXT: store i32 [[SC]], ptr [[SC_ADDR]], align 41578// CHECK3-NEXT: store i32 [[D]], ptr [[D_ADDR]], align 41579// CHECK3-NEXT: store ptr [[GD]], ptr [[GD_ADDR]], align 41580// CHECK3-NEXT: store i32 [[SD]], ptr [[SD_ADDR]], align 41581// CHECK3-NEXT: [[TMP1:%.*]] = load ptr, ptr [[GB_ADDR]], align 41582// CHECK3-NEXT: [[TMP2:%.*]] = load ptr, ptr [[GC_ADDR]], align 41583// CHECK3-NEXT: [[TMP3:%.*]] = load ptr, ptr [[GD_ADDR]], align 41584// CHECK3-NEXT: [[TMP4:%.*]] = load double, ptr [[TMP1]], align 81585// CHECK3-NEXT: store double [[TMP4]], ptr [[GB1]], align 81586// CHECK3-NEXT: [[TMP5:%.*]] = load double, ptr [[TMP2]], align 81587// CHECK3-NEXT: store double [[TMP5]], ptr [[GC2]], align 81588// CHECK3-NEXT: [[TMP6:%.*]] = load double, ptr [[TMP3]], align 81589// CHECK3-NEXT: store double [[TMP6]], ptr [[GD3]], align 81590// CHECK3-NEXT: [[TMP7:%.*]] = load i16, ptr [[B_ADDR]], align 21591// CHECK3-NEXT: [[CONV:%.*]] = sext i16 [[TMP7]] to i321592// CHECK3-NEXT: [[ADD:%.*]] = add nsw i32 [[CONV]], 11593// CHECK3-NEXT: [[CONV4:%.*]] = trunc i32 [[ADD]] to i161594// CHECK3-NEXT: store i16 [[CONV4]], ptr [[B_ADDR]], align 21595// CHECK3-NEXT: [[TMP8:%.*]] = load double, ptr [[GB1]], align 81596// CHECK3-NEXT: [[ADD5:%.*]] = fadd double [[TMP8]], 1.000000e+001597// CHECK3-NEXT: store double [[ADD5]], ptr [[GB1]], align 81598// CHECK3-NEXT: [[TMP9:%.*]] = load float, ptr [[SB_ADDR]], align 41599// CHECK3-NEXT: [[CONV6:%.*]] = fpext float [[TMP9]] to double1600// CHECK3-NEXT: [[ADD7:%.*]] = fadd double [[CONV6]], 1.000000e+001601// CHECK3-NEXT: [[CONV8:%.*]] = fptrunc double [[ADD7]] to float1602// CHECK3-NEXT: store float [[CONV8]], ptr [[SB_ADDR]], align 41603// CHECK3-NEXT: [[TMP10:%.*]] = load double, ptr [[GC2]], align 81604// CHECK3-NEXT: [[CMP:%.*]] = fcmp ogt double [[TMP10]], 0.000000e+001605// CHECK3-NEXT: br i1 [[CMP]], label [[LAND_LHS_TRUE:%.*]], label [[OMP_IF_ELSE:%.*]]1606// CHECK3: land.lhs.true:1607// CHECK3-NEXT: [[TMP11:%.*]] = load i16, ptr [[C_ADDR]], align 21608// CHECK3-NEXT: [[CONV9:%.*]] = sext i16 [[TMP11]] to i321609// CHECK3-NEXT: [[CMP10:%.*]] = icmp sgt i32 [[CONV9]], 01610// CHECK3-NEXT: br i1 [[CMP10]], label [[LAND_LHS_TRUE11:%.*]], label [[OMP_IF_ELSE]]1611// CHECK3: land.lhs.true11:1612// CHECK3-NEXT: [[TMP12:%.*]] = load float, ptr [[SC_ADDR]], align 41613// CHECK3-NEXT: [[CONV12:%.*]] = fpext float [[TMP12]] to double1614// CHECK3-NEXT: [[CMP13:%.*]] = fcmp ogt double [[CONV12]], 0.000000e+001615// CHECK3-NEXT: br i1 [[CMP13]], label [[OMP_IF_THEN:%.*]], label [[OMP_IF_ELSE]]1616// CHECK3: omp_if.then:1617// CHECK3-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB1]], i32 3, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3barssss_l94.omp_outlined, ptr [[D_ADDR]], ptr [[GD3]], ptr [[SD_ADDR]])1618// CHECK3-NEXT: br label [[OMP_IF_END:%.*]]1619// CHECK3: omp_if.else:1620// CHECK3-NEXT: call void @__kmpc_serialized_parallel(ptr @[[GLOB1]], i32 [[TMP0]])1621// CHECK3-NEXT: store i32 [[TMP0]], ptr [[DOTTHREADID_TEMP_]], align 41622// CHECK3-NEXT: store i32 0, ptr [[DOTBOUND_ZERO_ADDR]], align 41623// CHECK3-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3barssss_l94.omp_outlined(ptr [[DOTTHREADID_TEMP_]], ptr [[DOTBOUND_ZERO_ADDR]], ptr [[D_ADDR]], ptr [[GD3]], ptr [[SD_ADDR]]) #[[ATTR2]]1624// CHECK3-NEXT: call void @__kmpc_end_serialized_parallel(ptr @[[GLOB1]], i32 [[TMP0]])1625// CHECK3-NEXT: br label [[OMP_IF_END]]1626// CHECK3: omp_if.end:1627// CHECK3-NEXT: ret void1628//1629//1630// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3barssss_l94.omp_outlined1631// CHECK3-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 2 dereferenceable(2) [[D:%.*]], ptr noundef nonnull align 4 dereferenceable(8) [[GD:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[SD:%.*]]) #[[ATTR1]] {1632// CHECK3-NEXT: entry:1633// CHECK3-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 41634// CHECK3-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 41635// CHECK3-NEXT: [[D_ADDR:%.*]] = alloca ptr, align 41636// CHECK3-NEXT: [[GD_ADDR:%.*]] = alloca ptr, align 41637// CHECK3-NEXT: [[SD_ADDR:%.*]] = alloca ptr, align 41638// CHECK3-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 41639// CHECK3-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 41640// CHECK3-NEXT: store ptr [[D]], ptr [[D_ADDR]], align 41641// CHECK3-NEXT: store ptr [[GD]], ptr [[GD_ADDR]], align 41642// CHECK3-NEXT: store ptr [[SD]], ptr [[SD_ADDR]], align 41643// CHECK3-NEXT: [[TMP0:%.*]] = load ptr, ptr [[D_ADDR]], align 41644// CHECK3-NEXT: [[TMP1:%.*]] = load ptr, ptr [[GD_ADDR]], align 41645// CHECK3-NEXT: [[TMP2:%.*]] = load ptr, ptr [[SD_ADDR]], align 41646// CHECK3-NEXT: [[TMP3:%.*]] = load i16, ptr [[TMP0]], align 21647// CHECK3-NEXT: [[CONV:%.*]] = sext i16 [[TMP3]] to i321648// CHECK3-NEXT: [[ADD:%.*]] = add nsw i32 [[CONV]], 11649// CHECK3-NEXT: [[CONV1:%.*]] = trunc i32 [[ADD]] to i161650// CHECK3-NEXT: store i16 [[CONV1]], ptr [[TMP0]], align 21651// CHECK3-NEXT: [[TMP4:%.*]] = load double, ptr [[TMP1]], align 81652// CHECK3-NEXT: [[ADD2:%.*]] = fadd double [[TMP4]], 1.000000e+001653// CHECK3-NEXT: store double [[ADD2]], ptr [[TMP1]], align 81654// CHECK3-NEXT: [[TMP5:%.*]] = load float, ptr [[TMP2]], align 41655// CHECK3-NEXT: [[CONV3:%.*]] = fpext float [[TMP5]] to double1656// CHECK3-NEXT: [[ADD4:%.*]] = fadd double [[CONV3]], 1.000000e+001657// CHECK3-NEXT: [[CONV5:%.*]] = fptrunc double [[ADD4]] to float1658// CHECK3-NEXT: store float [[CONV5]], ptr [[TMP2]], align 41659// CHECK3-NEXT: ret void1660//1661//1662// CHECK3-LABEL: define {{[^@]+}}@_Z5tbar2ssss1663// CHECK3-SAME: (i16 noundef signext [[A:%.*]], i16 noundef signext [[B:%.*]], i16 noundef signext [[C:%.*]], i16 noundef signext [[D:%.*]]) #[[ATTR0]] {1664// CHECK3-NEXT: entry:1665// CHECK3-NEXT: [[A_ADDR:%.*]] = alloca i16, align 21666// CHECK3-NEXT: [[B_ADDR:%.*]] = alloca i16, align 21667// CHECK3-NEXT: [[C_ADDR:%.*]] = alloca i16, align 21668// CHECK3-NEXT: [[D_ADDR:%.*]] = alloca i16, align 21669// CHECK3-NEXT: store i16 [[A]], ptr [[A_ADDR]], align 21670// CHECK3-NEXT: store i16 [[B]], ptr [[B_ADDR]], align 21671// CHECK3-NEXT: store i16 [[C]], ptr [[C_ADDR]], align 21672// CHECK3-NEXT: store i16 [[D]], ptr [[D_ADDR]], align 21673// CHECK3-NEXT: [[TMP0:%.*]] = load i16, ptr [[A_ADDR]], align 21674// CHECK3-NEXT: [[TMP1:%.*]] = load i16, ptr [[B_ADDR]], align 21675// CHECK3-NEXT: [[TMP2:%.*]] = load i16, ptr [[C_ADDR]], align 21676// CHECK3-NEXT: [[TMP3:%.*]] = load i16, ptr [[D_ADDR]], align 21677// CHECK3-NEXT: [[CALL:%.*]] = call noundef i32 @_Z4tbarIsEiT_S0_S0_S0_(i16 noundef signext [[TMP0]], i16 noundef signext [[TMP1]], i16 noundef signext [[TMP2]], i16 noundef signext [[TMP3]])1678// CHECK3-NEXT: ret i32 [[CALL]]1679//1680//1681// CHECK3-LABEL: define {{[^@]+}}@_Z4tbarIsEiT_S0_S0_S0_1682// CHECK3-SAME: (i16 noundef signext [[A:%.*]], i16 noundef signext [[B:%.*]], i16 noundef signext [[C:%.*]], i16 noundef signext [[D:%.*]]) #[[ATTR0]] comdat {1683// CHECK3-NEXT: entry:1684// CHECK3-NEXT: [[A_ADDR:%.*]] = alloca i16, align 21685// CHECK3-NEXT: [[B_ADDR:%.*]] = alloca i16, align 21686// CHECK3-NEXT: [[C_ADDR:%.*]] = alloca i16, align 21687// CHECK3-NEXT: [[D_ADDR:%.*]] = alloca i16, align 21688// CHECK3-NEXT: store i16 [[A]], ptr [[A_ADDR]], align 21689// CHECK3-NEXT: store i16 [[B]], ptr [[B_ADDR]], align 21690// CHECK3-NEXT: store i16 [[C]], ptr [[C_ADDR]], align 21691// CHECK3-NEXT: store i16 [[D]], ptr [[D_ADDR]], align 21692// CHECK3-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB1]], i32 4, ptr @_Z4tbarIsEiT_S0_S0_S0_.omp_outlined, ptr [[A_ADDR]], ptr [[B_ADDR]], ptr [[C_ADDR]], ptr [[D_ADDR]])1693// CHECK3-NEXT: [[TMP0:%.*]] = load i16, ptr [[A_ADDR]], align 21694// CHECK3-NEXT: [[CONV:%.*]] = sext i16 [[TMP0]] to i321695// CHECK3-NEXT: [[TMP1:%.*]] = load i16, ptr [[B_ADDR]], align 21696// CHECK3-NEXT: [[CONV1:%.*]] = sext i16 [[TMP1]] to i321697// CHECK3-NEXT: [[ADD:%.*]] = add nsw i32 [[CONV]], [[CONV1]]1698// CHECK3-NEXT: [[TMP2:%.*]] = load i16, ptr [[C_ADDR]], align 21699// CHECK3-NEXT: [[CONV2:%.*]] = sext i16 [[TMP2]] to i321700// CHECK3-NEXT: [[ADD3:%.*]] = add nsw i32 [[ADD]], [[CONV2]]1701// CHECK3-NEXT: [[TMP3:%.*]] = load i16, ptr [[D_ADDR]], align 21702// CHECK3-NEXT: [[CONV4:%.*]] = sext i16 [[TMP3]] to i321703// CHECK3-NEXT: [[ADD5:%.*]] = add nsw i32 [[ADD3]], [[CONV4]]1704// CHECK3-NEXT: [[TMP4:%.*]] = load float, ptr @_ZZ4tbarIsEiT_S0_S0_S0_E2Sa, align 41705// CHECK3-NEXT: [[CONV6:%.*]] = fptosi float [[TMP4]] to i321706// CHECK3-NEXT: [[ADD7:%.*]] = add nsw i32 [[ADD5]], [[CONV6]]1707// CHECK3-NEXT: [[TMP5:%.*]] = load float, ptr @_ZZ4tbarIsEiT_S0_S0_S0_E2Sb, align 41708// CHECK3-NEXT: [[CONV8:%.*]] = fptosi float [[TMP5]] to i321709// CHECK3-NEXT: [[ADD9:%.*]] = add nsw i32 [[ADD7]], [[CONV8]]1710// CHECK3-NEXT: [[TMP6:%.*]] = load float, ptr @_ZZ4tbarIsEiT_S0_S0_S0_E2Sc, align 41711// CHECK3-NEXT: [[CONV10:%.*]] = fptosi float [[TMP6]] to i321712// CHECK3-NEXT: [[ADD11:%.*]] = add nsw i32 [[ADD9]], [[CONV10]]1713// CHECK3-NEXT: [[TMP7:%.*]] = load float, ptr @_ZZ4tbarIsEiT_S0_S0_S0_E2Sd, align 41714// CHECK3-NEXT: [[CONV12:%.*]] = fptosi float [[TMP7]] to i321715// CHECK3-NEXT: [[ADD13:%.*]] = add nsw i32 [[ADD11]], [[CONV12]]1716// CHECK3-NEXT: ret i32 [[ADD13]]1717//1718//1719// CHECK3-LABEL: define {{[^@]+}}@_Z4tbarIsEiT_S0_S0_S0_.omp_outlined1720// CHECK3-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 2 dereferenceable(2) [[A:%.*]], ptr noundef nonnull align 2 dereferenceable(2) [[B:%.*]], ptr noundef nonnull align 2 dereferenceable(2) [[C:%.*]], ptr noundef nonnull align 2 dereferenceable(2) [[D:%.*]]) #[[ATTR1]] {1721// CHECK3-NEXT: entry:1722// CHECK3-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 41723// CHECK3-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 41724// CHECK3-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 41725// CHECK3-NEXT: [[B_ADDR:%.*]] = alloca ptr, align 41726// CHECK3-NEXT: [[C_ADDR:%.*]] = alloca ptr, align 41727// CHECK3-NEXT: [[D_ADDR:%.*]] = alloca ptr, align 41728// CHECK3-NEXT: [[B_CASTED:%.*]] = alloca i32, align 41729// CHECK3-NEXT: [[SB_CASTED:%.*]] = alloca i32, align 41730// CHECK3-NEXT: [[C_CASTED:%.*]] = alloca i32, align 41731// CHECK3-NEXT: [[SC_CASTED:%.*]] = alloca i32, align 41732// CHECK3-NEXT: [[D_CASTED:%.*]] = alloca i32, align 41733// CHECK3-NEXT: [[SD_CASTED:%.*]] = alloca i32, align 41734// CHECK3-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [9 x ptr], align 41735// CHECK3-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [9 x ptr], align 41736// CHECK3-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [9 x ptr], align 41737// CHECK3-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 81738// CHECK3-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 41739// CHECK3-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 41740// CHECK3-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 41741// CHECK3-NEXT: store ptr [[B]], ptr [[B_ADDR]], align 41742// CHECK3-NEXT: store ptr [[C]], ptr [[C_ADDR]], align 41743// CHECK3-NEXT: store ptr [[D]], ptr [[D_ADDR]], align 41744// CHECK3-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 41745// CHECK3-NEXT: [[TMP1:%.*]] = load ptr, ptr [[B_ADDR]], align 41746// CHECK3-NEXT: [[TMP2:%.*]] = load ptr, ptr [[C_ADDR]], align 41747// CHECK3-NEXT: [[TMP3:%.*]] = load ptr, ptr [[D_ADDR]], align 41748// CHECK3-NEXT: [[TMP4:%.*]] = load i16, ptr [[TMP1]], align 21749// CHECK3-NEXT: store i16 [[TMP4]], ptr [[B_CASTED]], align 21750// CHECK3-NEXT: [[TMP5:%.*]] = load i32, ptr [[B_CASTED]], align 41751// CHECK3-NEXT: [[TMP6:%.*]] = load float, ptr @_ZZ4tbarIsEiT_S0_S0_S0_E2Sb, align 41752// CHECK3-NEXT: store float [[TMP6]], ptr [[SB_CASTED]], align 41753// CHECK3-NEXT: [[TMP7:%.*]] = load i32, ptr [[SB_CASTED]], align 41754// CHECK3-NEXT: [[TMP8:%.*]] = load i16, ptr [[TMP2]], align 21755// CHECK3-NEXT: store i16 [[TMP8]], ptr [[C_CASTED]], align 21756// CHECK3-NEXT: [[TMP9:%.*]] = load i32, ptr [[C_CASTED]], align 41757// CHECK3-NEXT: [[TMP10:%.*]] = load float, ptr @_ZZ4tbarIsEiT_S0_S0_S0_E2Sc, align 41758// CHECK3-NEXT: store float [[TMP10]], ptr [[SC_CASTED]], align 41759// CHECK3-NEXT: [[TMP11:%.*]] = load i32, ptr [[SC_CASTED]], align 41760// CHECK3-NEXT: [[TMP12:%.*]] = load i16, ptr [[TMP3]], align 21761// CHECK3-NEXT: store i16 [[TMP12]], ptr [[D_CASTED]], align 21762// CHECK3-NEXT: [[TMP13:%.*]] = load i32, ptr [[D_CASTED]], align 41763// CHECK3-NEXT: [[TMP14:%.*]] = load float, ptr @_ZZ4tbarIsEiT_S0_S0_S0_E2Sd, align 41764// CHECK3-NEXT: store float [[TMP14]], ptr [[SD_CASTED]], align 41765// CHECK3-NEXT: [[TMP15:%.*]] = load i32, ptr [[SD_CASTED]], align 41766// CHECK3-NEXT: [[TMP16:%.*]] = load double, ptr @Ga, align 81767// CHECK3-NEXT: [[CMP:%.*]] = fcmp ogt double [[TMP16]], 0.000000e+001768// CHECK3-NEXT: br i1 [[CMP]], label [[LAND_LHS_TRUE:%.*]], label [[OMP_IF_ELSE:%.*]]1769// CHECK3: land.lhs.true:1770// CHECK3-NEXT: [[TMP17:%.*]] = load i16, ptr [[TMP0]], align 21771// CHECK3-NEXT: [[CONV:%.*]] = sext i16 [[TMP17]] to i321772// CHECK3-NEXT: [[CMP1:%.*]] = icmp sgt i32 [[CONV]], 01773// CHECK3-NEXT: br i1 [[CMP1]], label [[LAND_LHS_TRUE2:%.*]], label [[OMP_IF_ELSE]]1774// CHECK3: land.lhs.true2:1775// CHECK3-NEXT: [[TMP18:%.*]] = load float, ptr @_ZZ4tbarIsEiT_S0_S0_S0_E2Sa, align 41776// CHECK3-NEXT: [[CONV3:%.*]] = fpext float [[TMP18]] to double1777// CHECK3-NEXT: [[CMP4:%.*]] = fcmp ogt double [[CONV3]], 0.000000e+001778// CHECK3-NEXT: br i1 [[CMP4]], label [[OMP_IF_THEN:%.*]], label [[OMP_IF_ELSE]]1779// CHECK3: omp_if.then:1780// CHECK3-NEXT: [[TMP19:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 01781// CHECK3-NEXT: store i32 [[TMP5]], ptr [[TMP19]], align 41782// CHECK3-NEXT: [[TMP20:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 01783// CHECK3-NEXT: store i32 [[TMP5]], ptr [[TMP20]], align 41784// CHECK3-NEXT: [[TMP21:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 01785// CHECK3-NEXT: store ptr null, ptr [[TMP21]], align 41786// CHECK3-NEXT: [[TMP22:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 11787// CHECK3-NEXT: store ptr @Gb, ptr [[TMP22]], align 41788// CHECK3-NEXT: [[TMP23:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 11789// CHECK3-NEXT: store ptr @Gb, ptr [[TMP23]], align 41790// CHECK3-NEXT: [[TMP24:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 11791// CHECK3-NEXT: store ptr null, ptr [[TMP24]], align 41792// CHECK3-NEXT: [[TMP25:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 21793// CHECK3-NEXT: store i32 [[TMP7]], ptr [[TMP25]], align 41794// CHECK3-NEXT: [[TMP26:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 21795// CHECK3-NEXT: store i32 [[TMP7]], ptr [[TMP26]], align 41796// CHECK3-NEXT: [[TMP27:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 21797// CHECK3-NEXT: store ptr null, ptr [[TMP27]], align 41798// CHECK3-NEXT: [[TMP28:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 31799// CHECK3-NEXT: store ptr @Gc, ptr [[TMP28]], align 41800// CHECK3-NEXT: [[TMP29:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 31801// CHECK3-NEXT: store ptr @Gc, ptr [[TMP29]], align 41802// CHECK3-NEXT: [[TMP30:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 31803// CHECK3-NEXT: store ptr null, ptr [[TMP30]], align 41804// CHECK3-NEXT: [[TMP31:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 41805// CHECK3-NEXT: store i32 [[TMP9]], ptr [[TMP31]], align 41806// CHECK3-NEXT: [[TMP32:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 41807// CHECK3-NEXT: store i32 [[TMP9]], ptr [[TMP32]], align 41808// CHECK3-NEXT: [[TMP33:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 41809// CHECK3-NEXT: store ptr null, ptr [[TMP33]], align 41810// CHECK3-NEXT: [[TMP34:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 51811// CHECK3-NEXT: store i32 [[TMP11]], ptr [[TMP34]], align 41812// CHECK3-NEXT: [[TMP35:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 51813// CHECK3-NEXT: store i32 [[TMP11]], ptr [[TMP35]], align 41814// CHECK3-NEXT: [[TMP36:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 51815// CHECK3-NEXT: store ptr null, ptr [[TMP36]], align 41816// CHECK3-NEXT: [[TMP37:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 61817// CHECK3-NEXT: store i32 [[TMP13]], ptr [[TMP37]], align 41818// CHECK3-NEXT: [[TMP38:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 61819// CHECK3-NEXT: store i32 [[TMP13]], ptr [[TMP38]], align 41820// CHECK3-NEXT: [[TMP39:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 61821// CHECK3-NEXT: store ptr null, ptr [[TMP39]], align 41822// CHECK3-NEXT: [[TMP40:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 71823// CHECK3-NEXT: store ptr @Gd, ptr [[TMP40]], align 41824// CHECK3-NEXT: [[TMP41:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 71825// CHECK3-NEXT: store ptr @Gd, ptr [[TMP41]], align 41826// CHECK3-NEXT: [[TMP42:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 71827// CHECK3-NEXT: store ptr null, ptr [[TMP42]], align 41828// CHECK3-NEXT: [[TMP43:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 81829// CHECK3-NEXT: store i32 [[TMP15]], ptr [[TMP43]], align 41830// CHECK3-NEXT: [[TMP44:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 81831// CHECK3-NEXT: store i32 [[TMP15]], ptr [[TMP44]], align 41832// CHECK3-NEXT: [[TMP45:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 81833// CHECK3-NEXT: store ptr null, ptr [[TMP45]], align 41834// CHECK3-NEXT: [[TMP46:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 01835// CHECK3-NEXT: [[TMP47:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 01836// CHECK3-NEXT: [[TMP48:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 01837// CHECK3-NEXT: store i32 3, ptr [[TMP48]], align 41838// CHECK3-NEXT: [[TMP49:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11839// CHECK3-NEXT: store i32 9, ptr [[TMP49]], align 41840// CHECK3-NEXT: [[TMP50:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 21841// CHECK3-NEXT: store ptr [[TMP46]], ptr [[TMP50]], align 41842// CHECK3-NEXT: [[TMP51:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 31843// CHECK3-NEXT: store ptr [[TMP47]], ptr [[TMP51]], align 41844// CHECK3-NEXT: [[TMP52:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 41845// CHECK3-NEXT: store ptr @.offload_sizes.3, ptr [[TMP52]], align 41846// CHECK3-NEXT: [[TMP53:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 51847// CHECK3-NEXT: store ptr @.offload_maptypes.4, ptr [[TMP53]], align 41848// CHECK3-NEXT: [[TMP54:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 61849// CHECK3-NEXT: store ptr null, ptr [[TMP54]], align 41850// CHECK3-NEXT: [[TMP55:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 71851// CHECK3-NEXT: store ptr null, ptr [[TMP55]], align 41852// CHECK3-NEXT: [[TMP56:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 81853// CHECK3-NEXT: store i64 0, ptr [[TMP56]], align 81854// CHECK3-NEXT: [[TMP57:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 91855// CHECK3-NEXT: store i64 0, ptr [[TMP57]], align 81856// CHECK3-NEXT: [[TMP58:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 101857// CHECK3-NEXT: store [3 x i32] [i32 -1, i32 0, i32 0], ptr [[TMP58]], align 41858// CHECK3-NEXT: [[TMP59:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 111859// CHECK3-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP59]], align 41860// CHECK3-NEXT: [[TMP60:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 121861// CHECK3-NEXT: store i32 0, ptr [[TMP60]], align 41862// CHECK3-NEXT: [[TMP61:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 -1, i32 -1, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z4tbarIsEiT_S0_S0_S0__l145.region_id, ptr [[KERNEL_ARGS]])1863// CHECK3-NEXT: [[TMP62:%.*]] = icmp ne i32 [[TMP61]], 01864// CHECK3-NEXT: br i1 [[TMP62]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]1865// CHECK3: omp_offload.failed:1866// CHECK3-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z4tbarIsEiT_S0_S0_S0__l145(i32 [[TMP5]], ptr @Gb, i32 [[TMP7]], ptr @Gc, i32 [[TMP9]], i32 [[TMP11]], i32 [[TMP13]], ptr @Gd, i32 [[TMP15]]) #[[ATTR2]]1867// CHECK3-NEXT: br label [[OMP_OFFLOAD_CONT]]1868// CHECK3: omp_offload.cont:1869// CHECK3-NEXT: br label [[OMP_IF_END:%.*]]1870// CHECK3: omp_if.else:1871// CHECK3-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z4tbarIsEiT_S0_S0_S0__l145(i32 [[TMP5]], ptr @Gb, i32 [[TMP7]], ptr @Gc, i32 [[TMP9]], i32 [[TMP11]], i32 [[TMP13]], ptr @Gd, i32 [[TMP15]]) #[[ATTR2]]1872// CHECK3-NEXT: br label [[OMP_IF_END]]1873// CHECK3: omp_if.end:1874// CHECK3-NEXT: ret void1875//1876//1877// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z4tbarIsEiT_S0_S0_S0__l1451878// CHECK3-SAME: (i32 noundef [[B:%.*]], ptr noundef nonnull align 4 dereferenceable(8) [[GB:%.*]], i32 noundef [[SB:%.*]], ptr noundef nonnull align 4 dereferenceable(8) [[GC:%.*]], i32 noundef [[C:%.*]], i32 noundef [[SC:%.*]], i32 noundef [[D:%.*]], ptr noundef nonnull align 4 dereferenceable(8) [[GD:%.*]], i32 noundef [[SD:%.*]]) #[[ATTR1]] {1879// CHECK3-NEXT: entry:1880// CHECK3-NEXT: [[B_ADDR:%.*]] = alloca i32, align 41881// CHECK3-NEXT: [[GB_ADDR:%.*]] = alloca ptr, align 41882// CHECK3-NEXT: [[SB_ADDR:%.*]] = alloca i32, align 41883// CHECK3-NEXT: [[GC_ADDR:%.*]] = alloca ptr, align 41884// CHECK3-NEXT: [[C_ADDR:%.*]] = alloca i32, align 41885// CHECK3-NEXT: [[SC_ADDR:%.*]] = alloca i32, align 41886// CHECK3-NEXT: [[D_ADDR:%.*]] = alloca i32, align 41887// CHECK3-NEXT: [[GD_ADDR:%.*]] = alloca ptr, align 41888// CHECK3-NEXT: [[SD_ADDR:%.*]] = alloca i32, align 41889// CHECK3-NEXT: [[GB1:%.*]] = alloca double, align 81890// CHECK3-NEXT: [[GC2:%.*]] = alloca double, align 81891// CHECK3-NEXT: [[GD3:%.*]] = alloca double, align 81892// CHECK3-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 41893// CHECK3-NEXT: [[DOTBOUND_ZERO_ADDR:%.*]] = alloca i32, align 41894// CHECK3-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])1895// CHECK3-NEXT: store i32 [[B]], ptr [[B_ADDR]], align 41896// CHECK3-NEXT: store ptr [[GB]], ptr [[GB_ADDR]], align 41897// CHECK3-NEXT: store i32 [[SB]], ptr [[SB_ADDR]], align 41898// CHECK3-NEXT: store ptr [[GC]], ptr [[GC_ADDR]], align 41899// CHECK3-NEXT: store i32 [[C]], ptr [[C_ADDR]], align 41900// CHECK3-NEXT: store i32 [[SC]], ptr [[SC_ADDR]], align 41901// CHECK3-NEXT: store i32 [[D]], ptr [[D_ADDR]], align 41902// CHECK3-NEXT: store ptr [[GD]], ptr [[GD_ADDR]], align 41903// CHECK3-NEXT: store i32 [[SD]], ptr [[SD_ADDR]], align 41904// CHECK3-NEXT: [[TMP1:%.*]] = load ptr, ptr [[GB_ADDR]], align 41905// CHECK3-NEXT: [[TMP2:%.*]] = load ptr, ptr [[GC_ADDR]], align 41906// CHECK3-NEXT: [[TMP3:%.*]] = load ptr, ptr [[GD_ADDR]], align 41907// CHECK3-NEXT: [[TMP4:%.*]] = load double, ptr [[TMP1]], align 81908// CHECK3-NEXT: store double [[TMP4]], ptr [[GB1]], align 81909// CHECK3-NEXT: [[TMP5:%.*]] = load double, ptr [[TMP2]], align 81910// CHECK3-NEXT: store double [[TMP5]], ptr [[GC2]], align 81911// CHECK3-NEXT: [[TMP6:%.*]] = load double, ptr [[TMP3]], align 81912// CHECK3-NEXT: store double [[TMP6]], ptr [[GD3]], align 81913// CHECK3-NEXT: [[TMP7:%.*]] = load i16, ptr [[B_ADDR]], align 21914// CHECK3-NEXT: [[CONV:%.*]] = sext i16 [[TMP7]] to i321915// CHECK3-NEXT: [[ADD:%.*]] = add nsw i32 [[CONV]], 11916// CHECK3-NEXT: [[CONV4:%.*]] = trunc i32 [[ADD]] to i161917// CHECK3-NEXT: store i16 [[CONV4]], ptr [[B_ADDR]], align 21918// CHECK3-NEXT: [[TMP8:%.*]] = load double, ptr [[GB1]], align 81919// CHECK3-NEXT: [[ADD5:%.*]] = fadd double [[TMP8]], 1.000000e+001920// CHECK3-NEXT: store double [[ADD5]], ptr [[GB1]], align 81921// CHECK3-NEXT: [[TMP9:%.*]] = load float, ptr [[SB_ADDR]], align 41922// CHECK3-NEXT: [[CONV6:%.*]] = fpext float [[TMP9]] to double1923// CHECK3-NEXT: [[ADD7:%.*]] = fadd double [[CONV6]], 1.000000e+001924// CHECK3-NEXT: [[CONV8:%.*]] = fptrunc double [[ADD7]] to float1925// CHECK3-NEXT: store float [[CONV8]], ptr [[SB_ADDR]], align 41926// CHECK3-NEXT: [[TMP10:%.*]] = load double, ptr [[GC2]], align 81927// CHECK3-NEXT: [[CMP:%.*]] = fcmp ogt double [[TMP10]], 0.000000e+001928// CHECK3-NEXT: br i1 [[CMP]], label [[LAND_LHS_TRUE:%.*]], label [[OMP_IF_ELSE:%.*]]1929// CHECK3: land.lhs.true:1930// CHECK3-NEXT: [[TMP11:%.*]] = load i16, ptr [[C_ADDR]], align 21931// CHECK3-NEXT: [[CONV9:%.*]] = sext i16 [[TMP11]] to i321932// CHECK3-NEXT: [[CMP10:%.*]] = icmp sgt i32 [[CONV9]], 01933// CHECK3-NEXT: br i1 [[CMP10]], label [[LAND_LHS_TRUE11:%.*]], label [[OMP_IF_ELSE]]1934// CHECK3: land.lhs.true11:1935// CHECK3-NEXT: [[TMP12:%.*]] = load float, ptr [[SC_ADDR]], align 41936// CHECK3-NEXT: [[CONV12:%.*]] = fpext float [[TMP12]] to double1937// CHECK3-NEXT: [[CMP13:%.*]] = fcmp ogt double [[CONV12]], 0.000000e+001938// CHECK3-NEXT: br i1 [[CMP13]], label [[OMP_IF_THEN:%.*]], label [[OMP_IF_ELSE]]1939// CHECK3: omp_if.then:1940// CHECK3-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB1]], i32 3, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z4tbarIsEiT_S0_S0_S0__l145.omp_outlined, ptr [[D_ADDR]], ptr [[GD3]], ptr [[SD_ADDR]])1941// CHECK3-NEXT: br label [[OMP_IF_END:%.*]]1942// CHECK3: omp_if.else:1943// CHECK3-NEXT: call void @__kmpc_serialized_parallel(ptr @[[GLOB1]], i32 [[TMP0]])1944// CHECK3-NEXT: store i32 [[TMP0]], ptr [[DOTTHREADID_TEMP_]], align 41945// CHECK3-NEXT: store i32 0, ptr [[DOTBOUND_ZERO_ADDR]], align 41946// CHECK3-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z4tbarIsEiT_S0_S0_S0__l145.omp_outlined(ptr [[DOTTHREADID_TEMP_]], ptr [[DOTBOUND_ZERO_ADDR]], ptr [[D_ADDR]], ptr [[GD3]], ptr [[SD_ADDR]]) #[[ATTR2]]1947// CHECK3-NEXT: call void @__kmpc_end_serialized_parallel(ptr @[[GLOB1]], i32 [[TMP0]])1948// CHECK3-NEXT: br label [[OMP_IF_END]]1949// CHECK3: omp_if.end:1950// CHECK3-NEXT: ret void1951//1952//1953// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z4tbarIsEiT_S0_S0_S0__l145.omp_outlined1954// CHECK3-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 2 dereferenceable(2) [[D:%.*]], ptr noundef nonnull align 4 dereferenceable(8) [[GD:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[SD:%.*]]) #[[ATTR1]] {1955// CHECK3-NEXT: entry:1956// CHECK3-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 41957// CHECK3-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 41958// CHECK3-NEXT: [[D_ADDR:%.*]] = alloca ptr, align 41959// CHECK3-NEXT: [[GD_ADDR:%.*]] = alloca ptr, align 41960// CHECK3-NEXT: [[SD_ADDR:%.*]] = alloca ptr, align 41961// CHECK3-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 41962// CHECK3-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 41963// CHECK3-NEXT: store ptr [[D]], ptr [[D_ADDR]], align 41964// CHECK3-NEXT: store ptr [[GD]], ptr [[GD_ADDR]], align 41965// CHECK3-NEXT: store ptr [[SD]], ptr [[SD_ADDR]], align 41966// CHECK3-NEXT: [[TMP0:%.*]] = load ptr, ptr [[D_ADDR]], align 41967// CHECK3-NEXT: [[TMP1:%.*]] = load ptr, ptr [[GD_ADDR]], align 41968// CHECK3-NEXT: [[TMP2:%.*]] = load ptr, ptr [[SD_ADDR]], align 41969// CHECK3-NEXT: [[TMP3:%.*]] = load i16, ptr [[TMP0]], align 21970// CHECK3-NEXT: [[CONV:%.*]] = sext i16 [[TMP3]] to i321971// CHECK3-NEXT: [[ADD:%.*]] = add nsw i32 [[CONV]], 11972// CHECK3-NEXT: [[CONV1:%.*]] = trunc i32 [[ADD]] to i161973// CHECK3-NEXT: store i16 [[CONV1]], ptr [[TMP0]], align 21974// CHECK3-NEXT: [[TMP4:%.*]] = load double, ptr [[TMP1]], align 81975// CHECK3-NEXT: [[ADD2:%.*]] = fadd double [[TMP4]], 1.000000e+001976// CHECK3-NEXT: store double [[ADD2]], ptr [[TMP1]], align 81977// CHECK3-NEXT: [[TMP5:%.*]] = load float, ptr [[TMP2]], align 41978// CHECK3-NEXT: [[CONV3:%.*]] = fpext float [[TMP5]] to double1979// CHECK3-NEXT: [[ADD4:%.*]] = fadd double [[CONV3]], 1.000000e+001980// CHECK3-NEXT: [[CONV5:%.*]] = fptrunc double [[ADD4]] to float1981// CHECK3-NEXT: store float [[CONV5]], ptr [[TMP2]], align 41982// CHECK3-NEXT: ret void1983//1984