1976 lines · cpp
1// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --function-signature --include-generated-funcs --replace-value-regex "__omp_offloading_[0-9a-z]+_[0-9a-z]+" "reduction_size[.].+[.]" "pl_cond[.].+[.|,]" --prefix-filecheck-ir-name _2// Test host codegen.3// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix=CHECK14// RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s5// RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -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=CHECK16// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix=CHECK37// RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s8// RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -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=CHECK39 10// RUN: %clang_cc1 -verify -fopenmp-simd -fopenmp-version=45 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"11// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=45 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s12// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=45 -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}}"13// RUN: %clang_cc1 -verify -fopenmp-simd -fopenmp-version=45 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"14// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=45 -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s15// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=45 -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}}"16 17// Test target codegen - host bc file has to be created first.18// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm-bc %s -o %t-ppc-host.bc19// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --check-prefix=CHECK920// RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o %t %s21// RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix=CHECK922// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm-bc %s -o %t-x86-host.bc23// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s --check-prefix=CHECK1124// RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-x86-host.bc -o %t %s25// RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-x86-host.bc -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix=CHECK1126 27// RUN: %clang_cc1 -verify -fopenmp-simd -fopenmp-version=45 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm-bc %s -o %t-ppc-host.bc28// RUN: %clang_cc1 -verify -fopenmp-simd -fopenmp-version=45 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"29// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=45 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o %t %s30// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=45 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"31// RUN: %clang_cc1 -verify -fopenmp-simd -fopenmp-version=45 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm-bc %s -o %t-x86-host.bc32// RUN: %clang_cc1 -verify -fopenmp-simd -fopenmp-version=45 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"33// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=45 -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-x86-host.bc -o %t %s34// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=45 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-x86-host.bc -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"35 36// 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=CHECK137// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s38// 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=CHECK139// 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=CHECK340// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s41// 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=CHECK342 43// 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}}"44// RUN: %clang_cc1 -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s45// 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}}"46// 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}}"47// RUN: %clang_cc1 -fopenmp-simd -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s48// 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}}"49 50// Test target codegen - host bc file has to be created first.51// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm-bc %s -o %t-ppc-host.bc52// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --check-prefix=CHECK953// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o %t %s54// RUN: %clang_cc1 -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix=CHECK955// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm-bc %s -o %t-x86-host.bc56// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s --check-prefix=CHECK1157// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-x86-host.bc -o %t %s58// RUN: %clang_cc1 -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-x86-host.bc -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix=CHECK1159 60// RUN: %clang_cc1 -verify -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm-bc %s -o %t-ppc-host.bc61// RUN: %clang_cc1 -verify -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"62// RUN: %clang_cc1 -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o %t %s63// RUN: %clang_cc1 -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"64// RUN: %clang_cc1 -verify -fopenmp-simd -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm-bc %s -o %t-x86-host.bc65// RUN: %clang_cc1 -verify -fopenmp-simd -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"66// RUN: %clang_cc1 -fopenmp-simd -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-x86-host.bc -o %t %s67// RUN: %clang_cc1 -fopenmp-simd -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-x86-host.bc -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"68 69// expected-no-diagnostics70#ifndef HEADER71#define HEADER72 73 74 75 76// We have 6 target regions77 78 79 80// Check target registration is registered as a Ctor.81 82 83template<typename tx>84tx ftemplate(int n) {85 tx a = 0;86 87 #pragma omp target parallel if(parallel: 0)88 {89 a += 1;90 }91 92 short b = 1;93 #pragma omp target parallel if(parallel: 1)94 {95 a += b;96 }97 98 return a;99}100 101static102int fstatic(int n) {103 104 #pragma omp target parallel if(n>1)105 {106 }107 108 #pragma omp target parallel if(target: n-2>2)109 {110 }111 112 return n+1;113}114 115struct S1 {116 double a;117 118 int r1(int n){119 int b = 1;120 121 #pragma omp target parallel if(parallel: n>3)122 {123 this->a = (double)b + 1.5;124 }125 126 #pragma omp target parallel if(target: n>4) if(parallel: n>5)127 {128 this->a = 2.5;129 }130 131 return (int)a;132 }133};134 135int bar(int n){136 int a = 0;137 138 S1 S;139 a += S.r1(n);140 141 a += fstatic(n);142 143 a += ftemplate<int>(n);144 145 return a;146}147 148 149 150 151 152 153 154 155 156 157 158 159 160 161// Check that the offloading functions are emitted and that the parallel function162// is appropriately guarded.163 164 165 166 167 168 169 170 171 172 173 174 175 176 177 178 179 180 181 182 183 184#endif185// CHECK1-LABEL: define {{[^@]+}}@_Z3bari186// CHECK1-SAME: (i32 noundef signext [[N:%.*]]) #[[ATTR0:[0-9]+]] {187// CHECK1-NEXT: entry:188// CHECK1-NEXT: [[N_ADDR:%.*]] = alloca i32, align 4189// CHECK1-NEXT: [[A:%.*]] = alloca i32, align 4190// CHECK1-NEXT: [[S:%.*]] = alloca [[STRUCT_S1:%.*]], align 8191// CHECK1-NEXT: store i32 [[N]], ptr [[N_ADDR]], align 4192// CHECK1-NEXT: store i32 0, ptr [[A]], align 4193// CHECK1-NEXT: [[TMP0:%.*]] = load i32, ptr [[N_ADDR]], align 4194// CHECK1-NEXT: [[CALL:%.*]] = call noundef signext i32 @_ZN2S12r1Ei(ptr noundef nonnull align 8 dereferenceable(8) [[S]], i32 noundef signext [[TMP0]])195// CHECK1-NEXT: [[TMP1:%.*]] = load i32, ptr [[A]], align 4196// CHECK1-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP1]], [[CALL]]197// CHECK1-NEXT: store i32 [[ADD]], ptr [[A]], align 4198// CHECK1-NEXT: [[TMP2:%.*]] = load i32, ptr [[N_ADDR]], align 4199// CHECK1-NEXT: [[CALL1:%.*]] = call noundef signext i32 @_ZL7fstatici(i32 noundef signext [[TMP2]])200// CHECK1-NEXT: [[TMP3:%.*]] = load i32, ptr [[A]], align 4201// CHECK1-NEXT: [[ADD2:%.*]] = add nsw i32 [[TMP3]], [[CALL1]]202// CHECK1-NEXT: store i32 [[ADD2]], ptr [[A]], align 4203// CHECK1-NEXT: [[TMP4:%.*]] = load i32, ptr [[N_ADDR]], align 4204// CHECK1-NEXT: [[CALL3:%.*]] = call noundef signext i32 @_Z9ftemplateIiET_i(i32 noundef signext [[TMP4]])205// CHECK1-NEXT: [[TMP5:%.*]] = load i32, ptr [[A]], align 4206// CHECK1-NEXT: [[ADD4:%.*]] = add nsw i32 [[TMP5]], [[CALL3]]207// CHECK1-NEXT: store i32 [[ADD4]], ptr [[A]], align 4208// CHECK1-NEXT: [[TMP6:%.*]] = load i32, ptr [[A]], align 4209// CHECK1-NEXT: ret i32 [[TMP6]]210//211//212// CHECK1-LABEL: define {{[^@]+}}@_ZN2S12r1Ei213// CHECK1-SAME: (ptr noundef nonnull align 8 dereferenceable(8) [[THIS:%.*]], i32 noundef signext [[N:%.*]]) #[[ATTR0]] comdat {214// CHECK1-NEXT: entry:215// CHECK1-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 8216// CHECK1-NEXT: [[N_ADDR:%.*]] = alloca i32, align 4217// CHECK1-NEXT: [[B:%.*]] = alloca i32, align 4218// CHECK1-NEXT: [[DOTCAPTURE_EXPR_:%.*]] = alloca i8, align 1219// CHECK1-NEXT: [[B_CASTED:%.*]] = alloca i64, align 8220// CHECK1-NEXT: [[DOTCAPTURE_EXPR__CASTED:%.*]] = alloca i64, align 8221// CHECK1-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [3 x ptr], align 8222// CHECK1-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [3 x ptr], align 8223// CHECK1-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [3 x ptr], align 8224// CHECK1-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8225// CHECK1-NEXT: [[DOTCAPTURE_EXPR_4:%.*]] = alloca i8, align 1226// CHECK1-NEXT: [[DOTCAPTURE_EXPR__CASTED8:%.*]] = alloca i64, align 8227// CHECK1-NEXT: [[DOTOFFLOAD_BASEPTRS12:%.*]] = alloca [2 x ptr], align 8228// CHECK1-NEXT: [[DOTOFFLOAD_PTRS13:%.*]] = alloca [2 x ptr], align 8229// CHECK1-NEXT: [[DOTOFFLOAD_MAPPERS14:%.*]] = alloca [2 x ptr], align 8230// CHECK1-NEXT: [[KERNEL_ARGS16:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 8231// CHECK1-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 8232// CHECK1-NEXT: store i32 [[N]], ptr [[N_ADDR]], align 4233// CHECK1-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8234// CHECK1-NEXT: store i32 1, ptr [[B]], align 4235// CHECK1-NEXT: [[TMP0:%.*]] = load i32, ptr [[N_ADDR]], align 4236// CHECK1-NEXT: [[CMP:%.*]] = icmp sgt i32 [[TMP0]], 3237// CHECK1-NEXT: [[STOREDV:%.*]] = zext i1 [[CMP]] to i8238// CHECK1-NEXT: store i8 [[STOREDV]], ptr [[DOTCAPTURE_EXPR_]], align 1239// CHECK1-NEXT: [[TMP1:%.*]] = load i32, ptr [[B]], align 4240// CHECK1-NEXT: store i32 [[TMP1]], ptr [[B_CASTED]], align 4241// CHECK1-NEXT: [[TMP2:%.*]] = load i64, ptr [[B_CASTED]], align 8242// CHECK1-NEXT: [[TMP3:%.*]] = load i8, ptr [[DOTCAPTURE_EXPR_]], align 1243// CHECK1-NEXT: [[LOADEDV:%.*]] = trunc i8 [[TMP3]] to i1244// CHECK1-NEXT: [[STOREDV2:%.*]] = zext i1 [[LOADEDV]] to i8245// CHECK1-NEXT: store i8 [[STOREDV2]], ptr [[DOTCAPTURE_EXPR__CASTED]], align 1246// CHECK1-NEXT: [[TMP4:%.*]] = load i64, ptr [[DOTCAPTURE_EXPR__CASTED]], align 8247// CHECK1-NEXT: [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_S1:%.*]], ptr [[THIS1]], i32 0, i32 0248// CHECK1-NEXT: [[TMP5:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0249// CHECK1-NEXT: store ptr [[THIS1]], ptr [[TMP5]], align 8250// CHECK1-NEXT: [[TMP6:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0251// CHECK1-NEXT: store ptr [[A]], ptr [[TMP6]], align 8252// CHECK1-NEXT: [[TMP7:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 0253// CHECK1-NEXT: store ptr null, ptr [[TMP7]], align 8254// CHECK1-NEXT: [[TMP8:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 1255// CHECK1-NEXT: store i64 [[TMP2]], ptr [[TMP8]], align 8256// CHECK1-NEXT: [[TMP9:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 1257// CHECK1-NEXT: store i64 [[TMP2]], ptr [[TMP9]], align 8258// CHECK1-NEXT: [[TMP10:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 1259// CHECK1-NEXT: store ptr null, ptr [[TMP10]], align 8260// CHECK1-NEXT: [[TMP11:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 2261// CHECK1-NEXT: store i64 [[TMP4]], ptr [[TMP11]], align 8262// CHECK1-NEXT: [[TMP12:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 2263// CHECK1-NEXT: store i64 [[TMP4]], ptr [[TMP12]], align 8264// CHECK1-NEXT: [[TMP13:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 2265// CHECK1-NEXT: store ptr null, ptr [[TMP13]], align 8266// CHECK1-NEXT: [[TMP14:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0267// CHECK1-NEXT: [[TMP15:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0268// CHECK1-NEXT: [[TMP16:%.*]] = load i8, ptr [[DOTCAPTURE_EXPR_]], align 1269// CHECK1-NEXT: [[LOADEDV3:%.*]] = trunc i8 [[TMP16]] to i1270// CHECK1-NEXT: [[TMP17:%.*]] = select i1 [[LOADEDV3]], i32 0, i32 1271// CHECK1-NEXT: [[TMP18:%.*]] = insertvalue [3 x i32] zeroinitializer, i32 [[TMP17]], 0272// CHECK1-NEXT: [[TMP19:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0273// CHECK1-NEXT: store i32 3, ptr [[TMP19]], align 4274// CHECK1-NEXT: [[TMP20:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1275// CHECK1-NEXT: store i32 3, ptr [[TMP20]], align 4276// CHECK1-NEXT: [[TMP21:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2277// CHECK1-NEXT: store ptr [[TMP14]], ptr [[TMP21]], align 8278// CHECK1-NEXT: [[TMP22:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3279// CHECK1-NEXT: store ptr [[TMP15]], ptr [[TMP22]], align 8280// CHECK1-NEXT: [[TMP23:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4281// CHECK1-NEXT: store ptr @.offload_sizes, ptr [[TMP23]], align 8282// CHECK1-NEXT: [[TMP24:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5283// CHECK1-NEXT: store ptr @.offload_maptypes, ptr [[TMP24]], align 8284// CHECK1-NEXT: [[TMP25:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6285// CHECK1-NEXT: store ptr null, ptr [[TMP25]], align 8286// CHECK1-NEXT: [[TMP26:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7287// CHECK1-NEXT: store ptr null, ptr [[TMP26]], align 8288// CHECK1-NEXT: [[TMP27:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8289// CHECK1-NEXT: store i64 0, ptr [[TMP27]], align 8290// CHECK1-NEXT: [[TMP28:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9291// CHECK1-NEXT: store i64 0, ptr [[TMP28]], align 8292// CHECK1-NEXT: [[TMP29:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10293// CHECK1-NEXT: store [3 x i32] [i32 1, i32 0, i32 0], ptr [[TMP29]], align 4294// CHECK1-NEXT: [[TMP30:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11295// CHECK1-NEXT: store [3 x i32] [[TMP18]], ptr [[TMP30]], align 4296// CHECK1-NEXT: [[TMP31:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12297// CHECK1-NEXT: store i32 0, ptr [[TMP31]], align 4298// CHECK1-NEXT: [[TMP32:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1:[0-9]+]], i64 -1, i32 1, i32 [[TMP17]], ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2S12r1Ei_l121.region_id, ptr [[KERNEL_ARGS]])299// CHECK1-NEXT: [[TMP33:%.*]] = icmp ne i32 [[TMP32]], 0300// CHECK1-NEXT: br i1 [[TMP33]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]301// CHECK1: omp_offload.failed:302// CHECK1-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2S12r1Ei_l121(ptr [[THIS1]], i64 [[TMP2]], i64 [[TMP4]]) #[[ATTR2:[0-9]+]]303// CHECK1-NEXT: br label [[OMP_OFFLOAD_CONT]]304// CHECK1: omp_offload.cont:305// CHECK1-NEXT: [[TMP34:%.*]] = load i32, ptr [[N_ADDR]], align 4306// CHECK1-NEXT: [[CMP5:%.*]] = icmp sgt i32 [[TMP34]], 5307// CHECK1-NEXT: [[STOREDV6:%.*]] = zext i1 [[CMP5]] to i8308// CHECK1-NEXT: store i8 [[STOREDV6]], ptr [[DOTCAPTURE_EXPR_4]], align 1309// CHECK1-NEXT: [[TMP35:%.*]] = load i8, ptr [[DOTCAPTURE_EXPR_4]], align 1310// CHECK1-NEXT: [[LOADEDV7:%.*]] = trunc i8 [[TMP35]] to i1311// CHECK1-NEXT: [[STOREDV9:%.*]] = zext i1 [[LOADEDV7]] to i8312// CHECK1-NEXT: store i8 [[STOREDV9]], ptr [[DOTCAPTURE_EXPR__CASTED8]], align 1313// CHECK1-NEXT: [[TMP36:%.*]] = load i64, ptr [[DOTCAPTURE_EXPR__CASTED8]], align 8314// CHECK1-NEXT: [[TMP37:%.*]] = load i32, ptr [[N_ADDR]], align 4315// CHECK1-NEXT: [[CMP10:%.*]] = icmp sgt i32 [[TMP37]], 4316// CHECK1-NEXT: br i1 [[CMP10]], label [[OMP_IF_THEN:%.*]], label [[OMP_IF_ELSE:%.*]]317// CHECK1: omp_if.then:318// CHECK1-NEXT: [[A11:%.*]] = getelementptr inbounds nuw [[STRUCT_S1]], ptr [[THIS1]], i32 0, i32 0319// CHECK1-NEXT: [[TMP38:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS12]], i32 0, i32 0320// CHECK1-NEXT: store ptr [[THIS1]], ptr [[TMP38]], align 8321// CHECK1-NEXT: [[TMP39:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS13]], i32 0, i32 0322// CHECK1-NEXT: store ptr [[A11]], ptr [[TMP39]], align 8323// CHECK1-NEXT: [[TMP40:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS14]], i64 0, i64 0324// CHECK1-NEXT: store ptr null, ptr [[TMP40]], align 8325// CHECK1-NEXT: [[TMP41:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS12]], i32 0, i32 1326// CHECK1-NEXT: store i64 [[TMP36]], ptr [[TMP41]], align 8327// CHECK1-NEXT: [[TMP42:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS13]], i32 0, i32 1328// CHECK1-NEXT: store i64 [[TMP36]], ptr [[TMP42]], align 8329// CHECK1-NEXT: [[TMP43:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS14]], i64 0, i64 1330// CHECK1-NEXT: store ptr null, ptr [[TMP43]], align 8331// CHECK1-NEXT: [[TMP44:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS12]], i32 0, i32 0332// CHECK1-NEXT: [[TMP45:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS13]], i32 0, i32 0333// CHECK1-NEXT: [[TMP46:%.*]] = load i8, ptr [[DOTCAPTURE_EXPR_4]], align 1334// CHECK1-NEXT: [[LOADEDV15:%.*]] = trunc i8 [[TMP46]] to i1335// CHECK1-NEXT: [[TMP47:%.*]] = select i1 [[LOADEDV15]], i32 0, i32 1336// CHECK1-NEXT: [[TMP48:%.*]] = insertvalue [3 x i32] zeroinitializer, i32 [[TMP47]], 0337// CHECK1-NEXT: [[TMP49:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS16]], i32 0, i32 0338// CHECK1-NEXT: store i32 3, ptr [[TMP49]], align 4339// CHECK1-NEXT: [[TMP50:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS16]], i32 0, i32 1340// CHECK1-NEXT: store i32 2, ptr [[TMP50]], align 4341// CHECK1-NEXT: [[TMP51:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS16]], i32 0, i32 2342// CHECK1-NEXT: store ptr [[TMP44]], ptr [[TMP51]], align 8343// CHECK1-NEXT: [[TMP52:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS16]], i32 0, i32 3344// CHECK1-NEXT: store ptr [[TMP45]], ptr [[TMP52]], align 8345// CHECK1-NEXT: [[TMP53:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS16]], i32 0, i32 4346// CHECK1-NEXT: store ptr @.offload_sizes.1, ptr [[TMP53]], align 8347// CHECK1-NEXT: [[TMP54:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS16]], i32 0, i32 5348// CHECK1-NEXT: store ptr @.offload_maptypes.2, ptr [[TMP54]], align 8349// CHECK1-NEXT: [[TMP55:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS16]], i32 0, i32 6350// CHECK1-NEXT: store ptr null, ptr [[TMP55]], align 8351// CHECK1-NEXT: [[TMP56:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS16]], i32 0, i32 7352// CHECK1-NEXT: store ptr null, ptr [[TMP56]], align 8353// CHECK1-NEXT: [[TMP57:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS16]], i32 0, i32 8354// CHECK1-NEXT: store i64 0, ptr [[TMP57]], align 8355// CHECK1-NEXT: [[TMP58:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS16]], i32 0, i32 9356// CHECK1-NEXT: store i64 0, ptr [[TMP58]], align 8357// CHECK1-NEXT: [[TMP59:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS16]], i32 0, i32 10358// CHECK1-NEXT: store [3 x i32] [i32 1, i32 0, i32 0], ptr [[TMP59]], align 4359// CHECK1-NEXT: [[TMP60:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS16]], i32 0, i32 11360// CHECK1-NEXT: store [3 x i32] [[TMP48]], ptr [[TMP60]], align 4361// CHECK1-NEXT: [[TMP61:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS16]], i32 0, i32 12362// CHECK1-NEXT: store i32 0, ptr [[TMP61]], align 4363// CHECK1-NEXT: [[TMP62:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 -1, i32 1, i32 [[TMP47]], ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2S12r1Ei_l126.region_id, ptr [[KERNEL_ARGS16]])364// CHECK1-NEXT: [[TMP63:%.*]] = icmp ne i32 [[TMP62]], 0365// CHECK1-NEXT: br i1 [[TMP63]], label [[OMP_OFFLOAD_FAILED17:%.*]], label [[OMP_OFFLOAD_CONT18:%.*]]366// CHECK1: omp_offload.failed17:367// CHECK1-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2S12r1Ei_l126(ptr [[THIS1]], i64 [[TMP36]]) #[[ATTR2]]368// CHECK1-NEXT: br label [[OMP_OFFLOAD_CONT18]]369// CHECK1: omp_offload.cont18:370// CHECK1-NEXT: br label [[OMP_IF_END:%.*]]371// CHECK1: omp_if.else:372// CHECK1-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2S12r1Ei_l126(ptr [[THIS1]], i64 [[TMP36]]) #[[ATTR2]]373// CHECK1-NEXT: br label [[OMP_IF_END]]374// CHECK1: omp_if.end:375// CHECK1-NEXT: [[A19:%.*]] = getelementptr inbounds nuw [[STRUCT_S1]], ptr [[THIS1]], i32 0, i32 0376// CHECK1-NEXT: [[TMP64:%.*]] = load double, ptr [[A19]], align 8377// CHECK1-NEXT: [[CONV:%.*]] = fptosi double [[TMP64]] to i32378// CHECK1-NEXT: ret i32 [[CONV]]379//380//381// CHECK1-LABEL: define {{[^@]+}}@_ZL7fstatici382// CHECK1-SAME: (i32 noundef signext [[N:%.*]]) #[[ATTR0]] {383// CHECK1-NEXT: entry:384// CHECK1-NEXT: [[N_ADDR:%.*]] = alloca i32, align 4385// CHECK1-NEXT: [[DOTCAPTURE_EXPR_:%.*]] = alloca i8, align 1386// CHECK1-NEXT: [[DOTCAPTURE_EXPR__CASTED:%.*]] = alloca i64, align 8387// CHECK1-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x ptr], align 8388// CHECK1-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x ptr], align 8389// CHECK1-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [1 x ptr], align 8390// CHECK1-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8391// CHECK1-NEXT: [[KERNEL_ARGS6:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 8392// CHECK1-NEXT: store i32 [[N]], ptr [[N_ADDR]], align 4393// CHECK1-NEXT: [[TMP0:%.*]] = load i32, ptr [[N_ADDR]], align 4394// CHECK1-NEXT: [[CMP:%.*]] = icmp sgt i32 [[TMP0]], 1395// CHECK1-NEXT: [[STOREDV:%.*]] = zext i1 [[CMP]] to i8396// CHECK1-NEXT: store i8 [[STOREDV]], ptr [[DOTCAPTURE_EXPR_]], align 1397// CHECK1-NEXT: [[TMP1:%.*]] = load i8, ptr [[DOTCAPTURE_EXPR_]], align 1398// CHECK1-NEXT: [[LOADEDV:%.*]] = trunc i8 [[TMP1]] to i1399// CHECK1-NEXT: [[STOREDV1:%.*]] = zext i1 [[LOADEDV]] to i8400// CHECK1-NEXT: store i8 [[STOREDV1]], ptr [[DOTCAPTURE_EXPR__CASTED]], align 1401// CHECK1-NEXT: [[TMP2:%.*]] = load i64, ptr [[DOTCAPTURE_EXPR__CASTED]], align 8402// CHECK1-NEXT: [[TMP3:%.*]] = load i8, ptr [[DOTCAPTURE_EXPR_]], align 1403// CHECK1-NEXT: [[LOADEDV2:%.*]] = trunc i8 [[TMP3]] to i1404// CHECK1-NEXT: br i1 [[LOADEDV2]], label [[OMP_IF_THEN:%.*]], label [[OMP_IF_ELSE:%.*]]405// CHECK1: omp_if.then:406// CHECK1-NEXT: [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0407// CHECK1-NEXT: store i64 [[TMP2]], ptr [[TMP4]], align 8408// CHECK1-NEXT: [[TMP5:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0409// CHECK1-NEXT: store i64 [[TMP2]], ptr [[TMP5]], align 8410// CHECK1-NEXT: [[TMP6:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 0411// CHECK1-NEXT: store ptr null, ptr [[TMP6]], align 8412// CHECK1-NEXT: [[TMP7:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0413// CHECK1-NEXT: [[TMP8:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0414// CHECK1-NEXT: [[TMP9:%.*]] = load i8, ptr [[DOTCAPTURE_EXPR_]], align 1415// CHECK1-NEXT: [[LOADEDV3:%.*]] = trunc i8 [[TMP9]] to i1416// CHECK1-NEXT: [[TMP10:%.*]] = select i1 [[LOADEDV3]], i32 0, i32 1417// CHECK1-NEXT: [[TMP11:%.*]] = insertvalue [3 x i32] zeroinitializer, i32 [[TMP10]], 0418// CHECK1-NEXT: [[TMP12:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0419// CHECK1-NEXT: store i32 3, ptr [[TMP12]], align 4420// CHECK1-NEXT: [[TMP13:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1421// CHECK1-NEXT: store i32 1, ptr [[TMP13]], align 4422// CHECK1-NEXT: [[TMP14:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2423// CHECK1-NEXT: store ptr [[TMP7]], ptr [[TMP14]], align 8424// CHECK1-NEXT: [[TMP15:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3425// CHECK1-NEXT: store ptr [[TMP8]], ptr [[TMP15]], align 8426// CHECK1-NEXT: [[TMP16:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4427// CHECK1-NEXT: store ptr @.offload_sizes.3, ptr [[TMP16]], align 8428// CHECK1-NEXT: [[TMP17:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5429// CHECK1-NEXT: store ptr @.offload_maptypes.4, ptr [[TMP17]], align 8430// CHECK1-NEXT: [[TMP18:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6431// CHECK1-NEXT: store ptr null, ptr [[TMP18]], align 8432// CHECK1-NEXT: [[TMP19:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7433// CHECK1-NEXT: store ptr null, ptr [[TMP19]], align 8434// CHECK1-NEXT: [[TMP20:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8435// CHECK1-NEXT: store i64 0, ptr [[TMP20]], align 8436// CHECK1-NEXT: [[TMP21:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9437// CHECK1-NEXT: store i64 0, ptr [[TMP21]], align 8438// CHECK1-NEXT: [[TMP22:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10439// CHECK1-NEXT: store [3 x i32] [i32 1, i32 0, i32 0], ptr [[TMP22]], align 4440// CHECK1-NEXT: [[TMP23:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11441// CHECK1-NEXT: store [3 x i32] [[TMP11]], ptr [[TMP23]], align 4442// CHECK1-NEXT: [[TMP24:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12443// CHECK1-NEXT: store i32 0, ptr [[TMP24]], align 4444// CHECK1-NEXT: [[TMP25:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 -1, i32 1, i32 [[TMP10]], ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l104.region_id, ptr [[KERNEL_ARGS]])445// CHECK1-NEXT: [[TMP26:%.*]] = icmp ne i32 [[TMP25]], 0446// CHECK1-NEXT: br i1 [[TMP26]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]447// CHECK1: omp_offload.failed:448// CHECK1-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l104(i64 [[TMP2]]) #[[ATTR2]]449// CHECK1-NEXT: br label [[OMP_OFFLOAD_CONT]]450// CHECK1: omp_offload.cont:451// CHECK1-NEXT: br label [[OMP_IF_END:%.*]]452// CHECK1: omp_if.else:453// CHECK1-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l104(i64 [[TMP2]]) #[[ATTR2]]454// CHECK1-NEXT: br label [[OMP_IF_END]]455// CHECK1: omp_if.end:456// CHECK1-NEXT: [[TMP27:%.*]] = load i32, ptr [[N_ADDR]], align 4457// CHECK1-NEXT: [[SUB:%.*]] = sub nsw i32 [[TMP27]], 2458// CHECK1-NEXT: [[CMP4:%.*]] = icmp sgt i32 [[SUB]], 2459// CHECK1-NEXT: br i1 [[CMP4]], label [[OMP_IF_THEN5:%.*]], label [[OMP_IF_ELSE9:%.*]]460// CHECK1: omp_if.then5:461// CHECK1-NEXT: [[TMP28:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 0462// CHECK1-NEXT: store i32 3, ptr [[TMP28]], align 4463// CHECK1-NEXT: [[TMP29:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 1464// CHECK1-NEXT: store i32 0, ptr [[TMP29]], align 4465// CHECK1-NEXT: [[TMP30:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 2466// CHECK1-NEXT: store ptr null, ptr [[TMP30]], align 8467// CHECK1-NEXT: [[TMP31:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 3468// CHECK1-NEXT: store ptr null, ptr [[TMP31]], align 8469// CHECK1-NEXT: [[TMP32:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 4470// CHECK1-NEXT: store ptr null, ptr [[TMP32]], align 8471// CHECK1-NEXT: [[TMP33:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 5472// CHECK1-NEXT: store ptr null, ptr [[TMP33]], align 8473// CHECK1-NEXT: [[TMP34:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 6474// CHECK1-NEXT: store ptr null, ptr [[TMP34]], align 8475// CHECK1-NEXT: [[TMP35:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 7476// CHECK1-NEXT: store ptr null, ptr [[TMP35]], align 8477// CHECK1-NEXT: [[TMP36:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 8478// CHECK1-NEXT: store i64 0, ptr [[TMP36]], align 8479// CHECK1-NEXT: [[TMP37:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 9480// CHECK1-NEXT: store i64 0, ptr [[TMP37]], align 8481// CHECK1-NEXT: [[TMP38:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 10482// CHECK1-NEXT: store [3 x i32] [i32 1, i32 0, i32 0], ptr [[TMP38]], align 4483// CHECK1-NEXT: [[TMP39:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 11484// CHECK1-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP39]], align 4485// CHECK1-NEXT: [[TMP40:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 12486// CHECK1-NEXT: store i32 0, ptr [[TMP40]], align 4487// CHECK1-NEXT: [[TMP41:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 -1, i32 1, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l108.region_id, ptr [[KERNEL_ARGS6]])488// CHECK1-NEXT: [[TMP42:%.*]] = icmp ne i32 [[TMP41]], 0489// CHECK1-NEXT: br i1 [[TMP42]], label [[OMP_OFFLOAD_FAILED7:%.*]], label [[OMP_OFFLOAD_CONT8:%.*]]490// CHECK1: omp_offload.failed7:491// CHECK1-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l108() #[[ATTR2]]492// CHECK1-NEXT: br label [[OMP_OFFLOAD_CONT8]]493// CHECK1: omp_offload.cont8:494// CHECK1-NEXT: br label [[OMP_IF_END10:%.*]]495// CHECK1: omp_if.else9:496// CHECK1-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l108() #[[ATTR2]]497// CHECK1-NEXT: br label [[OMP_IF_END10]]498// CHECK1: omp_if.end10:499// CHECK1-NEXT: [[TMP43:%.*]] = load i32, ptr [[N_ADDR]], align 4500// CHECK1-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP43]], 1501// CHECK1-NEXT: ret i32 [[ADD]]502//503//504// CHECK1-LABEL: define {{[^@]+}}@_Z9ftemplateIiET_i505// CHECK1-SAME: (i32 noundef signext [[N:%.*]]) #[[ATTR0]] comdat {506// CHECK1-NEXT: entry:507// CHECK1-NEXT: [[N_ADDR:%.*]] = alloca i32, align 4508// CHECK1-NEXT: [[A:%.*]] = alloca i32, align 4509// CHECK1-NEXT: [[A_CASTED:%.*]] = alloca i64, align 8510// CHECK1-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x ptr], align 8511// CHECK1-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x ptr], align 8512// CHECK1-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [1 x ptr], align 8513// CHECK1-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8514// CHECK1-NEXT: [[B:%.*]] = alloca i16, align 2515// CHECK1-NEXT: [[A_CASTED1:%.*]] = alloca i64, align 8516// CHECK1-NEXT: [[B_CASTED:%.*]] = alloca i64, align 8517// CHECK1-NEXT: [[DOTOFFLOAD_BASEPTRS2:%.*]] = alloca [2 x ptr], align 8518// CHECK1-NEXT: [[DOTOFFLOAD_PTRS3:%.*]] = alloca [2 x ptr], align 8519// CHECK1-NEXT: [[DOTOFFLOAD_MAPPERS4:%.*]] = alloca [2 x ptr], align 8520// CHECK1-NEXT: [[KERNEL_ARGS5:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 8521// CHECK1-NEXT: store i32 [[N]], ptr [[N_ADDR]], align 4522// CHECK1-NEXT: store i32 0, ptr [[A]], align 4523// CHECK1-NEXT: [[TMP0:%.*]] = load i32, ptr [[A]], align 4524// CHECK1-NEXT: store i32 [[TMP0]], ptr [[A_CASTED]], align 4525// CHECK1-NEXT: [[TMP1:%.*]] = load i64, ptr [[A_CASTED]], align 8526// CHECK1-NEXT: [[TMP2:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0527// CHECK1-NEXT: store i64 [[TMP1]], ptr [[TMP2]], align 8528// CHECK1-NEXT: [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0529// CHECK1-NEXT: store i64 [[TMP1]], ptr [[TMP3]], align 8530// CHECK1-NEXT: [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 0531// CHECK1-NEXT: store ptr null, ptr [[TMP4]], align 8532// CHECK1-NEXT: [[TMP5:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0533// CHECK1-NEXT: [[TMP6:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0534// CHECK1-NEXT: [[TMP7:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0535// CHECK1-NEXT: store i32 3, ptr [[TMP7]], align 4536// CHECK1-NEXT: [[TMP8:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1537// CHECK1-NEXT: store i32 1, ptr [[TMP8]], align 4538// CHECK1-NEXT: [[TMP9:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2539// CHECK1-NEXT: store ptr [[TMP5]], ptr [[TMP9]], align 8540// CHECK1-NEXT: [[TMP10:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3541// CHECK1-NEXT: store ptr [[TMP6]], ptr [[TMP10]], align 8542// CHECK1-NEXT: [[TMP11:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4543// CHECK1-NEXT: store ptr @.offload_sizes.5, ptr [[TMP11]], align 8544// CHECK1-NEXT: [[TMP12:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5545// CHECK1-NEXT: store ptr @.offload_maptypes.6, ptr [[TMP12]], align 8546// CHECK1-NEXT: [[TMP13:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6547// CHECK1-NEXT: store ptr null, ptr [[TMP13]], align 8548// CHECK1-NEXT: [[TMP14:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7549// CHECK1-NEXT: store ptr null, ptr [[TMP14]], align 8550// CHECK1-NEXT: [[TMP15:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8551// CHECK1-NEXT: store i64 0, ptr [[TMP15]], align 8552// CHECK1-NEXT: [[TMP16:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9553// CHECK1-NEXT: store i64 0, ptr [[TMP16]], align 8554// CHECK1-NEXT: [[TMP17:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10555// CHECK1-NEXT: store [3 x i32] [i32 1, i32 0, i32 0], ptr [[TMP17]], align 4556// CHECK1-NEXT: [[TMP18:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11557// CHECK1-NEXT: store [3 x i32] [i32 1, i32 0, i32 0], ptr [[TMP18]], align 4558// CHECK1-NEXT: [[TMP19:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12559// CHECK1-NEXT: store i32 0, ptr [[TMP19]], align 4560// CHECK1-NEXT: [[TMP20:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 -1, i32 1, i32 1, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l87.region_id, ptr [[KERNEL_ARGS]])561// CHECK1-NEXT: [[TMP21:%.*]] = icmp ne i32 [[TMP20]], 0562// CHECK1-NEXT: br i1 [[TMP21]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]563// CHECK1: omp_offload.failed:564// CHECK1-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l87(i64 [[TMP1]]) #[[ATTR2]]565// CHECK1-NEXT: br label [[OMP_OFFLOAD_CONT]]566// CHECK1: omp_offload.cont:567// CHECK1-NEXT: store i16 1, ptr [[B]], align 2568// CHECK1-NEXT: [[TMP22:%.*]] = load i32, ptr [[A]], align 4569// CHECK1-NEXT: store i32 [[TMP22]], ptr [[A_CASTED1]], align 4570// CHECK1-NEXT: [[TMP23:%.*]] = load i64, ptr [[A_CASTED1]], align 8571// CHECK1-NEXT: [[TMP24:%.*]] = load i16, ptr [[B]], align 2572// CHECK1-NEXT: store i16 [[TMP24]], ptr [[B_CASTED]], align 2573// CHECK1-NEXT: [[TMP25:%.*]] = load i64, ptr [[B_CASTED]], align 8574// CHECK1-NEXT: [[TMP26:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS2]], i32 0, i32 0575// CHECK1-NEXT: store i64 [[TMP23]], ptr [[TMP26]], align 8576// CHECK1-NEXT: [[TMP27:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS3]], i32 0, i32 0577// CHECK1-NEXT: store i64 [[TMP23]], ptr [[TMP27]], align 8578// CHECK1-NEXT: [[TMP28:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS4]], i64 0, i64 0579// CHECK1-NEXT: store ptr null, ptr [[TMP28]], align 8580// CHECK1-NEXT: [[TMP29:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS2]], i32 0, i32 1581// CHECK1-NEXT: store i64 [[TMP25]], ptr [[TMP29]], align 8582// CHECK1-NEXT: [[TMP30:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS3]], i32 0, i32 1583// CHECK1-NEXT: store i64 [[TMP25]], ptr [[TMP30]], align 8584// CHECK1-NEXT: [[TMP31:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS4]], i64 0, i64 1585// CHECK1-NEXT: store ptr null, ptr [[TMP31]], align 8586// CHECK1-NEXT: [[TMP32:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS2]], i32 0, i32 0587// CHECK1-NEXT: [[TMP33:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS3]], i32 0, i32 0588// CHECK1-NEXT: [[TMP34:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 0589// CHECK1-NEXT: store i32 3, ptr [[TMP34]], align 4590// CHECK1-NEXT: [[TMP35:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 1591// CHECK1-NEXT: store i32 2, ptr [[TMP35]], align 4592// CHECK1-NEXT: [[TMP36:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 2593// CHECK1-NEXT: store ptr [[TMP32]], ptr [[TMP36]], align 8594// CHECK1-NEXT: [[TMP37:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 3595// CHECK1-NEXT: store ptr [[TMP33]], ptr [[TMP37]], align 8596// CHECK1-NEXT: [[TMP38:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 4597// CHECK1-NEXT: store ptr @.offload_sizes.7, ptr [[TMP38]], align 8598// CHECK1-NEXT: [[TMP39:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 5599// CHECK1-NEXT: store ptr @.offload_maptypes.8, ptr [[TMP39]], align 8600// CHECK1-NEXT: [[TMP40:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 6601// CHECK1-NEXT: store ptr null, ptr [[TMP40]], align 8602// CHECK1-NEXT: [[TMP41:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 7603// CHECK1-NEXT: store ptr null, ptr [[TMP41]], align 8604// CHECK1-NEXT: [[TMP42:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 8605// CHECK1-NEXT: store i64 0, ptr [[TMP42]], align 8606// CHECK1-NEXT: [[TMP43:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 9607// CHECK1-NEXT: store i64 0, ptr [[TMP43]], align 8608// CHECK1-NEXT: [[TMP44:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 10609// CHECK1-NEXT: store [3 x i32] [i32 1, i32 0, i32 0], ptr [[TMP44]], align 4610// CHECK1-NEXT: [[TMP45:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 11611// CHECK1-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP45]], align 4612// CHECK1-NEXT: [[TMP46:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 12613// CHECK1-NEXT: store i32 0, ptr [[TMP46]], align 4614// CHECK1-NEXT: [[TMP47:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 -1, i32 1, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l93.region_id, ptr [[KERNEL_ARGS5]])615// CHECK1-NEXT: [[TMP48:%.*]] = icmp ne i32 [[TMP47]], 0616// CHECK1-NEXT: br i1 [[TMP48]], label [[OMP_OFFLOAD_FAILED6:%.*]], label [[OMP_OFFLOAD_CONT7:%.*]]617// CHECK1: omp_offload.failed6:618// CHECK1-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l93(i64 [[TMP23]], i64 [[TMP25]]) #[[ATTR2]]619// CHECK1-NEXT: br label [[OMP_OFFLOAD_CONT7]]620// CHECK1: omp_offload.cont7:621// CHECK1-NEXT: [[TMP49:%.*]] = load i32, ptr [[A]], align 4622// CHECK1-NEXT: ret i32 [[TMP49]]623//624//625// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2S12r1Ei_l121626// CHECK1-SAME: (ptr noundef [[THIS:%.*]], i64 noundef [[B:%.*]], i64 noundef [[DOTCAPTURE_EXPR_:%.*]]) #[[ATTR1:[0-9]+]] {627// CHECK1-NEXT: entry:628// CHECK1-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 8629// CHECK1-NEXT: [[B_ADDR:%.*]] = alloca i64, align 8630// CHECK1-NEXT: [[DOTCAPTURE_EXPR__ADDR:%.*]] = alloca i64, align 8631// CHECK1-NEXT: [[B_CASTED:%.*]] = alloca i64, align 8632// CHECK1-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 4633// CHECK1-NEXT: [[DOTBOUND_ZERO_ADDR:%.*]] = alloca i32, align 4634// CHECK1-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])635// CHECK1-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 8636// CHECK1-NEXT: store i64 [[B]], ptr [[B_ADDR]], align 8637// CHECK1-NEXT: store i64 [[DOTCAPTURE_EXPR_]], ptr [[DOTCAPTURE_EXPR__ADDR]], align 8638// CHECK1-NEXT: [[TMP1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8639// CHECK1-NEXT: [[TMP2:%.*]] = load i32, ptr [[B_ADDR]], align 4640// CHECK1-NEXT: store i32 [[TMP2]], ptr [[B_CASTED]], align 4641// CHECK1-NEXT: [[TMP3:%.*]] = load i64, ptr [[B_CASTED]], align 8642// CHECK1-NEXT: [[TMP4:%.*]] = load i8, ptr [[DOTCAPTURE_EXPR__ADDR]], align 1643// CHECK1-NEXT: [[LOADEDV:%.*]] = trunc i8 [[TMP4]] to i1644// CHECK1-NEXT: br i1 [[LOADEDV]], label [[OMP_IF_THEN:%.*]], label [[OMP_IF_ELSE:%.*]]645// CHECK1: omp_if.then:646// CHECK1-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB1]], i32 2, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2S12r1Ei_l121.omp_outlined, ptr [[TMP1]], i64 [[TMP3]])647// CHECK1-NEXT: br label [[OMP_IF_END:%.*]]648// CHECK1: omp_if.else:649// CHECK1-NEXT: call void @__kmpc_serialized_parallel(ptr @[[GLOB1]], i32 [[TMP0]])650// CHECK1-NEXT: store i32 [[TMP0]], ptr [[DOTTHREADID_TEMP_]], align 4651// CHECK1-NEXT: store i32 0, ptr [[DOTBOUND_ZERO_ADDR]], align 4652// CHECK1-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2S12r1Ei_l121.omp_outlined(ptr [[DOTTHREADID_TEMP_]], ptr [[DOTBOUND_ZERO_ADDR]], ptr [[TMP1]], i64 [[TMP3]]) #[[ATTR2]]653// CHECK1-NEXT: call void @__kmpc_end_serialized_parallel(ptr @[[GLOB1]], i32 [[TMP0]])654// CHECK1-NEXT: br label [[OMP_IF_END]]655// CHECK1: omp_if.end:656// CHECK1-NEXT: ret void657//658//659// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2S12r1Ei_l121.omp_outlined660// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef [[THIS:%.*]], i64 noundef [[B:%.*]]) #[[ATTR1]] {661// CHECK1-NEXT: entry:662// CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8663// CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8664// CHECK1-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 8665// CHECK1-NEXT: [[B_ADDR:%.*]] = alloca i64, align 8666// CHECK1-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8667// CHECK1-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8668// CHECK1-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 8669// CHECK1-NEXT: store i64 [[B]], ptr [[B_ADDR]], align 8670// CHECK1-NEXT: [[TMP0:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8671// CHECK1-NEXT: [[TMP1:%.*]] = load i32, ptr [[B_ADDR]], align 4672// CHECK1-NEXT: [[CONV:%.*]] = sitofp i32 [[TMP1]] to double673// CHECK1-NEXT: [[ADD:%.*]] = fadd double [[CONV]], 1.500000e+00674// CHECK1-NEXT: [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_S1:%.*]], ptr [[TMP0]], i32 0, i32 0675// CHECK1-NEXT: store double [[ADD]], ptr [[A]], align 8676// CHECK1-NEXT: ret void677//678//679// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2S12r1Ei_l126680// CHECK1-SAME: (ptr noundef [[THIS:%.*]], i64 noundef [[DOTCAPTURE_EXPR_:%.*]]) #[[ATTR1]] {681// CHECK1-NEXT: entry:682// CHECK1-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 8683// CHECK1-NEXT: [[DOTCAPTURE_EXPR__ADDR:%.*]] = alloca i64, align 8684// CHECK1-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 4685// CHECK1-NEXT: [[DOTBOUND_ZERO_ADDR:%.*]] = alloca i32, align 4686// CHECK1-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])687// CHECK1-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 8688// CHECK1-NEXT: store i64 [[DOTCAPTURE_EXPR_]], ptr [[DOTCAPTURE_EXPR__ADDR]], align 8689// CHECK1-NEXT: [[TMP1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8690// CHECK1-NEXT: [[TMP2:%.*]] = load i8, ptr [[DOTCAPTURE_EXPR__ADDR]], align 1691// CHECK1-NEXT: [[LOADEDV:%.*]] = trunc i8 [[TMP2]] to i1692// CHECK1-NEXT: br i1 [[LOADEDV]], label [[OMP_IF_THEN:%.*]], label [[OMP_IF_ELSE:%.*]]693// CHECK1: omp_if.then:694// CHECK1-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB1]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2S12r1Ei_l126.omp_outlined, ptr [[TMP1]])695// CHECK1-NEXT: br label [[OMP_IF_END:%.*]]696// CHECK1: omp_if.else:697// CHECK1-NEXT: call void @__kmpc_serialized_parallel(ptr @[[GLOB1]], i32 [[TMP0]])698// CHECK1-NEXT: store i32 [[TMP0]], ptr [[DOTTHREADID_TEMP_]], align 4699// CHECK1-NEXT: store i32 0, ptr [[DOTBOUND_ZERO_ADDR]], align 4700// CHECK1-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2S12r1Ei_l126.omp_outlined(ptr [[DOTTHREADID_TEMP_]], ptr [[DOTBOUND_ZERO_ADDR]], ptr [[TMP1]]) #[[ATTR2]]701// CHECK1-NEXT: call void @__kmpc_end_serialized_parallel(ptr @[[GLOB1]], i32 [[TMP0]])702// CHECK1-NEXT: br label [[OMP_IF_END]]703// CHECK1: omp_if.end:704// CHECK1-NEXT: ret void705//706//707// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2S12r1Ei_l126.omp_outlined708// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef [[THIS:%.*]]) #[[ATTR1]] {709// CHECK1-NEXT: entry:710// CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8711// CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8712// CHECK1-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 8713// CHECK1-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8714// CHECK1-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8715// CHECK1-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 8716// CHECK1-NEXT: [[TMP0:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8717// CHECK1-NEXT: [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_S1:%.*]], ptr [[TMP0]], i32 0, i32 0718// CHECK1-NEXT: store double 2.500000e+00, ptr [[A]], align 8719// CHECK1-NEXT: ret void720//721//722// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l104723// CHECK1-SAME: (i64 noundef [[DOTCAPTURE_EXPR_:%.*]]) #[[ATTR1]] {724// CHECK1-NEXT: entry:725// CHECK1-NEXT: [[DOTCAPTURE_EXPR__ADDR:%.*]] = alloca i64, align 8726// CHECK1-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 4727// CHECK1-NEXT: [[DOTBOUND_ZERO_ADDR:%.*]] = alloca i32, align 4728// CHECK1-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])729// CHECK1-NEXT: store i64 [[DOTCAPTURE_EXPR_]], ptr [[DOTCAPTURE_EXPR__ADDR]], align 8730// CHECK1-NEXT: [[TMP1:%.*]] = load i8, ptr [[DOTCAPTURE_EXPR__ADDR]], align 1731// CHECK1-NEXT: [[LOADEDV:%.*]] = trunc i8 [[TMP1]] to i1732// CHECK1-NEXT: br i1 [[LOADEDV]], label [[OMP_IF_THEN:%.*]], label [[OMP_IF_ELSE:%.*]]733// CHECK1: omp_if.then:734// CHECK1-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB1]], i32 0, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l104.omp_outlined)735// CHECK1-NEXT: br label [[OMP_IF_END:%.*]]736// CHECK1: omp_if.else:737// CHECK1-NEXT: call void @__kmpc_serialized_parallel(ptr @[[GLOB1]], i32 [[TMP0]])738// CHECK1-NEXT: store i32 [[TMP0]], ptr [[DOTTHREADID_TEMP_]], align 4739// CHECK1-NEXT: store i32 0, ptr [[DOTBOUND_ZERO_ADDR]], align 4740// CHECK1-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l104.omp_outlined(ptr [[DOTTHREADID_TEMP_]], ptr [[DOTBOUND_ZERO_ADDR]]) #[[ATTR2]]741// CHECK1-NEXT: call void @__kmpc_end_serialized_parallel(ptr @[[GLOB1]], i32 [[TMP0]])742// CHECK1-NEXT: br label [[OMP_IF_END]]743// CHECK1: omp_if.end:744// CHECK1-NEXT: ret void745//746//747// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l104.omp_outlined748// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR1]] {749// CHECK1-NEXT: entry:750// CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8751// CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8752// CHECK1-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8753// CHECK1-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8754// CHECK1-NEXT: ret void755//756//757// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l108758// CHECK1-SAME: () #[[ATTR1]] {759// CHECK1-NEXT: entry:760// CHECK1-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB1]], i32 0, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l108.omp_outlined)761// CHECK1-NEXT: ret void762//763//764// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l108.omp_outlined765// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR1]] {766// CHECK1-NEXT: entry:767// CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8768// CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8769// CHECK1-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8770// CHECK1-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8771// CHECK1-NEXT: ret void772//773//774// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l87775// CHECK1-SAME: (i64 noundef [[A:%.*]]) #[[ATTR1]] {776// CHECK1-NEXT: entry:777// CHECK1-NEXT: [[A_ADDR:%.*]] = alloca i64, align 8778// CHECK1-NEXT: [[A_CASTED:%.*]] = alloca i64, align 8779// CHECK1-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 4780// CHECK1-NEXT: [[DOTBOUND_ZERO_ADDR:%.*]] = alloca i32, align 4781// CHECK1-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])782// CHECK1-NEXT: store i64 [[A]], ptr [[A_ADDR]], align 8783// CHECK1-NEXT: [[TMP1:%.*]] = load i32, ptr [[A_ADDR]], align 4784// CHECK1-NEXT: store i32 [[TMP1]], ptr [[A_CASTED]], align 4785// CHECK1-NEXT: [[TMP2:%.*]] = load i64, ptr [[A_CASTED]], align 8786// CHECK1-NEXT: call void @__kmpc_serialized_parallel(ptr @[[GLOB1]], i32 [[TMP0]])787// CHECK1-NEXT: store i32 [[TMP0]], ptr [[DOTTHREADID_TEMP_]], align 4788// CHECK1-NEXT: store i32 0, ptr [[DOTBOUND_ZERO_ADDR]], align 4789// CHECK1-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l87.omp_outlined(ptr [[DOTTHREADID_TEMP_]], ptr [[DOTBOUND_ZERO_ADDR]], i64 [[TMP2]]) #[[ATTR2]]790// CHECK1-NEXT: call void @__kmpc_end_serialized_parallel(ptr @[[GLOB1]], i32 [[TMP0]])791// CHECK1-NEXT: ret void792//793//794// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l87.omp_outlined795// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i64 noundef [[A:%.*]]) #[[ATTR1]] {796// CHECK1-NEXT: entry:797// CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8798// CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8799// CHECK1-NEXT: [[A_ADDR:%.*]] = alloca i64, align 8800// CHECK1-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8801// CHECK1-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8802// CHECK1-NEXT: store i64 [[A]], ptr [[A_ADDR]], align 8803// CHECK1-NEXT: [[TMP0:%.*]] = load i32, ptr [[A_ADDR]], align 4804// CHECK1-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP0]], 1805// CHECK1-NEXT: store i32 [[ADD]], ptr [[A_ADDR]], align 4806// CHECK1-NEXT: ret void807//808//809// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l93810// CHECK1-SAME: (i64 noundef [[A:%.*]], i64 noundef [[B:%.*]]) #[[ATTR1]] {811// CHECK1-NEXT: entry:812// CHECK1-NEXT: [[A_ADDR:%.*]] = alloca i64, align 8813// CHECK1-NEXT: [[B_ADDR:%.*]] = alloca i64, align 8814// CHECK1-NEXT: [[A_CASTED:%.*]] = alloca i64, align 8815// CHECK1-NEXT: [[B_CASTED:%.*]] = alloca i64, align 8816// CHECK1-NEXT: store i64 [[A]], ptr [[A_ADDR]], align 8817// CHECK1-NEXT: store i64 [[B]], ptr [[B_ADDR]], align 8818// CHECK1-NEXT: [[TMP0:%.*]] = load i32, ptr [[A_ADDR]], align 4819// CHECK1-NEXT: store i32 [[TMP0]], ptr [[A_CASTED]], align 4820// CHECK1-NEXT: [[TMP1:%.*]] = load i64, ptr [[A_CASTED]], align 8821// CHECK1-NEXT: [[TMP2:%.*]] = load i16, ptr [[B_ADDR]], align 2822// CHECK1-NEXT: store i16 [[TMP2]], ptr [[B_CASTED]], align 2823// CHECK1-NEXT: [[TMP3:%.*]] = load i64, ptr [[B_CASTED]], align 8824// CHECK1-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB1]], i32 2, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l93.omp_outlined, i64 [[TMP1]], i64 [[TMP3]])825// CHECK1-NEXT: ret void826//827//828// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l93.omp_outlined829// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i64 noundef [[A:%.*]], i64 noundef [[B:%.*]]) #[[ATTR1]] {830// CHECK1-NEXT: entry:831// CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8832// CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8833// CHECK1-NEXT: [[A_ADDR:%.*]] = alloca i64, align 8834// CHECK1-NEXT: [[B_ADDR:%.*]] = alloca i64, align 8835// CHECK1-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8836// CHECK1-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8837// CHECK1-NEXT: store i64 [[A]], ptr [[A_ADDR]], align 8838// CHECK1-NEXT: store i64 [[B]], ptr [[B_ADDR]], align 8839// CHECK1-NEXT: [[TMP0:%.*]] = load i16, ptr [[B_ADDR]], align 2840// CHECK1-NEXT: [[CONV:%.*]] = sext i16 [[TMP0]] to i32841// CHECK1-NEXT: [[TMP1:%.*]] = load i32, ptr [[A_ADDR]], align 4842// CHECK1-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP1]], [[CONV]]843// CHECK1-NEXT: store i32 [[ADD]], ptr [[A_ADDR]], align 4844// CHECK1-NEXT: ret void845//846//847// CHECK3-LABEL: define {{[^@]+}}@_Z3bari848// CHECK3-SAME: (i32 noundef [[N:%.*]]) #[[ATTR0:[0-9]+]] {849// CHECK3-NEXT: entry:850// CHECK3-NEXT: [[N_ADDR:%.*]] = alloca i32, align 4851// CHECK3-NEXT: [[A:%.*]] = alloca i32, align 4852// CHECK3-NEXT: [[S:%.*]] = alloca [[STRUCT_S1:%.*]], align 4853// CHECK3-NEXT: store i32 [[N]], ptr [[N_ADDR]], align 4854// CHECK3-NEXT: store i32 0, ptr [[A]], align 4855// CHECK3-NEXT: [[TMP0:%.*]] = load i32, ptr [[N_ADDR]], align 4856// CHECK3-NEXT: [[CALL:%.*]] = call noundef i32 @_ZN2S12r1Ei(ptr noundef nonnull align 4 dereferenceable(8) [[S]], i32 noundef [[TMP0]])857// CHECK3-NEXT: [[TMP1:%.*]] = load i32, ptr [[A]], align 4858// CHECK3-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP1]], [[CALL]]859// CHECK3-NEXT: store i32 [[ADD]], ptr [[A]], align 4860// CHECK3-NEXT: [[TMP2:%.*]] = load i32, ptr [[N_ADDR]], align 4861// CHECK3-NEXT: [[CALL1:%.*]] = call noundef i32 @_ZL7fstatici(i32 noundef [[TMP2]])862// CHECK3-NEXT: [[TMP3:%.*]] = load i32, ptr [[A]], align 4863// CHECK3-NEXT: [[ADD2:%.*]] = add nsw i32 [[TMP3]], [[CALL1]]864// CHECK3-NEXT: store i32 [[ADD2]], ptr [[A]], align 4865// CHECK3-NEXT: [[TMP4:%.*]] = load i32, ptr [[N_ADDR]], align 4866// CHECK3-NEXT: [[CALL3:%.*]] = call noundef i32 @_Z9ftemplateIiET_i(i32 noundef [[TMP4]])867// CHECK3-NEXT: [[TMP5:%.*]] = load i32, ptr [[A]], align 4868// CHECK3-NEXT: [[ADD4:%.*]] = add nsw i32 [[TMP5]], [[CALL3]]869// CHECK3-NEXT: store i32 [[ADD4]], ptr [[A]], align 4870// CHECK3-NEXT: [[TMP6:%.*]] = load i32, ptr [[A]], align 4871// CHECK3-NEXT: ret i32 [[TMP6]]872//873//874// CHECK3-LABEL: define {{[^@]+}}@_ZN2S12r1Ei875// CHECK3-SAME: (ptr noundef nonnull align 4 dereferenceable(8) [[THIS:%.*]], i32 noundef [[N:%.*]]) #[[ATTR0]] comdat align 2 {876// CHECK3-NEXT: entry:877// CHECK3-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 4878// CHECK3-NEXT: [[N_ADDR:%.*]] = alloca i32, align 4879// CHECK3-NEXT: [[B:%.*]] = alloca i32, align 4880// CHECK3-NEXT: [[DOTCAPTURE_EXPR_:%.*]] = alloca i8, align 1881// CHECK3-NEXT: [[B_CASTED:%.*]] = alloca i32, align 4882// CHECK3-NEXT: [[DOTCAPTURE_EXPR__CASTED:%.*]] = alloca i32, align 4883// CHECK3-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [3 x ptr], align 4884// CHECK3-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [3 x ptr], align 4885// CHECK3-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [3 x ptr], align 4886// CHECK3-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8887// CHECK3-NEXT: [[DOTCAPTURE_EXPR_4:%.*]] = alloca i8, align 1888// CHECK3-NEXT: [[DOTCAPTURE_EXPR__CASTED8:%.*]] = alloca i32, align 4889// CHECK3-NEXT: [[DOTOFFLOAD_BASEPTRS12:%.*]] = alloca [2 x ptr], align 4890// CHECK3-NEXT: [[DOTOFFLOAD_PTRS13:%.*]] = alloca [2 x ptr], align 4891// CHECK3-NEXT: [[DOTOFFLOAD_MAPPERS14:%.*]] = alloca [2 x ptr], align 4892// CHECK3-NEXT: [[KERNEL_ARGS16:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 8893// CHECK3-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 4894// CHECK3-NEXT: store i32 [[N]], ptr [[N_ADDR]], align 4895// CHECK3-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 4896// CHECK3-NEXT: store i32 1, ptr [[B]], align 4897// CHECK3-NEXT: [[TMP0:%.*]] = load i32, ptr [[N_ADDR]], align 4898// CHECK3-NEXT: [[CMP:%.*]] = icmp sgt i32 [[TMP0]], 3899// CHECK3-NEXT: [[STOREDV:%.*]] = zext i1 [[CMP]] to i8900// CHECK3-NEXT: store i8 [[STOREDV]], ptr [[DOTCAPTURE_EXPR_]], align 1901// CHECK3-NEXT: [[TMP1:%.*]] = load i32, ptr [[B]], align 4902// CHECK3-NEXT: store i32 [[TMP1]], ptr [[B_CASTED]], align 4903// CHECK3-NEXT: [[TMP2:%.*]] = load i32, ptr [[B_CASTED]], align 4904// CHECK3-NEXT: [[TMP3:%.*]] = load i8, ptr [[DOTCAPTURE_EXPR_]], align 1905// CHECK3-NEXT: [[LOADEDV:%.*]] = trunc i8 [[TMP3]] to i1906// CHECK3-NEXT: [[STOREDV2:%.*]] = zext i1 [[LOADEDV]] to i8907// CHECK3-NEXT: store i8 [[STOREDV2]], ptr [[DOTCAPTURE_EXPR__CASTED]], align 1908// CHECK3-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR__CASTED]], align 4909// CHECK3-NEXT: [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_S1:%.*]], ptr [[THIS1]], i32 0, i32 0910// CHECK3-NEXT: [[TMP5:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0911// CHECK3-NEXT: store ptr [[THIS1]], ptr [[TMP5]], align 4912// CHECK3-NEXT: [[TMP6:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0913// CHECK3-NEXT: store ptr [[A]], ptr [[TMP6]], align 4914// CHECK3-NEXT: [[TMP7:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 0915// CHECK3-NEXT: store ptr null, ptr [[TMP7]], align 4916// CHECK3-NEXT: [[TMP8:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 1917// CHECK3-NEXT: store i32 [[TMP2]], ptr [[TMP8]], align 4918// CHECK3-NEXT: [[TMP9:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 1919// CHECK3-NEXT: store i32 [[TMP2]], ptr [[TMP9]], align 4920// CHECK3-NEXT: [[TMP10:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 1921// CHECK3-NEXT: store ptr null, ptr [[TMP10]], align 4922// CHECK3-NEXT: [[TMP11:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 2923// CHECK3-NEXT: store i32 [[TMP4]], ptr [[TMP11]], align 4924// CHECK3-NEXT: [[TMP12:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 2925// CHECK3-NEXT: store i32 [[TMP4]], ptr [[TMP12]], align 4926// CHECK3-NEXT: [[TMP13:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 2927// CHECK3-NEXT: store ptr null, ptr [[TMP13]], align 4928// CHECK3-NEXT: [[TMP14:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0929// CHECK3-NEXT: [[TMP15:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0930// CHECK3-NEXT: [[TMP16:%.*]] = load i8, ptr [[DOTCAPTURE_EXPR_]], align 1931// CHECK3-NEXT: [[LOADEDV3:%.*]] = trunc i8 [[TMP16]] to i1932// CHECK3-NEXT: [[TMP17:%.*]] = select i1 [[LOADEDV3]], i32 0, i32 1933// CHECK3-NEXT: [[TMP18:%.*]] = insertvalue [3 x i32] zeroinitializer, i32 [[TMP17]], 0934// CHECK3-NEXT: [[TMP19:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0935// CHECK3-NEXT: store i32 3, ptr [[TMP19]], align 4936// CHECK3-NEXT: [[TMP20:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1937// CHECK3-NEXT: store i32 3, ptr [[TMP20]], align 4938// CHECK3-NEXT: [[TMP21:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2939// CHECK3-NEXT: store ptr [[TMP14]], ptr [[TMP21]], align 4940// CHECK3-NEXT: [[TMP22:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3941// CHECK3-NEXT: store ptr [[TMP15]], ptr [[TMP22]], align 4942// CHECK3-NEXT: [[TMP23:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4943// CHECK3-NEXT: store ptr @.offload_sizes, ptr [[TMP23]], align 4944// CHECK3-NEXT: [[TMP24:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5945// CHECK3-NEXT: store ptr @.offload_maptypes, ptr [[TMP24]], align 4946// CHECK3-NEXT: [[TMP25:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6947// CHECK3-NEXT: store ptr null, ptr [[TMP25]], align 4948// CHECK3-NEXT: [[TMP26:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7949// CHECK3-NEXT: store ptr null, ptr [[TMP26]], align 4950// CHECK3-NEXT: [[TMP27:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8951// CHECK3-NEXT: store i64 0, ptr [[TMP27]], align 8952// CHECK3-NEXT: [[TMP28:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9953// CHECK3-NEXT: store i64 0, ptr [[TMP28]], align 8954// CHECK3-NEXT: [[TMP29:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10955// CHECK3-NEXT: store [3 x i32] [i32 1, i32 0, i32 0], ptr [[TMP29]], align 4956// CHECK3-NEXT: [[TMP30:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11957// CHECK3-NEXT: store [3 x i32] [[TMP18]], ptr [[TMP30]], align 4958// CHECK3-NEXT: [[TMP31:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12959// CHECK3-NEXT: store i32 0, ptr [[TMP31]], align 4960// CHECK3-NEXT: [[TMP32:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1:[0-9]+]], i64 -1, i32 1, i32 [[TMP17]], ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2S12r1Ei_l121.region_id, ptr [[KERNEL_ARGS]])961// CHECK3-NEXT: [[TMP33:%.*]] = icmp ne i32 [[TMP32]], 0962// CHECK3-NEXT: br i1 [[TMP33]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]963// CHECK3: omp_offload.failed:964// CHECK3-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2S12r1Ei_l121(ptr [[THIS1]], i32 [[TMP2]], i32 [[TMP4]]) #[[ATTR2:[0-9]+]]965// CHECK3-NEXT: br label [[OMP_OFFLOAD_CONT]]966// CHECK3: omp_offload.cont:967// CHECK3-NEXT: [[TMP34:%.*]] = load i32, ptr [[N_ADDR]], align 4968// CHECK3-NEXT: [[CMP5:%.*]] = icmp sgt i32 [[TMP34]], 5969// CHECK3-NEXT: [[STOREDV6:%.*]] = zext i1 [[CMP5]] to i8970// CHECK3-NEXT: store i8 [[STOREDV6]], ptr [[DOTCAPTURE_EXPR_4]], align 1971// CHECK3-NEXT: [[TMP35:%.*]] = load i8, ptr [[DOTCAPTURE_EXPR_4]], align 1972// CHECK3-NEXT: [[LOADEDV7:%.*]] = trunc i8 [[TMP35]] to i1973// CHECK3-NEXT: [[STOREDV9:%.*]] = zext i1 [[LOADEDV7]] to i8974// CHECK3-NEXT: store i8 [[STOREDV9]], ptr [[DOTCAPTURE_EXPR__CASTED8]], align 1975// CHECK3-NEXT: [[TMP36:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR__CASTED8]], align 4976// CHECK3-NEXT: [[TMP37:%.*]] = load i32, ptr [[N_ADDR]], align 4977// CHECK3-NEXT: [[CMP10:%.*]] = icmp sgt i32 [[TMP37]], 4978// CHECK3-NEXT: br i1 [[CMP10]], label [[OMP_IF_THEN:%.*]], label [[OMP_IF_ELSE:%.*]]979// CHECK3: omp_if.then:980// CHECK3-NEXT: [[A11:%.*]] = getelementptr inbounds nuw [[STRUCT_S1]], ptr [[THIS1]], i32 0, i32 0981// CHECK3-NEXT: [[TMP38:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS12]], i32 0, i32 0982// CHECK3-NEXT: store ptr [[THIS1]], ptr [[TMP38]], align 4983// CHECK3-NEXT: [[TMP39:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS13]], i32 0, i32 0984// CHECK3-NEXT: store ptr [[A11]], ptr [[TMP39]], align 4985// CHECK3-NEXT: [[TMP40:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS14]], i32 0, i32 0986// CHECK3-NEXT: store ptr null, ptr [[TMP40]], align 4987// CHECK3-NEXT: [[TMP41:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS12]], i32 0, i32 1988// CHECK3-NEXT: store i32 [[TMP36]], ptr [[TMP41]], align 4989// CHECK3-NEXT: [[TMP42:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS13]], i32 0, i32 1990// CHECK3-NEXT: store i32 [[TMP36]], ptr [[TMP42]], align 4991// CHECK3-NEXT: [[TMP43:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS14]], i32 0, i32 1992// CHECK3-NEXT: store ptr null, ptr [[TMP43]], align 4993// CHECK3-NEXT: [[TMP44:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS12]], i32 0, i32 0994// CHECK3-NEXT: [[TMP45:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS13]], i32 0, i32 0995// CHECK3-NEXT: [[TMP46:%.*]] = load i8, ptr [[DOTCAPTURE_EXPR_4]], align 1996// CHECK3-NEXT: [[LOADEDV15:%.*]] = trunc i8 [[TMP46]] to i1997// CHECK3-NEXT: [[TMP47:%.*]] = select i1 [[LOADEDV15]], i32 0, i32 1998// CHECK3-NEXT: [[TMP48:%.*]] = insertvalue [3 x i32] zeroinitializer, i32 [[TMP47]], 0999// CHECK3-NEXT: [[TMP49:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS16]], i32 0, i32 01000// CHECK3-NEXT: store i32 3, ptr [[TMP49]], align 41001// CHECK3-NEXT: [[TMP50:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS16]], i32 0, i32 11002// CHECK3-NEXT: store i32 2, ptr [[TMP50]], align 41003// CHECK3-NEXT: [[TMP51:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS16]], i32 0, i32 21004// CHECK3-NEXT: store ptr [[TMP44]], ptr [[TMP51]], align 41005// CHECK3-NEXT: [[TMP52:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS16]], i32 0, i32 31006// CHECK3-NEXT: store ptr [[TMP45]], ptr [[TMP52]], align 41007// CHECK3-NEXT: [[TMP53:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS16]], i32 0, i32 41008// CHECK3-NEXT: store ptr @.offload_sizes.1, ptr [[TMP53]], align 41009// CHECK3-NEXT: [[TMP54:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS16]], i32 0, i32 51010// CHECK3-NEXT: store ptr @.offload_maptypes.2, ptr [[TMP54]], align 41011// CHECK3-NEXT: [[TMP55:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS16]], i32 0, i32 61012// CHECK3-NEXT: store ptr null, ptr [[TMP55]], align 41013// CHECK3-NEXT: [[TMP56:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS16]], i32 0, i32 71014// CHECK3-NEXT: store ptr null, ptr [[TMP56]], align 41015// CHECK3-NEXT: [[TMP57:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS16]], i32 0, i32 81016// CHECK3-NEXT: store i64 0, ptr [[TMP57]], align 81017// CHECK3-NEXT: [[TMP58:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS16]], i32 0, i32 91018// CHECK3-NEXT: store i64 0, ptr [[TMP58]], align 81019// CHECK3-NEXT: [[TMP59:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS16]], i32 0, i32 101020// CHECK3-NEXT: store [3 x i32] [i32 1, i32 0, i32 0], ptr [[TMP59]], align 41021// CHECK3-NEXT: [[TMP60:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS16]], i32 0, i32 111022// CHECK3-NEXT: store [3 x i32] [[TMP48]], ptr [[TMP60]], align 41023// CHECK3-NEXT: [[TMP61:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS16]], i32 0, i32 121024// CHECK3-NEXT: store i32 0, ptr [[TMP61]], align 41025// CHECK3-NEXT: [[TMP62:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 -1, i32 1, i32 [[TMP47]], ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2S12r1Ei_l126.region_id, ptr [[KERNEL_ARGS16]])1026// CHECK3-NEXT: [[TMP63:%.*]] = icmp ne i32 [[TMP62]], 01027// CHECK3-NEXT: br i1 [[TMP63]], label [[OMP_OFFLOAD_FAILED17:%.*]], label [[OMP_OFFLOAD_CONT18:%.*]]1028// CHECK3: omp_offload.failed17:1029// CHECK3-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2S12r1Ei_l126(ptr [[THIS1]], i32 [[TMP36]]) #[[ATTR2]]1030// CHECK3-NEXT: br label [[OMP_OFFLOAD_CONT18]]1031// CHECK3: omp_offload.cont18:1032// CHECK3-NEXT: br label [[OMP_IF_END:%.*]]1033// CHECK3: omp_if.else:1034// CHECK3-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2S12r1Ei_l126(ptr [[THIS1]], i32 [[TMP36]]) #[[ATTR2]]1035// CHECK3-NEXT: br label [[OMP_IF_END]]1036// CHECK3: omp_if.end:1037// CHECK3-NEXT: [[A19:%.*]] = getelementptr inbounds nuw [[STRUCT_S1]], ptr [[THIS1]], i32 0, i32 01038// CHECK3-NEXT: [[TMP64:%.*]] = load double, ptr [[A19]], align 41039// CHECK3-NEXT: [[CONV:%.*]] = fptosi double [[TMP64]] to i321040// CHECK3-NEXT: ret i32 [[CONV]]1041//1042//1043// CHECK3-LABEL: define {{[^@]+}}@_ZL7fstatici1044// CHECK3-SAME: (i32 noundef [[N:%.*]]) #[[ATTR0]] {1045// CHECK3-NEXT: entry:1046// CHECK3-NEXT: [[N_ADDR:%.*]] = alloca i32, align 41047// CHECK3-NEXT: [[DOTCAPTURE_EXPR_:%.*]] = alloca i8, align 11048// CHECK3-NEXT: [[DOTCAPTURE_EXPR__CASTED:%.*]] = alloca i32, align 41049// CHECK3-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x ptr], align 41050// CHECK3-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x ptr], align 41051// CHECK3-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [1 x ptr], align 41052// CHECK3-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 81053// CHECK3-NEXT: [[KERNEL_ARGS6:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 81054// CHECK3-NEXT: store i32 [[N]], ptr [[N_ADDR]], align 41055// CHECK3-NEXT: [[TMP0:%.*]] = load i32, ptr [[N_ADDR]], align 41056// CHECK3-NEXT: [[CMP:%.*]] = icmp sgt i32 [[TMP0]], 11057// CHECK3-NEXT: [[STOREDV:%.*]] = zext i1 [[CMP]] to i81058// CHECK3-NEXT: store i8 [[STOREDV]], ptr [[DOTCAPTURE_EXPR_]], align 11059// CHECK3-NEXT: [[TMP1:%.*]] = load i8, ptr [[DOTCAPTURE_EXPR_]], align 11060// CHECK3-NEXT: [[LOADEDV:%.*]] = trunc i8 [[TMP1]] to i11061// CHECK3-NEXT: [[STOREDV1:%.*]] = zext i1 [[LOADEDV]] to i81062// CHECK3-NEXT: store i8 [[STOREDV1]], ptr [[DOTCAPTURE_EXPR__CASTED]], align 11063// CHECK3-NEXT: [[TMP2:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR__CASTED]], align 41064// CHECK3-NEXT: [[TMP3:%.*]] = load i8, ptr [[DOTCAPTURE_EXPR_]], align 11065// CHECK3-NEXT: [[LOADEDV2:%.*]] = trunc i8 [[TMP3]] to i11066// CHECK3-NEXT: br i1 [[LOADEDV2]], label [[OMP_IF_THEN:%.*]], label [[OMP_IF_ELSE:%.*]]1067// CHECK3: omp_if.then:1068// CHECK3-NEXT: [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 01069// CHECK3-NEXT: store i32 [[TMP2]], ptr [[TMP4]], align 41070// CHECK3-NEXT: [[TMP5:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 01071// CHECK3-NEXT: store i32 [[TMP2]], ptr [[TMP5]], align 41072// CHECK3-NEXT: [[TMP6:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 01073// CHECK3-NEXT: store ptr null, ptr [[TMP6]], align 41074// CHECK3-NEXT: [[TMP7:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 01075// CHECK3-NEXT: [[TMP8:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 01076// CHECK3-NEXT: [[TMP9:%.*]] = load i8, ptr [[DOTCAPTURE_EXPR_]], align 11077// CHECK3-NEXT: [[LOADEDV3:%.*]] = trunc i8 [[TMP9]] to i11078// CHECK3-NEXT: [[TMP10:%.*]] = select i1 [[LOADEDV3]], i32 0, i32 11079// CHECK3-NEXT: [[TMP11:%.*]] = insertvalue [3 x i32] zeroinitializer, i32 [[TMP10]], 01080// CHECK3-NEXT: [[TMP12:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 01081// CHECK3-NEXT: store i32 3, ptr [[TMP12]], align 41082// CHECK3-NEXT: [[TMP13:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11083// CHECK3-NEXT: store i32 1, ptr [[TMP13]], align 41084// CHECK3-NEXT: [[TMP14:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 21085// CHECK3-NEXT: store ptr [[TMP7]], ptr [[TMP14]], align 41086// CHECK3-NEXT: [[TMP15:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 31087// CHECK3-NEXT: store ptr [[TMP8]], ptr [[TMP15]], align 41088// CHECK3-NEXT: [[TMP16:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 41089// CHECK3-NEXT: store ptr @.offload_sizes.3, ptr [[TMP16]], align 41090// CHECK3-NEXT: [[TMP17:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 51091// CHECK3-NEXT: store ptr @.offload_maptypes.4, ptr [[TMP17]], align 41092// CHECK3-NEXT: [[TMP18:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 61093// CHECK3-NEXT: store ptr null, ptr [[TMP18]], align 41094// CHECK3-NEXT: [[TMP19:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 71095// CHECK3-NEXT: store ptr null, ptr [[TMP19]], align 41096// CHECK3-NEXT: [[TMP20:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 81097// CHECK3-NEXT: store i64 0, ptr [[TMP20]], align 81098// CHECK3-NEXT: [[TMP21:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 91099// CHECK3-NEXT: store i64 0, ptr [[TMP21]], align 81100// CHECK3-NEXT: [[TMP22:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 101101// CHECK3-NEXT: store [3 x i32] [i32 1, i32 0, i32 0], ptr [[TMP22]], align 41102// CHECK3-NEXT: [[TMP23:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 111103// CHECK3-NEXT: store [3 x i32] [[TMP11]], ptr [[TMP23]], align 41104// CHECK3-NEXT: [[TMP24:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 121105// CHECK3-NEXT: store i32 0, ptr [[TMP24]], align 41106// CHECK3-NEXT: [[TMP25:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 -1, i32 1, i32 [[TMP10]], ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l104.region_id, ptr [[KERNEL_ARGS]])1107// CHECK3-NEXT: [[TMP26:%.*]] = icmp ne i32 [[TMP25]], 01108// CHECK3-NEXT: br i1 [[TMP26]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]1109// CHECK3: omp_offload.failed:1110// CHECK3-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l104(i32 [[TMP2]]) #[[ATTR2]]1111// CHECK3-NEXT: br label [[OMP_OFFLOAD_CONT]]1112// CHECK3: omp_offload.cont:1113// CHECK3-NEXT: br label [[OMP_IF_END:%.*]]1114// CHECK3: omp_if.else:1115// CHECK3-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l104(i32 [[TMP2]]) #[[ATTR2]]1116// CHECK3-NEXT: br label [[OMP_IF_END]]1117// CHECK3: omp_if.end:1118// CHECK3-NEXT: [[TMP27:%.*]] = load i32, ptr [[N_ADDR]], align 41119// CHECK3-NEXT: [[SUB:%.*]] = sub nsw i32 [[TMP27]], 21120// CHECK3-NEXT: [[CMP4:%.*]] = icmp sgt i32 [[SUB]], 21121// CHECK3-NEXT: br i1 [[CMP4]], label [[OMP_IF_THEN5:%.*]], label [[OMP_IF_ELSE9:%.*]]1122// CHECK3: omp_if.then5:1123// CHECK3-NEXT: [[TMP28:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 01124// CHECK3-NEXT: store i32 3, ptr [[TMP28]], align 41125// CHECK3-NEXT: [[TMP29:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 11126// CHECK3-NEXT: store i32 0, ptr [[TMP29]], align 41127// CHECK3-NEXT: [[TMP30:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 21128// CHECK3-NEXT: store ptr null, ptr [[TMP30]], align 41129// CHECK3-NEXT: [[TMP31:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 31130// CHECK3-NEXT: store ptr null, ptr [[TMP31]], align 41131// CHECK3-NEXT: [[TMP32:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 41132// CHECK3-NEXT: store ptr null, ptr [[TMP32]], align 41133// CHECK3-NEXT: [[TMP33:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 51134// CHECK3-NEXT: store ptr null, ptr [[TMP33]], align 41135// CHECK3-NEXT: [[TMP34:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 61136// CHECK3-NEXT: store ptr null, ptr [[TMP34]], align 41137// CHECK3-NEXT: [[TMP35:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 71138// CHECK3-NEXT: store ptr null, ptr [[TMP35]], align 41139// CHECK3-NEXT: [[TMP36:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 81140// CHECK3-NEXT: store i64 0, ptr [[TMP36]], align 81141// CHECK3-NEXT: [[TMP37:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 91142// CHECK3-NEXT: store i64 0, ptr [[TMP37]], align 81143// CHECK3-NEXT: [[TMP38:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 101144// CHECK3-NEXT: store [3 x i32] [i32 1, i32 0, i32 0], ptr [[TMP38]], align 41145// CHECK3-NEXT: [[TMP39:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 111146// CHECK3-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP39]], align 41147// CHECK3-NEXT: [[TMP40:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 121148// CHECK3-NEXT: store i32 0, ptr [[TMP40]], align 41149// CHECK3-NEXT: [[TMP41:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 -1, i32 1, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l108.region_id, ptr [[KERNEL_ARGS6]])1150// CHECK3-NEXT: [[TMP42:%.*]] = icmp ne i32 [[TMP41]], 01151// CHECK3-NEXT: br i1 [[TMP42]], label [[OMP_OFFLOAD_FAILED7:%.*]], label [[OMP_OFFLOAD_CONT8:%.*]]1152// CHECK3: omp_offload.failed7:1153// CHECK3-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l108() #[[ATTR2]]1154// CHECK3-NEXT: br label [[OMP_OFFLOAD_CONT8]]1155// CHECK3: omp_offload.cont8:1156// CHECK3-NEXT: br label [[OMP_IF_END10:%.*]]1157// CHECK3: omp_if.else9:1158// CHECK3-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l108() #[[ATTR2]]1159// CHECK3-NEXT: br label [[OMP_IF_END10]]1160// CHECK3: omp_if.end10:1161// CHECK3-NEXT: [[TMP43:%.*]] = load i32, ptr [[N_ADDR]], align 41162// CHECK3-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP43]], 11163// CHECK3-NEXT: ret i32 [[ADD]]1164//1165//1166// CHECK3-LABEL: define {{[^@]+}}@_Z9ftemplateIiET_i1167// CHECK3-SAME: (i32 noundef [[N:%.*]]) #[[ATTR0]] comdat {1168// CHECK3-NEXT: entry:1169// CHECK3-NEXT: [[N_ADDR:%.*]] = alloca i32, align 41170// CHECK3-NEXT: [[A:%.*]] = alloca i32, align 41171// CHECK3-NEXT: [[A_CASTED:%.*]] = alloca i32, align 41172// CHECK3-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x ptr], align 41173// CHECK3-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x ptr], align 41174// CHECK3-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [1 x ptr], align 41175// CHECK3-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 81176// CHECK3-NEXT: [[B:%.*]] = alloca i16, align 21177// CHECK3-NEXT: [[A_CASTED1:%.*]] = alloca i32, align 41178// CHECK3-NEXT: [[B_CASTED:%.*]] = alloca i32, align 41179// CHECK3-NEXT: [[DOTOFFLOAD_BASEPTRS2:%.*]] = alloca [2 x ptr], align 41180// CHECK3-NEXT: [[DOTOFFLOAD_PTRS3:%.*]] = alloca [2 x ptr], align 41181// CHECK3-NEXT: [[DOTOFFLOAD_MAPPERS4:%.*]] = alloca [2 x ptr], align 41182// CHECK3-NEXT: [[KERNEL_ARGS5:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 81183// CHECK3-NEXT: store i32 [[N]], ptr [[N_ADDR]], align 41184// CHECK3-NEXT: store i32 0, ptr [[A]], align 41185// CHECK3-NEXT: [[TMP0:%.*]] = load i32, ptr [[A]], align 41186// CHECK3-NEXT: store i32 [[TMP0]], ptr [[A_CASTED]], align 41187// CHECK3-NEXT: [[TMP1:%.*]] = load i32, ptr [[A_CASTED]], align 41188// CHECK3-NEXT: [[TMP2:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 01189// CHECK3-NEXT: store i32 [[TMP1]], ptr [[TMP2]], align 41190// CHECK3-NEXT: [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 01191// CHECK3-NEXT: store i32 [[TMP1]], ptr [[TMP3]], align 41192// CHECK3-NEXT: [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 01193// CHECK3-NEXT: store ptr null, ptr [[TMP4]], align 41194// CHECK3-NEXT: [[TMP5:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 01195// CHECK3-NEXT: [[TMP6:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 01196// CHECK3-NEXT: [[TMP7:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 01197// CHECK3-NEXT: store i32 3, ptr [[TMP7]], align 41198// CHECK3-NEXT: [[TMP8:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11199// CHECK3-NEXT: store i32 1, ptr [[TMP8]], align 41200// CHECK3-NEXT: [[TMP9:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 21201// CHECK3-NEXT: store ptr [[TMP5]], ptr [[TMP9]], align 41202// CHECK3-NEXT: [[TMP10:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 31203// CHECK3-NEXT: store ptr [[TMP6]], ptr [[TMP10]], align 41204// CHECK3-NEXT: [[TMP11:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 41205// CHECK3-NEXT: store ptr @.offload_sizes.5, ptr [[TMP11]], align 41206// CHECK3-NEXT: [[TMP12:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 51207// CHECK3-NEXT: store ptr @.offload_maptypes.6, ptr [[TMP12]], align 41208// CHECK3-NEXT: [[TMP13:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 61209// CHECK3-NEXT: store ptr null, ptr [[TMP13]], align 41210// CHECK3-NEXT: [[TMP14:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 71211// CHECK3-NEXT: store ptr null, ptr [[TMP14]], align 41212// CHECK3-NEXT: [[TMP15:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 81213// CHECK3-NEXT: store i64 0, ptr [[TMP15]], align 81214// CHECK3-NEXT: [[TMP16:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 91215// CHECK3-NEXT: store i64 0, ptr [[TMP16]], align 81216// CHECK3-NEXT: [[TMP17:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 101217// CHECK3-NEXT: store [3 x i32] [i32 1, i32 0, i32 0], ptr [[TMP17]], align 41218// CHECK3-NEXT: [[TMP18:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 111219// CHECK3-NEXT: store [3 x i32] [i32 1, i32 0, i32 0], ptr [[TMP18]], align 41220// CHECK3-NEXT: [[TMP19:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 121221// CHECK3-NEXT: store i32 0, ptr [[TMP19]], align 41222// CHECK3-NEXT: [[TMP20:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 -1, i32 1, i32 1, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l87.region_id, ptr [[KERNEL_ARGS]])1223// CHECK3-NEXT: [[TMP21:%.*]] = icmp ne i32 [[TMP20]], 01224// CHECK3-NEXT: br i1 [[TMP21]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]1225// CHECK3: omp_offload.failed:1226// CHECK3-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l87(i32 [[TMP1]]) #[[ATTR2]]1227// CHECK3-NEXT: br label [[OMP_OFFLOAD_CONT]]1228// CHECK3: omp_offload.cont:1229// CHECK3-NEXT: store i16 1, ptr [[B]], align 21230// CHECK3-NEXT: [[TMP22:%.*]] = load i32, ptr [[A]], align 41231// CHECK3-NEXT: store i32 [[TMP22]], ptr [[A_CASTED1]], align 41232// CHECK3-NEXT: [[TMP23:%.*]] = load i32, ptr [[A_CASTED1]], align 41233// CHECK3-NEXT: [[TMP24:%.*]] = load i16, ptr [[B]], align 21234// CHECK3-NEXT: store i16 [[TMP24]], ptr [[B_CASTED]], align 21235// CHECK3-NEXT: [[TMP25:%.*]] = load i32, ptr [[B_CASTED]], align 41236// CHECK3-NEXT: [[TMP26:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS2]], i32 0, i32 01237// CHECK3-NEXT: store i32 [[TMP23]], ptr [[TMP26]], align 41238// CHECK3-NEXT: [[TMP27:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS3]], i32 0, i32 01239// CHECK3-NEXT: store i32 [[TMP23]], ptr [[TMP27]], align 41240// CHECK3-NEXT: [[TMP28:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS4]], i32 0, i32 01241// CHECK3-NEXT: store ptr null, ptr [[TMP28]], align 41242// CHECK3-NEXT: [[TMP29:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS2]], i32 0, i32 11243// CHECK3-NEXT: store i32 [[TMP25]], ptr [[TMP29]], align 41244// CHECK3-NEXT: [[TMP30:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS3]], i32 0, i32 11245// CHECK3-NEXT: store i32 [[TMP25]], ptr [[TMP30]], align 41246// CHECK3-NEXT: [[TMP31:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS4]], i32 0, i32 11247// CHECK3-NEXT: store ptr null, ptr [[TMP31]], align 41248// CHECK3-NEXT: [[TMP32:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS2]], i32 0, i32 01249// CHECK3-NEXT: [[TMP33:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS3]], i32 0, i32 01250// CHECK3-NEXT: [[TMP34:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 01251// CHECK3-NEXT: store i32 3, ptr [[TMP34]], align 41252// CHECK3-NEXT: [[TMP35:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 11253// CHECK3-NEXT: store i32 2, ptr [[TMP35]], align 41254// CHECK3-NEXT: [[TMP36:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 21255// CHECK3-NEXT: store ptr [[TMP32]], ptr [[TMP36]], align 41256// CHECK3-NEXT: [[TMP37:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 31257// CHECK3-NEXT: store ptr [[TMP33]], ptr [[TMP37]], align 41258// CHECK3-NEXT: [[TMP38:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 41259// CHECK3-NEXT: store ptr @.offload_sizes.7, ptr [[TMP38]], align 41260// CHECK3-NEXT: [[TMP39:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 51261// CHECK3-NEXT: store ptr @.offload_maptypes.8, ptr [[TMP39]], align 41262// CHECK3-NEXT: [[TMP40:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 61263// CHECK3-NEXT: store ptr null, ptr [[TMP40]], align 41264// CHECK3-NEXT: [[TMP41:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 71265// CHECK3-NEXT: store ptr null, ptr [[TMP41]], align 41266// CHECK3-NEXT: [[TMP42:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 81267// CHECK3-NEXT: store i64 0, ptr [[TMP42]], align 81268// CHECK3-NEXT: [[TMP43:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 91269// CHECK3-NEXT: store i64 0, ptr [[TMP43]], align 81270// CHECK3-NEXT: [[TMP44:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 101271// CHECK3-NEXT: store [3 x i32] [i32 1, i32 0, i32 0], ptr [[TMP44]], align 41272// CHECK3-NEXT: [[TMP45:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 111273// CHECK3-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP45]], align 41274// CHECK3-NEXT: [[TMP46:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 121275// CHECK3-NEXT: store i32 0, ptr [[TMP46]], align 41276// CHECK3-NEXT: [[TMP47:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 -1, i32 1, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l93.region_id, ptr [[KERNEL_ARGS5]])1277// CHECK3-NEXT: [[TMP48:%.*]] = icmp ne i32 [[TMP47]], 01278// CHECK3-NEXT: br i1 [[TMP48]], label [[OMP_OFFLOAD_FAILED6:%.*]], label [[OMP_OFFLOAD_CONT7:%.*]]1279// CHECK3: omp_offload.failed6:1280// CHECK3-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l93(i32 [[TMP23]], i32 [[TMP25]]) #[[ATTR2]]1281// CHECK3-NEXT: br label [[OMP_OFFLOAD_CONT7]]1282// CHECK3: omp_offload.cont7:1283// CHECK3-NEXT: [[TMP49:%.*]] = load i32, ptr [[A]], align 41284// CHECK3-NEXT: ret i32 [[TMP49]]1285//1286//1287// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2S12r1Ei_l1211288// CHECK3-SAME: (ptr noundef [[THIS:%.*]], i32 noundef [[B:%.*]], i32 noundef [[DOTCAPTURE_EXPR_:%.*]]) #[[ATTR1:[0-9]+]] {1289// CHECK3-NEXT: entry:1290// CHECK3-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 41291// CHECK3-NEXT: [[B_ADDR:%.*]] = alloca i32, align 41292// CHECK3-NEXT: [[DOTCAPTURE_EXPR__ADDR:%.*]] = alloca i32, align 41293// CHECK3-NEXT: [[B_CASTED:%.*]] = alloca i32, align 41294// CHECK3-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 41295// CHECK3-NEXT: [[DOTBOUND_ZERO_ADDR:%.*]] = alloca i32, align 41296// CHECK3-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])1297// CHECK3-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 41298// CHECK3-NEXT: store i32 [[B]], ptr [[B_ADDR]], align 41299// CHECK3-NEXT: store i32 [[DOTCAPTURE_EXPR_]], ptr [[DOTCAPTURE_EXPR__ADDR]], align 41300// CHECK3-NEXT: [[TMP1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 41301// CHECK3-NEXT: [[TMP2:%.*]] = load i32, ptr [[B_ADDR]], align 41302// CHECK3-NEXT: store i32 [[TMP2]], ptr [[B_CASTED]], align 41303// CHECK3-NEXT: [[TMP3:%.*]] = load i32, ptr [[B_CASTED]], align 41304// CHECK3-NEXT: [[TMP4:%.*]] = load i8, ptr [[DOTCAPTURE_EXPR__ADDR]], align 11305// CHECK3-NEXT: [[LOADEDV:%.*]] = trunc i8 [[TMP4]] to i11306// CHECK3-NEXT: br i1 [[LOADEDV]], label [[OMP_IF_THEN:%.*]], label [[OMP_IF_ELSE:%.*]]1307// CHECK3: omp_if.then:1308// CHECK3-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB1]], i32 2, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2S12r1Ei_l121.omp_outlined, ptr [[TMP1]], i32 [[TMP3]])1309// CHECK3-NEXT: br label [[OMP_IF_END:%.*]]1310// CHECK3: omp_if.else:1311// CHECK3-NEXT: call void @__kmpc_serialized_parallel(ptr @[[GLOB1]], i32 [[TMP0]])1312// CHECK3-NEXT: store i32 [[TMP0]], ptr [[DOTTHREADID_TEMP_]], align 41313// CHECK3-NEXT: store i32 0, ptr [[DOTBOUND_ZERO_ADDR]], align 41314// CHECK3-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2S12r1Ei_l121.omp_outlined(ptr [[DOTTHREADID_TEMP_]], ptr [[DOTBOUND_ZERO_ADDR]], ptr [[TMP1]], i32 [[TMP3]]) #[[ATTR2]]1315// CHECK3-NEXT: call void @__kmpc_end_serialized_parallel(ptr @[[GLOB1]], i32 [[TMP0]])1316// CHECK3-NEXT: br label [[OMP_IF_END]]1317// CHECK3: omp_if.end:1318// CHECK3-NEXT: ret void1319//1320//1321// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2S12r1Ei_l121.omp_outlined1322// CHECK3-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef [[THIS:%.*]], i32 noundef [[B:%.*]]) #[[ATTR1]] {1323// CHECK3-NEXT: entry:1324// CHECK3-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 41325// CHECK3-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 41326// CHECK3-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 41327// CHECK3-NEXT: [[B_ADDR:%.*]] = alloca i32, align 41328// CHECK3-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 41329// CHECK3-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 41330// CHECK3-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 41331// CHECK3-NEXT: store i32 [[B]], ptr [[B_ADDR]], align 41332// CHECK3-NEXT: [[TMP0:%.*]] = load ptr, ptr [[THIS_ADDR]], align 41333// CHECK3-NEXT: [[TMP1:%.*]] = load i32, ptr [[B_ADDR]], align 41334// CHECK3-NEXT: [[CONV:%.*]] = sitofp i32 [[TMP1]] to double1335// CHECK3-NEXT: [[ADD:%.*]] = fadd double [[CONV]], 1.500000e+001336// CHECK3-NEXT: [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_S1:%.*]], ptr [[TMP0]], i32 0, i32 01337// CHECK3-NEXT: store double [[ADD]], ptr [[A]], align 41338// CHECK3-NEXT: ret void1339//1340//1341// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2S12r1Ei_l1261342// CHECK3-SAME: (ptr noundef [[THIS:%.*]], i32 noundef [[DOTCAPTURE_EXPR_:%.*]]) #[[ATTR1]] {1343// CHECK3-NEXT: entry:1344// CHECK3-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 41345// CHECK3-NEXT: [[DOTCAPTURE_EXPR__ADDR:%.*]] = alloca i32, align 41346// CHECK3-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 41347// CHECK3-NEXT: [[DOTBOUND_ZERO_ADDR:%.*]] = alloca i32, align 41348// CHECK3-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])1349// CHECK3-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 41350// CHECK3-NEXT: store i32 [[DOTCAPTURE_EXPR_]], ptr [[DOTCAPTURE_EXPR__ADDR]], align 41351// CHECK3-NEXT: [[TMP1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 41352// CHECK3-NEXT: [[TMP2:%.*]] = load i8, ptr [[DOTCAPTURE_EXPR__ADDR]], align 11353// CHECK3-NEXT: [[LOADEDV:%.*]] = trunc i8 [[TMP2]] to i11354// CHECK3-NEXT: br i1 [[LOADEDV]], label [[OMP_IF_THEN:%.*]], label [[OMP_IF_ELSE:%.*]]1355// CHECK3: omp_if.then:1356// CHECK3-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB1]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2S12r1Ei_l126.omp_outlined, ptr [[TMP1]])1357// CHECK3-NEXT: br label [[OMP_IF_END:%.*]]1358// CHECK3: omp_if.else:1359// CHECK3-NEXT: call void @__kmpc_serialized_parallel(ptr @[[GLOB1]], i32 [[TMP0]])1360// CHECK3-NEXT: store i32 [[TMP0]], ptr [[DOTTHREADID_TEMP_]], align 41361// CHECK3-NEXT: store i32 0, ptr [[DOTBOUND_ZERO_ADDR]], align 41362// CHECK3-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2S12r1Ei_l126.omp_outlined(ptr [[DOTTHREADID_TEMP_]], ptr [[DOTBOUND_ZERO_ADDR]], ptr [[TMP1]]) #[[ATTR2]]1363// CHECK3-NEXT: call void @__kmpc_end_serialized_parallel(ptr @[[GLOB1]], i32 [[TMP0]])1364// CHECK3-NEXT: br label [[OMP_IF_END]]1365// CHECK3: omp_if.end:1366// CHECK3-NEXT: ret void1367//1368//1369// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2S12r1Ei_l126.omp_outlined1370// CHECK3-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef [[THIS:%.*]]) #[[ATTR1]] {1371// CHECK3-NEXT: entry:1372// CHECK3-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 41373// CHECK3-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 41374// CHECK3-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 41375// CHECK3-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 41376// CHECK3-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 41377// CHECK3-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 41378// CHECK3-NEXT: [[TMP0:%.*]] = load ptr, ptr [[THIS_ADDR]], align 41379// CHECK3-NEXT: [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_S1:%.*]], ptr [[TMP0]], i32 0, i32 01380// CHECK3-NEXT: store double 2.500000e+00, ptr [[A]], align 41381// CHECK3-NEXT: ret void1382//1383//1384// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l1041385// CHECK3-SAME: (i32 noundef [[DOTCAPTURE_EXPR_:%.*]]) #[[ATTR1]] {1386// CHECK3-NEXT: entry:1387// CHECK3-NEXT: [[DOTCAPTURE_EXPR__ADDR:%.*]] = alloca i32, align 41388// CHECK3-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 41389// CHECK3-NEXT: [[DOTBOUND_ZERO_ADDR:%.*]] = alloca i32, align 41390// CHECK3-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])1391// CHECK3-NEXT: store i32 [[DOTCAPTURE_EXPR_]], ptr [[DOTCAPTURE_EXPR__ADDR]], align 41392// CHECK3-NEXT: [[TMP1:%.*]] = load i8, ptr [[DOTCAPTURE_EXPR__ADDR]], align 11393// CHECK3-NEXT: [[LOADEDV:%.*]] = trunc i8 [[TMP1]] to i11394// CHECK3-NEXT: br i1 [[LOADEDV]], label [[OMP_IF_THEN:%.*]], label [[OMP_IF_ELSE:%.*]]1395// CHECK3: omp_if.then:1396// CHECK3-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB1]], i32 0, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l104.omp_outlined)1397// CHECK3-NEXT: br label [[OMP_IF_END:%.*]]1398// CHECK3: omp_if.else:1399// CHECK3-NEXT: call void @__kmpc_serialized_parallel(ptr @[[GLOB1]], i32 [[TMP0]])1400// CHECK3-NEXT: store i32 [[TMP0]], ptr [[DOTTHREADID_TEMP_]], align 41401// CHECK3-NEXT: store i32 0, ptr [[DOTBOUND_ZERO_ADDR]], align 41402// CHECK3-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l104.omp_outlined(ptr [[DOTTHREADID_TEMP_]], ptr [[DOTBOUND_ZERO_ADDR]]) #[[ATTR2]]1403// CHECK3-NEXT: call void @__kmpc_end_serialized_parallel(ptr @[[GLOB1]], i32 [[TMP0]])1404// CHECK3-NEXT: br label [[OMP_IF_END]]1405// CHECK3: omp_if.end:1406// CHECK3-NEXT: ret void1407//1408//1409// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l104.omp_outlined1410// CHECK3-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR1]] {1411// CHECK3-NEXT: entry:1412// CHECK3-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 41413// CHECK3-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 41414// CHECK3-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 41415// CHECK3-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 41416// CHECK3-NEXT: ret void1417//1418//1419// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l1081420// CHECK3-SAME: () #[[ATTR1]] {1421// CHECK3-NEXT: entry:1422// CHECK3-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB1]], i32 0, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l108.omp_outlined)1423// CHECK3-NEXT: ret void1424//1425//1426// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l108.omp_outlined1427// CHECK3-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR1]] {1428// CHECK3-NEXT: entry:1429// CHECK3-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 41430// CHECK3-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 41431// CHECK3-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 41432// CHECK3-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 41433// CHECK3-NEXT: ret void1434//1435//1436// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l871437// CHECK3-SAME: (i32 noundef [[A:%.*]]) #[[ATTR1]] {1438// CHECK3-NEXT: entry:1439// CHECK3-NEXT: [[A_ADDR:%.*]] = alloca i32, align 41440// CHECK3-NEXT: [[A_CASTED:%.*]] = alloca i32, align 41441// CHECK3-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 41442// CHECK3-NEXT: [[DOTBOUND_ZERO_ADDR:%.*]] = alloca i32, align 41443// CHECK3-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])1444// CHECK3-NEXT: store i32 [[A]], ptr [[A_ADDR]], align 41445// CHECK3-NEXT: [[TMP1:%.*]] = load i32, ptr [[A_ADDR]], align 41446// CHECK3-NEXT: store i32 [[TMP1]], ptr [[A_CASTED]], align 41447// CHECK3-NEXT: [[TMP2:%.*]] = load i32, ptr [[A_CASTED]], align 41448// CHECK3-NEXT: call void @__kmpc_serialized_parallel(ptr @[[GLOB1]], i32 [[TMP0]])1449// CHECK3-NEXT: store i32 [[TMP0]], ptr [[DOTTHREADID_TEMP_]], align 41450// CHECK3-NEXT: store i32 0, ptr [[DOTBOUND_ZERO_ADDR]], align 41451// CHECK3-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l87.omp_outlined(ptr [[DOTTHREADID_TEMP_]], ptr [[DOTBOUND_ZERO_ADDR]], i32 [[TMP2]]) #[[ATTR2]]1452// CHECK3-NEXT: call void @__kmpc_end_serialized_parallel(ptr @[[GLOB1]], i32 [[TMP0]])1453// CHECK3-NEXT: ret void1454//1455//1456// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l87.omp_outlined1457// CHECK3-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i32 noundef [[A:%.*]]) #[[ATTR1]] {1458// CHECK3-NEXT: entry:1459// CHECK3-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 41460// CHECK3-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 41461// CHECK3-NEXT: [[A_ADDR:%.*]] = alloca i32, align 41462// CHECK3-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 41463// CHECK3-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 41464// CHECK3-NEXT: store i32 [[A]], ptr [[A_ADDR]], align 41465// CHECK3-NEXT: [[TMP0:%.*]] = load i32, ptr [[A_ADDR]], align 41466// CHECK3-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP0]], 11467// CHECK3-NEXT: store i32 [[ADD]], ptr [[A_ADDR]], align 41468// CHECK3-NEXT: ret void1469//1470//1471// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l931472// CHECK3-SAME: (i32 noundef [[A:%.*]], i32 noundef [[B:%.*]]) #[[ATTR1]] {1473// CHECK3-NEXT: entry:1474// CHECK3-NEXT: [[A_ADDR:%.*]] = alloca i32, align 41475// CHECK3-NEXT: [[B_ADDR:%.*]] = alloca i32, align 41476// CHECK3-NEXT: [[A_CASTED:%.*]] = alloca i32, align 41477// CHECK3-NEXT: [[B_CASTED:%.*]] = alloca i32, align 41478// CHECK3-NEXT: store i32 [[A]], ptr [[A_ADDR]], align 41479// CHECK3-NEXT: store i32 [[B]], ptr [[B_ADDR]], align 41480// CHECK3-NEXT: [[TMP0:%.*]] = load i32, ptr [[A_ADDR]], align 41481// CHECK3-NEXT: store i32 [[TMP0]], ptr [[A_CASTED]], align 41482// CHECK3-NEXT: [[TMP1:%.*]] = load i32, ptr [[A_CASTED]], align 41483// CHECK3-NEXT: [[TMP2:%.*]] = load i16, ptr [[B_ADDR]], align 21484// CHECK3-NEXT: store i16 [[TMP2]], ptr [[B_CASTED]], align 21485// CHECK3-NEXT: [[TMP3:%.*]] = load i32, ptr [[B_CASTED]], align 41486// CHECK3-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB1]], i32 2, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l93.omp_outlined, i32 [[TMP1]], i32 [[TMP3]])1487// CHECK3-NEXT: ret void1488//1489//1490// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l93.omp_outlined1491// CHECK3-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i32 noundef [[A:%.*]], i32 noundef [[B:%.*]]) #[[ATTR1]] {1492// CHECK3-NEXT: entry:1493// CHECK3-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 41494// CHECK3-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 41495// CHECK3-NEXT: [[A_ADDR:%.*]] = alloca i32, align 41496// CHECK3-NEXT: [[B_ADDR:%.*]] = alloca i32, align 41497// CHECK3-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 41498// CHECK3-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 41499// CHECK3-NEXT: store i32 [[A]], ptr [[A_ADDR]], align 41500// CHECK3-NEXT: store i32 [[B]], ptr [[B_ADDR]], align 41501// CHECK3-NEXT: [[TMP0:%.*]] = load i16, ptr [[B_ADDR]], align 21502// CHECK3-NEXT: [[CONV:%.*]] = sext i16 [[TMP0]] to i321503// CHECK3-NEXT: [[TMP1:%.*]] = load i32, ptr [[A_ADDR]], align 41504// CHECK3-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP1]], [[CONV]]1505// CHECK3-NEXT: store i32 [[ADD]], ptr [[A_ADDR]], align 41506// CHECK3-NEXT: ret void1507//1508//1509// CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l1041510// CHECK9-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], i64 noundef [[DOTCAPTURE_EXPR_:%.*]]) #[[ATTR0:[0-9]+]] {1511// CHECK9-NEXT: entry:1512// CHECK9-NEXT: [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 81513// CHECK9-NEXT: [[DOTCAPTURE_EXPR__ADDR:%.*]] = alloca i64, align 81514// CHECK9-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 41515// CHECK9-NEXT: [[DOTBOUND_ZERO_ADDR:%.*]] = alloca i32, align 41516// CHECK9-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1:[0-9]+]])1517// CHECK9-NEXT: store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 81518// CHECK9-NEXT: store i64 [[DOTCAPTURE_EXPR_]], ptr [[DOTCAPTURE_EXPR__ADDR]], align 81519// CHECK9-NEXT: [[TMP1:%.*]] = load i8, ptr [[DOTCAPTURE_EXPR__ADDR]], align 11520// CHECK9-NEXT: [[LOADEDV:%.*]] = trunc i8 [[TMP1]] to i11521// CHECK9-NEXT: br i1 [[LOADEDV]], label [[OMP_IF_THEN:%.*]], label [[OMP_IF_ELSE:%.*]]1522// CHECK9: omp_if.then:1523// CHECK9-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB1]], i32 0, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l104.omp_outlined)1524// CHECK9-NEXT: br label [[OMP_IF_END:%.*]]1525// CHECK9: omp_if.else:1526// CHECK9-NEXT: call void @__kmpc_serialized_parallel(ptr @[[GLOB1]], i32 [[TMP0]])1527// CHECK9-NEXT: store i32 [[TMP0]], ptr [[DOTTHREADID_TEMP_]], align 41528// CHECK9-NEXT: store i32 0, ptr [[DOTBOUND_ZERO_ADDR]], align 41529// CHECK9-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l104.omp_outlined(ptr [[DOTTHREADID_TEMP_]], ptr [[DOTBOUND_ZERO_ADDR]]) #[[ATTR1:[0-9]+]]1530// CHECK9-NEXT: call void @__kmpc_end_serialized_parallel(ptr @[[GLOB1]], i32 [[TMP0]])1531// CHECK9-NEXT: br label [[OMP_IF_END]]1532// CHECK9: omp_if.end:1533// CHECK9-NEXT: ret void1534//1535//1536// CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l104.omp_outlined1537// CHECK9-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR0]] {1538// CHECK9-NEXT: entry:1539// CHECK9-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 81540// CHECK9-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 81541// CHECK9-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 81542// CHECK9-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 81543// CHECK9-NEXT: ret void1544//1545//1546// CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l1081547// CHECK9-SAME: (ptr noalias noundef [[DYN_PTR:%.*]]) #[[ATTR0]] {1548// CHECK9-NEXT: entry:1549// CHECK9-NEXT: [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 81550// CHECK9-NEXT: store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 81551// CHECK9-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB1]], i32 0, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l108.omp_outlined)1552// CHECK9-NEXT: ret void1553//1554//1555// CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l108.omp_outlined1556// CHECK9-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR0]] {1557// CHECK9-NEXT: entry:1558// CHECK9-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 81559// CHECK9-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 81560// CHECK9-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 81561// CHECK9-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 81562// CHECK9-NEXT: ret void1563//1564//1565// CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2S12r1Ei_l1211566// CHECK9-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], ptr noundef [[THIS:%.*]], i64 noundef [[B:%.*]], i64 noundef [[DOTCAPTURE_EXPR_:%.*]]) #[[ATTR0]] {1567// CHECK9-NEXT: entry:1568// CHECK9-NEXT: [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 81569// CHECK9-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 81570// CHECK9-NEXT: [[B_ADDR:%.*]] = alloca i64, align 81571// CHECK9-NEXT: [[DOTCAPTURE_EXPR__ADDR:%.*]] = alloca i64, align 81572// CHECK9-NEXT: [[B_CASTED:%.*]] = alloca i64, align 81573// CHECK9-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 41574// CHECK9-NEXT: [[DOTBOUND_ZERO_ADDR:%.*]] = alloca i32, align 41575// CHECK9-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])1576// CHECK9-NEXT: store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 81577// CHECK9-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 81578// CHECK9-NEXT: store i64 [[B]], ptr [[B_ADDR]], align 81579// CHECK9-NEXT: store i64 [[DOTCAPTURE_EXPR_]], ptr [[DOTCAPTURE_EXPR__ADDR]], align 81580// CHECK9-NEXT: [[TMP1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 81581// CHECK9-NEXT: [[TMP2:%.*]] = load i32, ptr [[B_ADDR]], align 41582// CHECK9-NEXT: store i32 [[TMP2]], ptr [[B_CASTED]], align 41583// CHECK9-NEXT: [[TMP3:%.*]] = load i64, ptr [[B_CASTED]], align 81584// CHECK9-NEXT: [[TMP4:%.*]] = load i8, ptr [[DOTCAPTURE_EXPR__ADDR]], align 11585// CHECK9-NEXT: [[LOADEDV:%.*]] = trunc i8 [[TMP4]] to i11586// CHECK9-NEXT: br i1 [[LOADEDV]], label [[OMP_IF_THEN:%.*]], label [[OMP_IF_ELSE:%.*]]1587// CHECK9: omp_if.then:1588// CHECK9-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB1]], i32 2, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2S12r1Ei_l121.omp_outlined, ptr [[TMP1]], i64 [[TMP3]])1589// CHECK9-NEXT: br label [[OMP_IF_END:%.*]]1590// CHECK9: omp_if.else:1591// CHECK9-NEXT: call void @__kmpc_serialized_parallel(ptr @[[GLOB1]], i32 [[TMP0]])1592// CHECK9-NEXT: store i32 [[TMP0]], ptr [[DOTTHREADID_TEMP_]], align 41593// CHECK9-NEXT: store i32 0, ptr [[DOTBOUND_ZERO_ADDR]], align 41594// CHECK9-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2S12r1Ei_l121.omp_outlined(ptr [[DOTTHREADID_TEMP_]], ptr [[DOTBOUND_ZERO_ADDR]], ptr [[TMP1]], i64 [[TMP3]]) #[[ATTR1]]1595// CHECK9-NEXT: call void @__kmpc_end_serialized_parallel(ptr @[[GLOB1]], i32 [[TMP0]])1596// CHECK9-NEXT: br label [[OMP_IF_END]]1597// CHECK9: omp_if.end:1598// CHECK9-NEXT: ret void1599//1600//1601// CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2S12r1Ei_l121.omp_outlined1602// CHECK9-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef [[THIS:%.*]], i64 noundef [[B:%.*]]) #[[ATTR0]] {1603// CHECK9-NEXT: entry:1604// CHECK9-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 81605// CHECK9-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 81606// CHECK9-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 81607// CHECK9-NEXT: [[B_ADDR:%.*]] = alloca i64, align 81608// CHECK9-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 81609// CHECK9-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 81610// CHECK9-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 81611// CHECK9-NEXT: store i64 [[B]], ptr [[B_ADDR]], align 81612// CHECK9-NEXT: [[TMP0:%.*]] = load ptr, ptr [[THIS_ADDR]], align 81613// CHECK9-NEXT: [[TMP1:%.*]] = load i32, ptr [[B_ADDR]], align 41614// CHECK9-NEXT: [[CONV:%.*]] = sitofp i32 [[TMP1]] to double1615// CHECK9-NEXT: [[ADD:%.*]] = fadd double [[CONV]], 1.500000e+001616// CHECK9-NEXT: [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_S1:%.*]], ptr [[TMP0]], i32 0, i32 01617// CHECK9-NEXT: store double [[ADD]], ptr [[A]], align 81618// CHECK9-NEXT: ret void1619//1620//1621// CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2S12r1Ei_l1261622// CHECK9-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], ptr noundef [[THIS:%.*]], i64 noundef [[DOTCAPTURE_EXPR_:%.*]]) #[[ATTR0]] {1623// CHECK9-NEXT: entry:1624// CHECK9-NEXT: [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 81625// CHECK9-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 81626// CHECK9-NEXT: [[DOTCAPTURE_EXPR__ADDR:%.*]] = alloca i64, align 81627// CHECK9-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 41628// CHECK9-NEXT: [[DOTBOUND_ZERO_ADDR:%.*]] = alloca i32, align 41629// CHECK9-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])1630// CHECK9-NEXT: store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 81631// CHECK9-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 81632// CHECK9-NEXT: store i64 [[DOTCAPTURE_EXPR_]], ptr [[DOTCAPTURE_EXPR__ADDR]], align 81633// CHECK9-NEXT: [[TMP1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 81634// CHECK9-NEXT: [[TMP2:%.*]] = load i8, ptr [[DOTCAPTURE_EXPR__ADDR]], align 11635// CHECK9-NEXT: [[LOADEDV:%.*]] = trunc i8 [[TMP2]] to i11636// CHECK9-NEXT: br i1 [[LOADEDV]], label [[OMP_IF_THEN:%.*]], label [[OMP_IF_ELSE:%.*]]1637// CHECK9: omp_if.then:1638// CHECK9-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB1]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2S12r1Ei_l126.omp_outlined, ptr [[TMP1]])1639// CHECK9-NEXT: br label [[OMP_IF_END:%.*]]1640// CHECK9: omp_if.else:1641// CHECK9-NEXT: call void @__kmpc_serialized_parallel(ptr @[[GLOB1]], i32 [[TMP0]])1642// CHECK9-NEXT: store i32 [[TMP0]], ptr [[DOTTHREADID_TEMP_]], align 41643// CHECK9-NEXT: store i32 0, ptr [[DOTBOUND_ZERO_ADDR]], align 41644// CHECK9-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2S12r1Ei_l126.omp_outlined(ptr [[DOTTHREADID_TEMP_]], ptr [[DOTBOUND_ZERO_ADDR]], ptr [[TMP1]]) #[[ATTR1]]1645// CHECK9-NEXT: call void @__kmpc_end_serialized_parallel(ptr @[[GLOB1]], i32 [[TMP0]])1646// CHECK9-NEXT: br label [[OMP_IF_END]]1647// CHECK9: omp_if.end:1648// CHECK9-NEXT: ret void1649//1650//1651// CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2S12r1Ei_l126.omp_outlined1652// CHECK9-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef [[THIS:%.*]]) #[[ATTR0]] {1653// CHECK9-NEXT: entry:1654// CHECK9-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 81655// CHECK9-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 81656// CHECK9-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 81657// CHECK9-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 81658// CHECK9-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 81659// CHECK9-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 81660// CHECK9-NEXT: [[TMP0:%.*]] = load ptr, ptr [[THIS_ADDR]], align 81661// CHECK9-NEXT: [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_S1:%.*]], ptr [[TMP0]], i32 0, i32 01662// CHECK9-NEXT: store double 2.500000e+00, ptr [[A]], align 81663// CHECK9-NEXT: ret void1664//1665//1666// CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l871667// CHECK9-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], i64 noundef [[A:%.*]]) #[[ATTR0]] {1668// CHECK9-NEXT: entry:1669// CHECK9-NEXT: [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 81670// CHECK9-NEXT: [[A_ADDR:%.*]] = alloca i64, align 81671// CHECK9-NEXT: [[A_CASTED:%.*]] = alloca i64, align 81672// CHECK9-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 41673// CHECK9-NEXT: [[DOTBOUND_ZERO_ADDR:%.*]] = alloca i32, align 41674// CHECK9-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])1675// CHECK9-NEXT: store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 81676// CHECK9-NEXT: store i64 [[A]], ptr [[A_ADDR]], align 81677// CHECK9-NEXT: [[TMP1:%.*]] = load i32, ptr [[A_ADDR]], align 41678// CHECK9-NEXT: store i32 [[TMP1]], ptr [[A_CASTED]], align 41679// CHECK9-NEXT: [[TMP2:%.*]] = load i64, ptr [[A_CASTED]], align 81680// CHECK9-NEXT: call void @__kmpc_serialized_parallel(ptr @[[GLOB1]], i32 [[TMP0]])1681// CHECK9-NEXT: store i32 [[TMP0]], ptr [[DOTTHREADID_TEMP_]], align 41682// CHECK9-NEXT: store i32 0, ptr [[DOTBOUND_ZERO_ADDR]], align 41683// CHECK9-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l87.omp_outlined(ptr [[DOTTHREADID_TEMP_]], ptr [[DOTBOUND_ZERO_ADDR]], i64 [[TMP2]]) #[[ATTR1]]1684// CHECK9-NEXT: call void @__kmpc_end_serialized_parallel(ptr @[[GLOB1]], i32 [[TMP0]])1685// CHECK9-NEXT: ret void1686//1687//1688// CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l87.omp_outlined1689// CHECK9-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i64 noundef [[A:%.*]]) #[[ATTR0]] {1690// CHECK9-NEXT: entry:1691// CHECK9-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 81692// CHECK9-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 81693// CHECK9-NEXT: [[A_ADDR:%.*]] = alloca i64, align 81694// CHECK9-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 81695// CHECK9-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 81696// CHECK9-NEXT: store i64 [[A]], ptr [[A_ADDR]], align 81697// CHECK9-NEXT: [[TMP0:%.*]] = load i32, ptr [[A_ADDR]], align 41698// CHECK9-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP0]], 11699// CHECK9-NEXT: store i32 [[ADD]], ptr [[A_ADDR]], align 41700// CHECK9-NEXT: ret void1701//1702//1703// CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l931704// CHECK9-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], i64 noundef [[A:%.*]], i64 noundef [[B:%.*]]) #[[ATTR0]] {1705// CHECK9-NEXT: entry:1706// CHECK9-NEXT: [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 81707// CHECK9-NEXT: [[A_ADDR:%.*]] = alloca i64, align 81708// CHECK9-NEXT: [[B_ADDR:%.*]] = alloca i64, align 81709// CHECK9-NEXT: [[A_CASTED:%.*]] = alloca i64, align 81710// CHECK9-NEXT: [[B_CASTED:%.*]] = alloca i64, align 81711// CHECK9-NEXT: store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 81712// CHECK9-NEXT: store i64 [[A]], ptr [[A_ADDR]], align 81713// CHECK9-NEXT: store i64 [[B]], ptr [[B_ADDR]], align 81714// CHECK9-NEXT: [[TMP0:%.*]] = load i32, ptr [[A_ADDR]], align 41715// CHECK9-NEXT: store i32 [[TMP0]], ptr [[A_CASTED]], align 41716// CHECK9-NEXT: [[TMP1:%.*]] = load i64, ptr [[A_CASTED]], align 81717// CHECK9-NEXT: [[TMP2:%.*]] = load i16, ptr [[B_ADDR]], align 21718// CHECK9-NEXT: store i16 [[TMP2]], ptr [[B_CASTED]], align 21719// CHECK9-NEXT: [[TMP3:%.*]] = load i64, ptr [[B_CASTED]], align 81720// CHECK9-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB1]], i32 2, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l93.omp_outlined, i64 [[TMP1]], i64 [[TMP3]])1721// CHECK9-NEXT: ret void1722//1723//1724// CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l93.omp_outlined1725// CHECK9-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i64 noundef [[A:%.*]], i64 noundef [[B:%.*]]) #[[ATTR0]] {1726// CHECK9-NEXT: entry:1727// CHECK9-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 81728// CHECK9-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 81729// CHECK9-NEXT: [[A_ADDR:%.*]] = alloca i64, align 81730// CHECK9-NEXT: [[B_ADDR:%.*]] = alloca i64, align 81731// CHECK9-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 81732// CHECK9-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 81733// CHECK9-NEXT: store i64 [[A]], ptr [[A_ADDR]], align 81734// CHECK9-NEXT: store i64 [[B]], ptr [[B_ADDR]], align 81735// CHECK9-NEXT: [[TMP0:%.*]] = load i16, ptr [[B_ADDR]], align 21736// CHECK9-NEXT: [[CONV:%.*]] = sext i16 [[TMP0]] to i321737// CHECK9-NEXT: [[TMP1:%.*]] = load i32, ptr [[A_ADDR]], align 41738// CHECK9-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP1]], [[CONV]]1739// CHECK9-NEXT: store i32 [[ADD]], ptr [[A_ADDR]], align 41740// CHECK9-NEXT: ret void1741//1742//1743// CHECK11-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l1041744// CHECK11-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], i32 noundef [[DOTCAPTURE_EXPR_:%.*]]) #[[ATTR0:[0-9]+]] {1745// CHECK11-NEXT: entry:1746// CHECK11-NEXT: [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 41747// CHECK11-NEXT: [[DOTCAPTURE_EXPR__ADDR:%.*]] = alloca i32, align 41748// CHECK11-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 41749// CHECK11-NEXT: [[DOTBOUND_ZERO_ADDR:%.*]] = alloca i32, align 41750// CHECK11-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1:[0-9]+]])1751// CHECK11-NEXT: store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 41752// CHECK11-NEXT: store i32 [[DOTCAPTURE_EXPR_]], ptr [[DOTCAPTURE_EXPR__ADDR]], align 41753// CHECK11-NEXT: [[TMP1:%.*]] = load i8, ptr [[DOTCAPTURE_EXPR__ADDR]], align 11754// CHECK11-NEXT: [[LOADEDV:%.*]] = trunc i8 [[TMP1]] to i11755// CHECK11-NEXT: br i1 [[LOADEDV]], label [[OMP_IF_THEN:%.*]], label [[OMP_IF_ELSE:%.*]]1756// CHECK11: omp_if.then:1757// CHECK11-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB1]], i32 0, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l104.omp_outlined)1758// CHECK11-NEXT: br label [[OMP_IF_END:%.*]]1759// CHECK11: omp_if.else:1760// CHECK11-NEXT: call void @__kmpc_serialized_parallel(ptr @[[GLOB1]], i32 [[TMP0]])1761// CHECK11-NEXT: store i32 [[TMP0]], ptr [[DOTTHREADID_TEMP_]], align 41762// CHECK11-NEXT: store i32 0, ptr [[DOTBOUND_ZERO_ADDR]], align 41763// CHECK11-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l104.omp_outlined(ptr [[DOTTHREADID_TEMP_]], ptr [[DOTBOUND_ZERO_ADDR]]) #[[ATTR1:[0-9]+]]1764// CHECK11-NEXT: call void @__kmpc_end_serialized_parallel(ptr @[[GLOB1]], i32 [[TMP0]])1765// CHECK11-NEXT: br label [[OMP_IF_END]]1766// CHECK11: omp_if.end:1767// CHECK11-NEXT: ret void1768//1769//1770// CHECK11-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l104.omp_outlined1771// CHECK11-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR0]] {1772// CHECK11-NEXT: entry:1773// CHECK11-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 41774// CHECK11-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 41775// CHECK11-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 41776// CHECK11-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 41777// CHECK11-NEXT: ret void1778//1779//1780// CHECK11-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l1081781// CHECK11-SAME: (ptr noalias noundef [[DYN_PTR:%.*]]) #[[ATTR0]] {1782// CHECK11-NEXT: entry:1783// CHECK11-NEXT: [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 41784// CHECK11-NEXT: store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 41785// CHECK11-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB1]], i32 0, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l108.omp_outlined)1786// CHECK11-NEXT: ret void1787//1788//1789// CHECK11-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l108.omp_outlined1790// CHECK11-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR0]] {1791// CHECK11-NEXT: entry:1792// CHECK11-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 41793// CHECK11-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 41794// CHECK11-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 41795// CHECK11-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 41796// CHECK11-NEXT: ret void1797//1798//1799// CHECK11-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2S12r1Ei_l1211800// CHECK11-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], ptr noundef [[THIS:%.*]], i32 noundef [[B:%.*]], i32 noundef [[DOTCAPTURE_EXPR_:%.*]]) #[[ATTR0]] {1801// CHECK11-NEXT: entry:1802// CHECK11-NEXT: [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 41803// CHECK11-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 41804// CHECK11-NEXT: [[B_ADDR:%.*]] = alloca i32, align 41805// CHECK11-NEXT: [[DOTCAPTURE_EXPR__ADDR:%.*]] = alloca i32, align 41806// CHECK11-NEXT: [[B_CASTED:%.*]] = alloca i32, align 41807// CHECK11-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 41808// CHECK11-NEXT: [[DOTBOUND_ZERO_ADDR:%.*]] = alloca i32, align 41809// CHECK11-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])1810// CHECK11-NEXT: store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 41811// CHECK11-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 41812// CHECK11-NEXT: store i32 [[B]], ptr [[B_ADDR]], align 41813// CHECK11-NEXT: store i32 [[DOTCAPTURE_EXPR_]], ptr [[DOTCAPTURE_EXPR__ADDR]], align 41814// CHECK11-NEXT: [[TMP1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 41815// CHECK11-NEXT: [[TMP2:%.*]] = load i32, ptr [[B_ADDR]], align 41816// CHECK11-NEXT: store i32 [[TMP2]], ptr [[B_CASTED]], align 41817// CHECK11-NEXT: [[TMP3:%.*]] = load i32, ptr [[B_CASTED]], align 41818// CHECK11-NEXT: [[TMP4:%.*]] = load i8, ptr [[DOTCAPTURE_EXPR__ADDR]], align 11819// CHECK11-NEXT: [[LOADEDV:%.*]] = trunc i8 [[TMP4]] to i11820// CHECK11-NEXT: br i1 [[LOADEDV]], label [[OMP_IF_THEN:%.*]], label [[OMP_IF_ELSE:%.*]]1821// CHECK11: omp_if.then:1822// CHECK11-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB1]], i32 2, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2S12r1Ei_l121.omp_outlined, ptr [[TMP1]], i32 [[TMP3]])1823// CHECK11-NEXT: br label [[OMP_IF_END:%.*]]1824// CHECK11: omp_if.else:1825// CHECK11-NEXT: call void @__kmpc_serialized_parallel(ptr @[[GLOB1]], i32 [[TMP0]])1826// CHECK11-NEXT: store i32 [[TMP0]], ptr [[DOTTHREADID_TEMP_]], align 41827// CHECK11-NEXT: store i32 0, ptr [[DOTBOUND_ZERO_ADDR]], align 41828// CHECK11-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2S12r1Ei_l121.omp_outlined(ptr [[DOTTHREADID_TEMP_]], ptr [[DOTBOUND_ZERO_ADDR]], ptr [[TMP1]], i32 [[TMP3]]) #[[ATTR1]]1829// CHECK11-NEXT: call void @__kmpc_end_serialized_parallel(ptr @[[GLOB1]], i32 [[TMP0]])1830// CHECK11-NEXT: br label [[OMP_IF_END]]1831// CHECK11: omp_if.end:1832// CHECK11-NEXT: ret void1833//1834//1835// CHECK11-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2S12r1Ei_l121.omp_outlined1836// CHECK11-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef [[THIS:%.*]], i32 noundef [[B:%.*]]) #[[ATTR0]] {1837// CHECK11-NEXT: entry:1838// CHECK11-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 41839// CHECK11-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 41840// CHECK11-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 41841// CHECK11-NEXT: [[B_ADDR:%.*]] = alloca i32, align 41842// CHECK11-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 41843// CHECK11-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 41844// CHECK11-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 41845// CHECK11-NEXT: store i32 [[B]], ptr [[B_ADDR]], align 41846// CHECK11-NEXT: [[TMP0:%.*]] = load ptr, ptr [[THIS_ADDR]], align 41847// CHECK11-NEXT: [[TMP1:%.*]] = load i32, ptr [[B_ADDR]], align 41848// CHECK11-NEXT: [[CONV:%.*]] = sitofp i32 [[TMP1]] to double1849// CHECK11-NEXT: [[ADD:%.*]] = fadd double [[CONV]], 1.500000e+001850// CHECK11-NEXT: [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_S1:%.*]], ptr [[TMP0]], i32 0, i32 01851// CHECK11-NEXT: store double [[ADD]], ptr [[A]], align 41852// CHECK11-NEXT: ret void1853//1854//1855// CHECK11-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2S12r1Ei_l1261856// CHECK11-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], ptr noundef [[THIS:%.*]], i32 noundef [[DOTCAPTURE_EXPR_:%.*]]) #[[ATTR0]] {1857// CHECK11-NEXT: entry:1858// CHECK11-NEXT: [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 41859// CHECK11-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 41860// CHECK11-NEXT: [[DOTCAPTURE_EXPR__ADDR:%.*]] = alloca i32, align 41861// CHECK11-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 41862// CHECK11-NEXT: [[DOTBOUND_ZERO_ADDR:%.*]] = alloca i32, align 41863// CHECK11-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])1864// CHECK11-NEXT: store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 41865// CHECK11-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 41866// CHECK11-NEXT: store i32 [[DOTCAPTURE_EXPR_]], ptr [[DOTCAPTURE_EXPR__ADDR]], align 41867// CHECK11-NEXT: [[TMP1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 41868// CHECK11-NEXT: [[TMP2:%.*]] = load i8, ptr [[DOTCAPTURE_EXPR__ADDR]], align 11869// CHECK11-NEXT: [[LOADEDV:%.*]] = trunc i8 [[TMP2]] to i11870// CHECK11-NEXT: br i1 [[LOADEDV]], label [[OMP_IF_THEN:%.*]], label [[OMP_IF_ELSE:%.*]]1871// CHECK11: omp_if.then:1872// CHECK11-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB1]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2S12r1Ei_l126.omp_outlined, ptr [[TMP1]])1873// CHECK11-NEXT: br label [[OMP_IF_END:%.*]]1874// CHECK11: omp_if.else:1875// CHECK11-NEXT: call void @__kmpc_serialized_parallel(ptr @[[GLOB1]], i32 [[TMP0]])1876// CHECK11-NEXT: store i32 [[TMP0]], ptr [[DOTTHREADID_TEMP_]], align 41877// CHECK11-NEXT: store i32 0, ptr [[DOTBOUND_ZERO_ADDR]], align 41878// CHECK11-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2S12r1Ei_l126.omp_outlined(ptr [[DOTTHREADID_TEMP_]], ptr [[DOTBOUND_ZERO_ADDR]], ptr [[TMP1]]) #[[ATTR1]]1879// CHECK11-NEXT: call void @__kmpc_end_serialized_parallel(ptr @[[GLOB1]], i32 [[TMP0]])1880// CHECK11-NEXT: br label [[OMP_IF_END]]1881// CHECK11: omp_if.end:1882// CHECK11-NEXT: ret void1883//1884//1885// CHECK11-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2S12r1Ei_l126.omp_outlined1886// CHECK11-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef [[THIS:%.*]]) #[[ATTR0]] {1887// CHECK11-NEXT: entry:1888// CHECK11-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 41889// CHECK11-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 41890// CHECK11-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 41891// CHECK11-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 41892// CHECK11-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 41893// CHECK11-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 41894// CHECK11-NEXT: [[TMP0:%.*]] = load ptr, ptr [[THIS_ADDR]], align 41895// CHECK11-NEXT: [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_S1:%.*]], ptr [[TMP0]], i32 0, i32 01896// CHECK11-NEXT: store double 2.500000e+00, ptr [[A]], align 41897// CHECK11-NEXT: ret void1898//1899//1900// CHECK11-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l871901// CHECK11-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], i32 noundef [[A:%.*]]) #[[ATTR0]] {1902// CHECK11-NEXT: entry:1903// CHECK11-NEXT: [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 41904// CHECK11-NEXT: [[A_ADDR:%.*]] = alloca i32, align 41905// CHECK11-NEXT: [[A_CASTED:%.*]] = alloca i32, align 41906// CHECK11-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 41907// CHECK11-NEXT: [[DOTBOUND_ZERO_ADDR:%.*]] = alloca i32, align 41908// CHECK11-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])1909// CHECK11-NEXT: store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 41910// CHECK11-NEXT: store i32 [[A]], ptr [[A_ADDR]], align 41911// CHECK11-NEXT: [[TMP1:%.*]] = load i32, ptr [[A_ADDR]], align 41912// CHECK11-NEXT: store i32 [[TMP1]], ptr [[A_CASTED]], align 41913// CHECK11-NEXT: [[TMP2:%.*]] = load i32, ptr [[A_CASTED]], align 41914// CHECK11-NEXT: call void @__kmpc_serialized_parallel(ptr @[[GLOB1]], i32 [[TMP0]])1915// CHECK11-NEXT: store i32 [[TMP0]], ptr [[DOTTHREADID_TEMP_]], align 41916// CHECK11-NEXT: store i32 0, ptr [[DOTBOUND_ZERO_ADDR]], align 41917// CHECK11-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l87.omp_outlined(ptr [[DOTTHREADID_TEMP_]], ptr [[DOTBOUND_ZERO_ADDR]], i32 [[TMP2]]) #[[ATTR1]]1918// CHECK11-NEXT: call void @__kmpc_end_serialized_parallel(ptr @[[GLOB1]], i32 [[TMP0]])1919// CHECK11-NEXT: ret void1920//1921//1922// CHECK11-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l87.omp_outlined1923// CHECK11-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i32 noundef [[A:%.*]]) #[[ATTR0]] {1924// CHECK11-NEXT: entry:1925// CHECK11-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 41926// CHECK11-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 41927// CHECK11-NEXT: [[A_ADDR:%.*]] = alloca i32, align 41928// CHECK11-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 41929// CHECK11-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 41930// CHECK11-NEXT: store i32 [[A]], ptr [[A_ADDR]], align 41931// CHECK11-NEXT: [[TMP0:%.*]] = load i32, ptr [[A_ADDR]], align 41932// CHECK11-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP0]], 11933// CHECK11-NEXT: store i32 [[ADD]], ptr [[A_ADDR]], align 41934// CHECK11-NEXT: ret void1935//1936//1937// CHECK11-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l931938// CHECK11-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], i32 noundef [[A:%.*]], i32 noundef [[B:%.*]]) #[[ATTR0]] {1939// CHECK11-NEXT: entry:1940// CHECK11-NEXT: [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 41941// CHECK11-NEXT: [[A_ADDR:%.*]] = alloca i32, align 41942// CHECK11-NEXT: [[B_ADDR:%.*]] = alloca i32, align 41943// CHECK11-NEXT: [[A_CASTED:%.*]] = alloca i32, align 41944// CHECK11-NEXT: [[B_CASTED:%.*]] = alloca i32, align 41945// CHECK11-NEXT: store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 41946// CHECK11-NEXT: store i32 [[A]], ptr [[A_ADDR]], align 41947// CHECK11-NEXT: store i32 [[B]], ptr [[B_ADDR]], align 41948// CHECK11-NEXT: [[TMP0:%.*]] = load i32, ptr [[A_ADDR]], align 41949// CHECK11-NEXT: store i32 [[TMP0]], ptr [[A_CASTED]], align 41950// CHECK11-NEXT: [[TMP1:%.*]] = load i32, ptr [[A_CASTED]], align 41951// CHECK11-NEXT: [[TMP2:%.*]] = load i16, ptr [[B_ADDR]], align 21952// CHECK11-NEXT: store i16 [[TMP2]], ptr [[B_CASTED]], align 21953// CHECK11-NEXT: [[TMP3:%.*]] = load i32, ptr [[B_CASTED]], align 41954// CHECK11-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB1]], i32 2, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l93.omp_outlined, i32 [[TMP1]], i32 [[TMP3]])1955// CHECK11-NEXT: ret void1956//1957//1958// CHECK11-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l93.omp_outlined1959// CHECK11-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i32 noundef [[A:%.*]], i32 noundef [[B:%.*]]) #[[ATTR0]] {1960// CHECK11-NEXT: entry:1961// CHECK11-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 41962// CHECK11-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 41963// CHECK11-NEXT: [[A_ADDR:%.*]] = alloca i32, align 41964// CHECK11-NEXT: [[B_ADDR:%.*]] = alloca i32, align 41965// CHECK11-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 41966// CHECK11-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 41967// CHECK11-NEXT: store i32 [[A]], ptr [[A_ADDR]], align 41968// CHECK11-NEXT: store i32 [[B]], ptr [[B_ADDR]], align 41969// CHECK11-NEXT: [[TMP0:%.*]] = load i16, ptr [[B_ADDR]], align 21970// CHECK11-NEXT: [[CONV:%.*]] = sext i16 [[TMP0]] to i321971// CHECK11-NEXT: [[TMP1:%.*]] = load i32, ptr [[A_ADDR]], align 41972// CHECK11-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP1]], [[CONV]]1973// CHECK11-NEXT: store i32 [[ADD]], ptr [[A_ADDR]], align 41974// CHECK11-NEXT: ret void1975//1976