2634 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=61 -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=61 -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=61 -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=61 -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=61 -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=61 -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=61 -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=61 -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=61 -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=61 -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=61 -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=61 -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=61 -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=61 -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=61 -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=61 -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=61 -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=61 -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=61 -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=61 -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=61 -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=61 -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=61 -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=61 -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=61 -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=61 -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=61 -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=61 -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 37// expected-no-diagnostics38#ifndef HEADER39#define HEADER40 41 42 43 44// We have 6 target regions45 46 47 48// Check target registration is registered as a Ctor.49 50 51template<typename tx>52tx ftemplate(int n) {53 tx a = 0;54 55 #pragma omp target teams dyn_groupprivate(tx(20))56 {57 }58 59 short b = 1;60 #pragma omp target teams num_teams(b) dyn_groupprivate(1024)61 {62 a += b;63 }64 65 return a;66}67 68static69int fstatic(int n) {70 71 #pragma omp target teams distribute parallel for simd num_teams(n) dyn_groupprivate(n*32)72 for (int i = 0; i < n ; ++i) {73 }74 75 #pragma omp target teams dyn_groupprivate(fallback(default_mem): 32+n) nowait76 {77 }78 79 return n+1;80}81 82struct S1 {83 double a;84 85 int r1(int n){86 int b = 1;87 88 #pragma omp target teams dyn_groupprivate(fallback(null): n-b)89 {90 this->a = (double)b + 1.5;91 }92 93 #pragma omp target dyn_groupprivate(fallback(abort): 1024)94 {95 this->a = 2.5;96 }97 98 return (int)a;99 }100};101 102int bar(int n){103 int a = 0;104 105 S1 S;106 a += S.r1(n);107 108 a += fstatic(n);109 110 a += ftemplate<int>(n);111 112 return a;113}114 115 116 117 118 119 120 121 122 123 124 125 126 127 128 129 130 131 132 133 134 135// Check that the offloading functions are emitted and that the parallel function136// is appropriately guarded.137 138 139 140 141 142 143#endif144 145// CHECK1-LABEL: define {{[^@]+}}@_Z3bari146// CHECK1-SAME: (i32 noundef signext [[N:%.*]]) #[[ATTR0:[0-9]+]] {147// CHECK1-NEXT: entry:148// CHECK1-NEXT: [[N_ADDR:%.*]] = alloca i32, align 4149// CHECK1-NEXT: [[A:%.*]] = alloca i32, align 4150// CHECK1-NEXT: [[S:%.*]] = alloca [[STRUCT_S1:%.*]], align 8151// CHECK1-NEXT: store i32 [[N]], ptr [[N_ADDR]], align 4152// CHECK1-NEXT: store i32 0, ptr [[A]], align 4153// CHECK1-NEXT: [[TMP0:%.*]] = load i32, ptr [[N_ADDR]], align 4154// CHECK1-NEXT: [[CALL:%.*]] = call noundef signext i32 @_ZN2S12r1Ei(ptr noundef nonnull align 8 dereferenceable(8) [[S]], i32 noundef signext [[TMP0]])155// CHECK1-NEXT: [[TMP1:%.*]] = load i32, ptr [[A]], align 4156// CHECK1-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP1]], [[CALL]]157// CHECK1-NEXT: store i32 [[ADD]], ptr [[A]], align 4158// CHECK1-NEXT: [[TMP2:%.*]] = load i32, ptr [[N_ADDR]], align 4159// CHECK1-NEXT: [[CALL1:%.*]] = call noundef signext i32 @_ZL7fstatici(i32 noundef signext [[TMP2]])160// CHECK1-NEXT: [[TMP3:%.*]] = load i32, ptr [[A]], align 4161// CHECK1-NEXT: [[ADD2:%.*]] = add nsw i32 [[TMP3]], [[CALL1]]162// CHECK1-NEXT: store i32 [[ADD2]], ptr [[A]], align 4163// CHECK1-NEXT: [[TMP4:%.*]] = load i32, ptr [[N_ADDR]], align 4164// CHECK1-NEXT: [[CALL3:%.*]] = call noundef signext i32 @_Z9ftemplateIiET_i(i32 noundef signext [[TMP4]])165// CHECK1-NEXT: [[TMP5:%.*]] = load i32, ptr [[A]], align 4166// CHECK1-NEXT: [[ADD4:%.*]] = add nsw i32 [[TMP5]], [[CALL3]]167// CHECK1-NEXT: store i32 [[ADD4]], ptr [[A]], align 4168// CHECK1-NEXT: [[TMP6:%.*]] = load i32, ptr [[A]], align 4169// CHECK1-NEXT: ret i32 [[TMP6]]170//171//172// CHECK1-LABEL: define {{[^@]+}}@_ZN2S12r1Ei173// CHECK1-SAME: (ptr noundef nonnull align 8 dereferenceable(8) [[THIS:%.*]], i32 noundef signext [[N:%.*]]) #[[ATTR0]] comdat {174// CHECK1-NEXT: entry:175// CHECK1-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 8176// CHECK1-NEXT: [[N_ADDR:%.*]] = alloca i32, align 4177// CHECK1-NEXT: [[B:%.*]] = alloca i32, align 4178// CHECK1-NEXT: [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 4179// CHECK1-NEXT: [[B_CASTED:%.*]] = alloca i64, align 8180// CHECK1-NEXT: [[DOTCAPTURE_EXPR__CASTED:%.*]] = alloca i64, align 8181// CHECK1-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [3 x ptr], align 8182// CHECK1-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [3 x ptr], align 8183// CHECK1-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [3 x ptr], align 8184// CHECK1-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8185// CHECK1-NEXT: [[DOTOFFLOAD_BASEPTRS3:%.*]] = alloca [1 x ptr], align 8186// CHECK1-NEXT: [[DOTOFFLOAD_PTRS4:%.*]] = alloca [1 x ptr], align 8187// CHECK1-NEXT: [[DOTOFFLOAD_MAPPERS5:%.*]] = alloca [1 x ptr], align 8188// CHECK1-NEXT: [[KERNEL_ARGS6:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 8189// CHECK1-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 8190// CHECK1-NEXT: store i32 [[N]], ptr [[N_ADDR]], align 4191// CHECK1-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8192// CHECK1-NEXT: store i32 1, ptr [[B]], align 4193// CHECK1-NEXT: [[TMP0:%.*]] = load i32, ptr [[N_ADDR]], align 4194// CHECK1-NEXT: [[TMP1:%.*]] = load i32, ptr [[B]], align 4195// CHECK1-NEXT: [[SUB:%.*]] = sub nsw i32 [[TMP0]], [[TMP1]]196// CHECK1-NEXT: store i32 [[SUB]], ptr [[DOTCAPTURE_EXPR_]], align 4197// CHECK1-NEXT: [[TMP2:%.*]] = load i32, ptr [[B]], align 4198// CHECK1-NEXT: store i32 [[TMP2]], ptr [[B_CASTED]], align 4199// CHECK1-NEXT: [[TMP3:%.*]] = load i64, ptr [[B_CASTED]], align 8200// CHECK1-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4201// CHECK1-NEXT: store i32 [[TMP4]], ptr [[DOTCAPTURE_EXPR__CASTED]], align 4202// CHECK1-NEXT: [[TMP5:%.*]] = load i64, ptr [[DOTCAPTURE_EXPR__CASTED]], align 8203// CHECK1-NEXT: [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_S1:%.*]], ptr [[THIS1]], i32 0, i32 0204// CHECK1-NEXT: [[TMP6:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0205// CHECK1-NEXT: store ptr [[THIS1]], ptr [[TMP6]], align 8206// CHECK1-NEXT: [[TMP7:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0207// CHECK1-NEXT: store ptr [[A]], ptr [[TMP7]], align 8208// CHECK1-NEXT: [[TMP8:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 0209// CHECK1-NEXT: store ptr null, ptr [[TMP8]], align 8210// CHECK1-NEXT: [[TMP9:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 1211// CHECK1-NEXT: store i64 [[TMP3]], ptr [[TMP9]], align 8212// CHECK1-NEXT: [[TMP10:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 1213// CHECK1-NEXT: store i64 [[TMP3]], ptr [[TMP10]], align 8214// CHECK1-NEXT: [[TMP11:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 1215// CHECK1-NEXT: store ptr null, ptr [[TMP11]], align 8216// CHECK1-NEXT: [[TMP12:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 2217// CHECK1-NEXT: store i64 [[TMP5]], ptr [[TMP12]], align 8218// CHECK1-NEXT: [[TMP13:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 2219// CHECK1-NEXT: store i64 [[TMP5]], ptr [[TMP13]], align 8220// CHECK1-NEXT: [[TMP14:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 2221// CHECK1-NEXT: store ptr null, ptr [[TMP14]], align 8222// CHECK1-NEXT: [[TMP15:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0223// CHECK1-NEXT: [[TMP16:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0224// CHECK1-NEXT: [[TMP17:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4225// CHECK1-NEXT: [[TMP18:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0226// CHECK1-NEXT: store i32 3, ptr [[TMP18]], align 4227// CHECK1-NEXT: [[TMP19:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1228// CHECK1-NEXT: store i32 3, ptr [[TMP19]], align 4229// CHECK1-NEXT: [[TMP20:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2230// CHECK1-NEXT: store ptr [[TMP15]], ptr [[TMP20]], align 8231// CHECK1-NEXT: [[TMP21:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3232// CHECK1-NEXT: store ptr [[TMP16]], ptr [[TMP21]], align 8233// CHECK1-NEXT: [[TMP22:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4234// CHECK1-NEXT: store ptr @.offload_sizes, ptr [[TMP22]], align 8235// CHECK1-NEXT: [[TMP23:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5236// CHECK1-NEXT: store ptr @.offload_maptypes, ptr [[TMP23]], align 8237// CHECK1-NEXT: [[TMP24:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6238// CHECK1-NEXT: store ptr null, ptr [[TMP24]], align 8239// CHECK1-NEXT: [[TMP25:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7240// CHECK1-NEXT: store ptr null, ptr [[TMP25]], align 8241// CHECK1-NEXT: [[TMP26:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8242// CHECK1-NEXT: store i64 0, ptr [[TMP26]], align 8243// CHECK1-NEXT: [[TMP27:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9244// CHECK1-NEXT: store i64 4, ptr [[TMP27]], align 8245// CHECK1-NEXT: [[TMP28:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10246// CHECK1-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP28]], align 4247// CHECK1-NEXT: [[TMP29:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11248// CHECK1-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP29]], align 4249// CHECK1-NEXT: [[TMP30:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12250// CHECK1-NEXT: store i32 [[TMP17]], ptr [[TMP30]], align 4251// CHECK1-NEXT: [[TMP31:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1:[0-9]+]], i64 -1, i32 0, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2S12r1Ei_l88.region_id, ptr [[KERNEL_ARGS]])252// CHECK1-NEXT: [[TMP32:%.*]] = icmp ne i32 [[TMP31]], 0253// CHECK1-NEXT: br i1 [[TMP32]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]254// CHECK1: omp_offload.failed:255// CHECK1-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2S12r1Ei_l88(ptr [[THIS1]], i64 [[TMP3]], i64 [[TMP5]]) #[[ATTR2:[0-9]+]]256// CHECK1-NEXT: br label [[OMP_OFFLOAD_CONT]]257// CHECK1: omp_offload.cont:258// CHECK1-NEXT: [[A2:%.*]] = getelementptr inbounds nuw [[STRUCT_S1]], ptr [[THIS1]], i32 0, i32 0259// CHECK1-NEXT: [[TMP33:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS3]], i32 0, i32 0260// CHECK1-NEXT: store ptr [[THIS1]], ptr [[TMP33]], align 8261// CHECK1-NEXT: [[TMP34:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS4]], i32 0, i32 0262// CHECK1-NEXT: store ptr [[A2]], ptr [[TMP34]], align 8263// CHECK1-NEXT: [[TMP35:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS5]], i64 0, i64 0264// CHECK1-NEXT: store ptr null, ptr [[TMP35]], align 8265// CHECK1-NEXT: [[TMP36:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS3]], i32 0, i32 0266// CHECK1-NEXT: [[TMP37:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS4]], i32 0, i32 0267// CHECK1-NEXT: [[TMP38:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 0268// CHECK1-NEXT: store i32 3, ptr [[TMP38]], align 4269// CHECK1-NEXT: [[TMP39:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 1270// CHECK1-NEXT: store i32 1, ptr [[TMP39]], align 4271// CHECK1-NEXT: [[TMP40:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 2272// CHECK1-NEXT: store ptr [[TMP36]], ptr [[TMP40]], align 8273// CHECK1-NEXT: [[TMP41:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 3274// CHECK1-NEXT: store ptr [[TMP37]], ptr [[TMP41]], align 8275// CHECK1-NEXT: [[TMP42:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 4276// CHECK1-NEXT: store ptr @.offload_sizes.1, ptr [[TMP42]], align 8277// CHECK1-NEXT: [[TMP43:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 5278// CHECK1-NEXT: store ptr @.offload_maptypes.2, ptr [[TMP43]], align 8279// CHECK1-NEXT: [[TMP44:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 6280// CHECK1-NEXT: store ptr null, ptr [[TMP44]], align 8281// CHECK1-NEXT: [[TMP45:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 7282// CHECK1-NEXT: store ptr null, ptr [[TMP45]], align 8283// CHECK1-NEXT: [[TMP46:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 8284// CHECK1-NEXT: store i64 0, ptr [[TMP46]], align 8285// CHECK1-NEXT: [[TMP47:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 9286// CHECK1-NEXT: store i64 0, ptr [[TMP47]], align 8287// CHECK1-NEXT: [[TMP48:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 10288// CHECK1-NEXT: store [3 x i32] [i32 -1, i32 0, i32 0], ptr [[TMP48]], align 4289// CHECK1-NEXT: [[TMP49:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 11290// CHECK1-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP49]], align 4291// CHECK1-NEXT: [[TMP50:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 12292// CHECK1-NEXT: store i32 1024, ptr [[TMP50]], align 4293// CHECK1-NEXT: [[TMP51:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 -1, i32 -1, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2S12r1Ei_l93.region_id, ptr [[KERNEL_ARGS6]])294// CHECK1-NEXT: [[TMP52:%.*]] = icmp ne i32 [[TMP51]], 0295// CHECK1-NEXT: br i1 [[TMP52]], label [[OMP_OFFLOAD_FAILED7:%.*]], label [[OMP_OFFLOAD_CONT8:%.*]]296// CHECK1: omp_offload.failed7:297// CHECK1-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2S12r1Ei_l93(ptr [[THIS1]]) #[[ATTR2]]298// CHECK1-NEXT: br label [[OMP_OFFLOAD_CONT8]]299// CHECK1: omp_offload.cont8:300// CHECK1-NEXT: [[A9:%.*]] = getelementptr inbounds nuw [[STRUCT_S1]], ptr [[THIS1]], i32 0, i32 0301// CHECK1-NEXT: [[TMP53:%.*]] = load double, ptr [[A9]], align 8302// CHECK1-NEXT: [[CONV:%.*]] = fptosi double [[TMP53]] to i32303// CHECK1-NEXT: ret i32 [[CONV]]304//305//306// CHECK1-LABEL: define {{[^@]+}}@_ZL7fstatici307// CHECK1-SAME: (i32 noundef signext [[N:%.*]]) #[[ATTR0]] {308// CHECK1-NEXT: entry:309// CHECK1-NEXT: [[N_ADDR:%.*]] = alloca i32, align 4310// CHECK1-NEXT: [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 4311// CHECK1-NEXT: [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 4312// CHECK1-NEXT: [[N_CASTED:%.*]] = alloca i64, align 8313// CHECK1-NEXT: [[DOTCAPTURE_EXPR__CASTED:%.*]] = alloca i64, align 8314// CHECK1-NEXT: [[DOTCAPTURE_EXPR__CASTED2:%.*]] = alloca i64, align 8315// CHECK1-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [3 x ptr], align 8316// CHECK1-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [3 x ptr], align 8317// CHECK1-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [3 x ptr], align 8318// CHECK1-NEXT: [[TMP:%.*]] = alloca i32, align 4319// CHECK1-NEXT: [[DOTCAPTURE_EXPR_3:%.*]] = alloca i32, align 4320// CHECK1-NEXT: [[DOTCAPTURE_EXPR_4:%.*]] = alloca i32, align 4321// CHECK1-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8322// CHECK1-NEXT: [[DOTCAPTURE_EXPR_6:%.*]] = alloca i32, align 4323// CHECK1-NEXT: [[DOTCAPTURE_EXPR__CASTED8:%.*]] = alloca i64, align 8324// CHECK1-NEXT: [[DOTOFFLOAD_BASEPTRS9:%.*]] = alloca [1 x ptr], align 8325// CHECK1-NEXT: [[DOTOFFLOAD_PTRS10:%.*]] = alloca [1 x ptr], align 8326// CHECK1-NEXT: [[DOTOFFLOAD_MAPPERS11:%.*]] = alloca [1 x ptr], align 8327// CHECK1-NEXT: [[AGG_CAPTURED:%.*]] = alloca [[STRUCT_ANON:%.*]], align 4328// CHECK1-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])329// CHECK1-NEXT: store i32 [[N]], ptr [[N_ADDR]], align 4330// CHECK1-NEXT: [[TMP1:%.*]] = load i32, ptr [[N_ADDR]], align 4331// CHECK1-NEXT: store i32 [[TMP1]], ptr [[DOTCAPTURE_EXPR_]], align 4332// CHECK1-NEXT: [[TMP2:%.*]] = load i32, ptr [[N_ADDR]], align 4333// CHECK1-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP2]], 32334// CHECK1-NEXT: store i32 [[MUL]], ptr [[DOTCAPTURE_EXPR_1]], align 4335// CHECK1-NEXT: [[TMP3:%.*]] = load i32, ptr [[N_ADDR]], align 4336// CHECK1-NEXT: store i32 [[TMP3]], ptr [[N_CASTED]], align 4337// CHECK1-NEXT: [[TMP4:%.*]] = load i64, ptr [[N_CASTED]], align 8338// CHECK1-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4339// CHECK1-NEXT: store i32 [[TMP5]], ptr [[DOTCAPTURE_EXPR__CASTED]], align 4340// CHECK1-NEXT: [[TMP6:%.*]] = load i64, ptr [[DOTCAPTURE_EXPR__CASTED]], align 8341// CHECK1-NEXT: [[TMP7:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4342// CHECK1-NEXT: store i32 [[TMP7]], ptr [[DOTCAPTURE_EXPR__CASTED2]], align 4343// CHECK1-NEXT: [[TMP8:%.*]] = load i64, ptr [[DOTCAPTURE_EXPR__CASTED2]], align 8344// CHECK1-NEXT: [[TMP9:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0345// CHECK1-NEXT: store i64 [[TMP4]], ptr [[TMP9]], align 8346// CHECK1-NEXT: [[TMP10:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0347// CHECK1-NEXT: store i64 [[TMP4]], ptr [[TMP10]], align 8348// CHECK1-NEXT: [[TMP11:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 0349// CHECK1-NEXT: store ptr null, ptr [[TMP11]], align 8350// CHECK1-NEXT: [[TMP12:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 1351// CHECK1-NEXT: store i64 [[TMP6]], ptr [[TMP12]], align 8352// CHECK1-NEXT: [[TMP13:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 1353// CHECK1-NEXT: store i64 [[TMP6]], ptr [[TMP13]], align 8354// CHECK1-NEXT: [[TMP14:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 1355// CHECK1-NEXT: store ptr null, ptr [[TMP14]], align 8356// CHECK1-NEXT: [[TMP15:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 2357// CHECK1-NEXT: store i64 [[TMP8]], ptr [[TMP15]], align 8358// CHECK1-NEXT: [[TMP16:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 2359// CHECK1-NEXT: store i64 [[TMP8]], ptr [[TMP16]], align 8360// CHECK1-NEXT: [[TMP17:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 2361// CHECK1-NEXT: store ptr null, ptr [[TMP17]], align 8362// CHECK1-NEXT: [[TMP18:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0363// CHECK1-NEXT: [[TMP19:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0364// CHECK1-NEXT: [[TMP20:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4365// CHECK1-NEXT: [[TMP21:%.*]] = load i32, ptr [[N_ADDR]], align 4366// CHECK1-NEXT: store i32 [[TMP21]], ptr [[DOTCAPTURE_EXPR_3]], align 4367// CHECK1-NEXT: [[TMP22:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_3]], align 4368// CHECK1-NEXT: [[SUB:%.*]] = sub nsw i32 [[TMP22]], 0369// CHECK1-NEXT: [[DIV:%.*]] = sdiv i32 [[SUB]], 1370// CHECK1-NEXT: [[SUB5:%.*]] = sub nsw i32 [[DIV]], 1371// CHECK1-NEXT: store i32 [[SUB5]], ptr [[DOTCAPTURE_EXPR_4]], align 4372// CHECK1-NEXT: [[TMP23:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_4]], align 4373// CHECK1-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP23]], 1374// CHECK1-NEXT: [[TMP24:%.*]] = zext i32 [[ADD]] to i64375// CHECK1-NEXT: [[TMP25:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4376// CHECK1-NEXT: [[TMP26:%.*]] = insertvalue [3 x i32] zeroinitializer, i32 [[TMP20]], 0377// CHECK1-NEXT: [[TMP27:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0378// CHECK1-NEXT: store i32 3, ptr [[TMP27]], align 4379// CHECK1-NEXT: [[TMP28:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1380// CHECK1-NEXT: store i32 3, ptr [[TMP28]], align 4381// CHECK1-NEXT: [[TMP29:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2382// CHECK1-NEXT: store ptr [[TMP18]], ptr [[TMP29]], align 8383// CHECK1-NEXT: [[TMP30:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3384// CHECK1-NEXT: store ptr [[TMP19]], ptr [[TMP30]], align 8385// CHECK1-NEXT: [[TMP31:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4386// CHECK1-NEXT: store ptr @.offload_sizes.3, ptr [[TMP31]], align 8387// CHECK1-NEXT: [[TMP32:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5388// CHECK1-NEXT: store ptr @.offload_maptypes.4, ptr [[TMP32]], align 8389// CHECK1-NEXT: [[TMP33:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6390// CHECK1-NEXT: store ptr null, ptr [[TMP33]], align 8391// CHECK1-NEXT: [[TMP34:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7392// CHECK1-NEXT: store ptr null, ptr [[TMP34]], align 8393// CHECK1-NEXT: [[TMP35:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8394// CHECK1-NEXT: store i64 [[TMP24]], ptr [[TMP35]], align 8395// CHECK1-NEXT: [[TMP36:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9396// CHECK1-NEXT: store i64 8, ptr [[TMP36]], align 8397// CHECK1-NEXT: [[TMP37:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10398// CHECK1-NEXT: store [3 x i32] [[TMP26]], ptr [[TMP37]], align 4399// CHECK1-NEXT: [[TMP38:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11400// CHECK1-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP38]], align 4401// CHECK1-NEXT: [[TMP39:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12402// CHECK1-NEXT: store i32 [[TMP25]], ptr [[TMP39]], align 4403// CHECK1-NEXT: [[TMP40:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 -1, i32 [[TMP20]], i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l71.region_id, ptr [[KERNEL_ARGS]])404// CHECK1-NEXT: [[TMP41:%.*]] = icmp ne i32 [[TMP40]], 0405// CHECK1-NEXT: br i1 [[TMP41]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]406// CHECK1: omp_offload.failed:407// CHECK1-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l71(i64 [[TMP4]], i64 [[TMP6]], i64 [[TMP8]]) #[[ATTR2]]408// CHECK1-NEXT: br label [[OMP_OFFLOAD_CONT]]409// CHECK1: omp_offload.cont:410// CHECK1-NEXT: [[TMP42:%.*]] = load i32, ptr [[N_ADDR]], align 4411// CHECK1-NEXT: [[ADD7:%.*]] = add nsw i32 32, [[TMP42]]412// CHECK1-NEXT: store i32 [[ADD7]], ptr [[DOTCAPTURE_EXPR_6]], align 4413// CHECK1-NEXT: [[TMP43:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_6]], align 4414// CHECK1-NEXT: store i32 [[TMP43]], ptr [[DOTCAPTURE_EXPR__CASTED8]], align 4415// CHECK1-NEXT: [[TMP44:%.*]] = load i64, ptr [[DOTCAPTURE_EXPR__CASTED8]], align 8416// CHECK1-NEXT: [[TMP45:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS9]], i32 0, i32 0417// CHECK1-NEXT: store i64 [[TMP44]], ptr [[TMP45]], align 8418// CHECK1-NEXT: [[TMP46:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS10]], i32 0, i32 0419// CHECK1-NEXT: store i64 [[TMP44]], ptr [[TMP46]], align 8420// CHECK1-NEXT: [[TMP47:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS11]], i64 0, i64 0421// CHECK1-NEXT: store ptr null, ptr [[TMP47]], align 8422// CHECK1-NEXT: [[TMP48:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS9]], i32 0, i32 0423// CHECK1-NEXT: [[TMP49:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS10]], i32 0, i32 0424// CHECK1-NEXT: [[TMP50:%.*]] = getelementptr inbounds nuw [[STRUCT_ANON]], ptr [[AGG_CAPTURED]], i32 0, i32 0425// CHECK1-NEXT: [[TMP51:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_6]], align 4426// CHECK1-NEXT: store i32 [[TMP51]], ptr [[TMP50]], align 4427// CHECK1-NEXT: [[TMP52:%.*]] = call ptr @__kmpc_omp_target_task_alloc(ptr @[[GLOB1]], i32 [[TMP0]], i32 1, i64 64, i64 4, ptr @.omp_task_entry., i64 -1)428// CHECK1-NEXT: [[TMP53:%.*]] = getelementptr inbounds nuw [[STRUCT_KMP_TASK_T_WITH_PRIVATES:%.*]], ptr [[TMP52]], i32 0, i32 0429// CHECK1-NEXT: [[TMP54:%.*]] = getelementptr inbounds nuw [[STRUCT_KMP_TASK_T:%.*]], ptr [[TMP53]], i32 0, i32 0430// CHECK1-NEXT: [[TMP55:%.*]] = load ptr, ptr [[TMP54]], align 8431// CHECK1-NEXT: call void @llvm.memcpy.p0.p0.i64(ptr align 4 [[TMP55]], ptr align 4 [[AGG_CAPTURED]], i64 4, i1 false)432// CHECK1-NEXT: [[TMP56:%.*]] = getelementptr inbounds nuw [[STRUCT_KMP_TASK_T_WITH_PRIVATES]], ptr [[TMP52]], i32 0, i32 1433// CHECK1-NEXT: [[TMP57:%.*]] = getelementptr inbounds nuw [[STRUCT__KMP_PRIVATES_T:%.*]], ptr [[TMP56]], i32 0, i32 0434// CHECK1-NEXT: call void @llvm.memcpy.p0.p0.i64(ptr align 8 [[TMP57]], ptr align 8 [[TMP48]], i64 8, i1 false)435// CHECK1-NEXT: [[TMP58:%.*]] = getelementptr inbounds nuw [[STRUCT__KMP_PRIVATES_T]], ptr [[TMP56]], i32 0, i32 1436// CHECK1-NEXT: call void @llvm.memcpy.p0.p0.i64(ptr align 8 [[TMP58]], ptr align 8 [[TMP49]], i64 8, i1 false)437// CHECK1-NEXT: [[TMP59:%.*]] = getelementptr inbounds nuw [[STRUCT__KMP_PRIVATES_T]], ptr [[TMP56]], i32 0, i32 2438// CHECK1-NEXT: call void @llvm.memcpy.p0.p0.i64(ptr align 8 [[TMP59]], ptr align 8 @.offload_sizes.5, i64 8, i1 false)439// CHECK1-NEXT: [[TMP60:%.*]] = call i32 @__kmpc_omp_task(ptr @[[GLOB1]], i32 [[TMP0]], ptr [[TMP52]])440// CHECK1-NEXT: [[TMP61:%.*]] = load i32, ptr [[N_ADDR]], align 4441// CHECK1-NEXT: [[ADD12:%.*]] = add nsw i32 [[TMP61]], 1442// CHECK1-NEXT: ret i32 [[ADD12]]443//444//445// CHECK1-LABEL: define {{[^@]+}}@_Z9ftemplateIiET_i446// CHECK1-SAME: (i32 noundef signext [[N:%.*]]) #[[ATTR0]] comdat {447// CHECK1-NEXT: entry:448// CHECK1-NEXT: [[N_ADDR:%.*]] = alloca i32, align 4449// CHECK1-NEXT: [[A:%.*]] = alloca i32, align 4450// CHECK1-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8451// CHECK1-NEXT: [[B:%.*]] = alloca i16, align 2452// CHECK1-NEXT: [[DOTCAPTURE_EXPR_:%.*]] = alloca i16, align 2453// CHECK1-NEXT: [[A_CASTED:%.*]] = alloca i64, align 8454// CHECK1-NEXT: [[B_CASTED:%.*]] = alloca i64, align 8455// CHECK1-NEXT: [[DOTCAPTURE_EXPR__CASTED:%.*]] = alloca i64, align 8456// CHECK1-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [3 x ptr], align 8457// CHECK1-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [3 x ptr], align 8458// CHECK1-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [3 x ptr], align 8459// CHECK1-NEXT: [[KERNEL_ARGS1:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 8460// CHECK1-NEXT: store i32 [[N]], ptr [[N_ADDR]], align 4461// CHECK1-NEXT: store i32 0, ptr [[A]], align 4462// CHECK1-NEXT: [[TMP0:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0463// CHECK1-NEXT: store i32 3, ptr [[TMP0]], align 4464// CHECK1-NEXT: [[TMP1:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1465// CHECK1-NEXT: store i32 0, ptr [[TMP1]], align 4466// CHECK1-NEXT: [[TMP2:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2467// CHECK1-NEXT: store ptr null, ptr [[TMP2]], align 8468// CHECK1-NEXT: [[TMP3:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3469// CHECK1-NEXT: store ptr null, ptr [[TMP3]], align 8470// CHECK1-NEXT: [[TMP4:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4471// CHECK1-NEXT: store ptr null, ptr [[TMP4]], align 8472// CHECK1-NEXT: [[TMP5:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5473// CHECK1-NEXT: store ptr null, ptr [[TMP5]], align 8474// CHECK1-NEXT: [[TMP6:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6475// CHECK1-NEXT: store ptr null, ptr [[TMP6]], align 8476// CHECK1-NEXT: [[TMP7:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7477// CHECK1-NEXT: store ptr null, ptr [[TMP7]], align 8478// CHECK1-NEXT: [[TMP8:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8479// CHECK1-NEXT: store i64 0, ptr [[TMP8]], align 8480// CHECK1-NEXT: [[TMP9:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9481// CHECK1-NEXT: store i64 8, ptr [[TMP9]], align 8482// CHECK1-NEXT: [[TMP10:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10483// CHECK1-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP10]], align 4484// CHECK1-NEXT: [[TMP11:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11485// CHECK1-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP11]], align 4486// CHECK1-NEXT: [[TMP12:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12487// CHECK1-NEXT: store i32 20, ptr [[TMP12]], align 4488// CHECK1-NEXT: [[TMP13:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 -1, i32 0, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l55.region_id, ptr [[KERNEL_ARGS]])489// CHECK1-NEXT: [[TMP14:%.*]] = icmp ne i32 [[TMP13]], 0490// CHECK1-NEXT: br i1 [[TMP14]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]491// CHECK1: omp_offload.failed:492// CHECK1-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l55() #[[ATTR2]]493// CHECK1-NEXT: br label [[OMP_OFFLOAD_CONT]]494// CHECK1: omp_offload.cont:495// CHECK1-NEXT: store i16 1, ptr [[B]], align 2496// CHECK1-NEXT: [[TMP15:%.*]] = load i16, ptr [[B]], align 2497// CHECK1-NEXT: store i16 [[TMP15]], ptr [[DOTCAPTURE_EXPR_]], align 2498// CHECK1-NEXT: [[TMP16:%.*]] = load i32, ptr [[A]], align 4499// CHECK1-NEXT: store i32 [[TMP16]], ptr [[A_CASTED]], align 4500// CHECK1-NEXT: [[TMP17:%.*]] = load i64, ptr [[A_CASTED]], align 8501// CHECK1-NEXT: [[TMP18:%.*]] = load i16, ptr [[B]], align 2502// CHECK1-NEXT: store i16 [[TMP18]], ptr [[B_CASTED]], align 2503// CHECK1-NEXT: [[TMP19:%.*]] = load i64, ptr [[B_CASTED]], align 8504// CHECK1-NEXT: [[TMP20:%.*]] = load i16, ptr [[DOTCAPTURE_EXPR_]], align 2505// CHECK1-NEXT: store i16 [[TMP20]], ptr [[DOTCAPTURE_EXPR__CASTED]], align 2506// CHECK1-NEXT: [[TMP21:%.*]] = load i64, ptr [[DOTCAPTURE_EXPR__CASTED]], align 8507// CHECK1-NEXT: [[TMP22:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0508// CHECK1-NEXT: store i64 [[TMP17]], ptr [[TMP22]], align 8509// CHECK1-NEXT: [[TMP23:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0510// CHECK1-NEXT: store i64 [[TMP17]], ptr [[TMP23]], align 8511// CHECK1-NEXT: [[TMP24:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 0512// CHECK1-NEXT: store ptr null, ptr [[TMP24]], align 8513// CHECK1-NEXT: [[TMP25:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 1514// CHECK1-NEXT: store i64 [[TMP19]], ptr [[TMP25]], align 8515// CHECK1-NEXT: [[TMP26:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 1516// CHECK1-NEXT: store i64 [[TMP19]], ptr [[TMP26]], align 8517// CHECK1-NEXT: [[TMP27:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 1518// CHECK1-NEXT: store ptr null, ptr [[TMP27]], align 8519// CHECK1-NEXT: [[TMP28:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 2520// CHECK1-NEXT: store i64 [[TMP21]], ptr [[TMP28]], align 8521// CHECK1-NEXT: [[TMP29:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 2522// CHECK1-NEXT: store i64 [[TMP21]], ptr [[TMP29]], align 8523// CHECK1-NEXT: [[TMP30:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 2524// CHECK1-NEXT: store ptr null, ptr [[TMP30]], align 8525// CHECK1-NEXT: [[TMP31:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0526// CHECK1-NEXT: [[TMP32:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0527// CHECK1-NEXT: [[TMP33:%.*]] = load i16, ptr [[DOTCAPTURE_EXPR_]], align 2528// CHECK1-NEXT: [[TMP34:%.*]] = sext i16 [[TMP33]] to i32529// CHECK1-NEXT: [[TMP35:%.*]] = insertvalue [3 x i32] zeroinitializer, i32 [[TMP34]], 0530// CHECK1-NEXT: [[TMP36:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS1]], i32 0, i32 0531// CHECK1-NEXT: store i32 3, ptr [[TMP36]], align 4532// CHECK1-NEXT: [[TMP37:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS1]], i32 0, i32 1533// CHECK1-NEXT: store i32 3, ptr [[TMP37]], align 4534// CHECK1-NEXT: [[TMP38:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS1]], i32 0, i32 2535// CHECK1-NEXT: store ptr [[TMP31]], ptr [[TMP38]], align 8536// CHECK1-NEXT: [[TMP39:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS1]], i32 0, i32 3537// CHECK1-NEXT: store ptr [[TMP32]], ptr [[TMP39]], align 8538// CHECK1-NEXT: [[TMP40:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS1]], i32 0, i32 4539// CHECK1-NEXT: store ptr @.offload_sizes.7, ptr [[TMP40]], align 8540// CHECK1-NEXT: [[TMP41:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS1]], i32 0, i32 5541// CHECK1-NEXT: store ptr @.offload_maptypes.8, ptr [[TMP41]], align 8542// CHECK1-NEXT: [[TMP42:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS1]], i32 0, i32 6543// CHECK1-NEXT: store ptr null, ptr [[TMP42]], align 8544// CHECK1-NEXT: [[TMP43:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS1]], i32 0, i32 7545// CHECK1-NEXT: store ptr null, ptr [[TMP43]], align 8546// CHECK1-NEXT: [[TMP44:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS1]], i32 0, i32 8547// CHECK1-NEXT: store i64 0, ptr [[TMP44]], align 8548// CHECK1-NEXT: [[TMP45:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS1]], i32 0, i32 9549// CHECK1-NEXT: store i64 8, ptr [[TMP45]], align 8550// CHECK1-NEXT: [[TMP46:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS1]], i32 0, i32 10551// CHECK1-NEXT: store [3 x i32] [[TMP35]], ptr [[TMP46]], align 4552// CHECK1-NEXT: [[TMP47:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS1]], i32 0, i32 11553// CHECK1-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP47]], align 4554// CHECK1-NEXT: [[TMP48:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS1]], i32 0, i32 12555// CHECK1-NEXT: store i32 1024, ptr [[TMP48]], align 4556// CHECK1-NEXT: [[TMP49:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 -1, i32 [[TMP34]], i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l60.region_id, ptr [[KERNEL_ARGS1]])557// CHECK1-NEXT: [[TMP50:%.*]] = icmp ne i32 [[TMP49]], 0558// CHECK1-NEXT: br i1 [[TMP50]], label [[OMP_OFFLOAD_FAILED2:%.*]], label [[OMP_OFFLOAD_CONT3:%.*]]559// CHECK1: omp_offload.failed2:560// CHECK1-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l60(i64 [[TMP17]], i64 [[TMP19]], i64 [[TMP21]]) #[[ATTR2]]561// CHECK1-NEXT: br label [[OMP_OFFLOAD_CONT3]]562// CHECK1: omp_offload.cont3:563// CHECK1-NEXT: [[TMP51:%.*]] = load i32, ptr [[A]], align 4564// CHECK1-NEXT: ret i32 [[TMP51]]565//566//567// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2S12r1Ei_l88568// CHECK1-SAME: (ptr noundef [[THIS:%.*]], i64 noundef [[B:%.*]], i64 noundef [[DOTCAPTURE_EXPR_:%.*]]) #[[ATTR1:[0-9]+]] {569// CHECK1-NEXT: entry:570// CHECK1-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 8571// CHECK1-NEXT: [[B_ADDR:%.*]] = alloca i64, align 8572// CHECK1-NEXT: [[DOTCAPTURE_EXPR__ADDR:%.*]] = alloca i64, align 8573// CHECK1-NEXT: [[B_CASTED:%.*]] = alloca i64, align 8574// CHECK1-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 8575// CHECK1-NEXT: store i64 [[B]], ptr [[B_ADDR]], align 8576// CHECK1-NEXT: store i64 [[DOTCAPTURE_EXPR_]], ptr [[DOTCAPTURE_EXPR__ADDR]], align 8577// CHECK1-NEXT: [[TMP0:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8578// CHECK1-NEXT: [[TMP1:%.*]] = load i32, ptr [[B_ADDR]], align 4579// CHECK1-NEXT: store i32 [[TMP1]], ptr [[B_CASTED]], align 4580// CHECK1-NEXT: [[TMP2:%.*]] = load i64, ptr [[B_CASTED]], align 8581// CHECK1-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB1]], i32 2, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2S12r1Ei_l88.omp_outlined, ptr [[TMP0]], i64 [[TMP2]])582// CHECK1-NEXT: ret void583//584//585// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2S12r1Ei_l88.omp_outlined586// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef [[THIS:%.*]], i64 noundef [[B:%.*]]) #[[ATTR1]] {587// CHECK1-NEXT: entry:588// CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8589// CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8590// CHECK1-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 8591// CHECK1-NEXT: [[B_ADDR:%.*]] = alloca i64, align 8592// CHECK1-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8593// CHECK1-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8594// CHECK1-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 8595// CHECK1-NEXT: store i64 [[B]], ptr [[B_ADDR]], align 8596// CHECK1-NEXT: [[TMP0:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8597// CHECK1-NEXT: [[TMP1:%.*]] = load i32, ptr [[B_ADDR]], align 4598// CHECK1-NEXT: [[CONV:%.*]] = sitofp i32 [[TMP1]] to double599// CHECK1-NEXT: [[ADD:%.*]] = fadd double [[CONV]], 1.500000e+00600// CHECK1-NEXT: [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_S1:%.*]], ptr [[TMP0]], i32 0, i32 0601// CHECK1-NEXT: store double [[ADD]], ptr [[A]], align 8602// CHECK1-NEXT: ret void603//604//605// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2S12r1Ei_l93606// CHECK1-SAME: (ptr noundef [[THIS:%.*]]) #[[ATTR1]] {607// CHECK1-NEXT: entry:608// CHECK1-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 8609// CHECK1-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 8610// CHECK1-NEXT: [[TMP0:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8611// CHECK1-NEXT: [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_S1:%.*]], ptr [[TMP0]], i32 0, i32 0612// CHECK1-NEXT: store double 2.500000e+00, ptr [[A]], align 8613// CHECK1-NEXT: ret void614//615//616// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l71617// CHECK1-SAME: (i64 noundef [[N:%.*]], i64 noundef [[DOTCAPTURE_EXPR_:%.*]], i64 noundef [[DOTCAPTURE_EXPR_1:%.*]]) #[[ATTR1]] {618// CHECK1-NEXT: entry:619// CHECK1-NEXT: [[N_ADDR:%.*]] = alloca i64, align 8620// CHECK1-NEXT: [[DOTCAPTURE_EXPR__ADDR:%.*]] = alloca i64, align 8621// CHECK1-NEXT: [[DOTCAPTURE_EXPR__ADDR2:%.*]] = alloca i64, align 8622// CHECK1-NEXT: [[N_CASTED:%.*]] = alloca i64, align 8623// CHECK1-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])624// CHECK1-NEXT: store i64 [[N]], ptr [[N_ADDR]], align 8625// CHECK1-NEXT: store i64 [[DOTCAPTURE_EXPR_]], ptr [[DOTCAPTURE_EXPR__ADDR]], align 8626// CHECK1-NEXT: store i64 [[DOTCAPTURE_EXPR_1]], ptr [[DOTCAPTURE_EXPR__ADDR2]], align 8627// CHECK1-NEXT: [[TMP1:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR__ADDR]], align 4628// CHECK1-NEXT: call void @__kmpc_push_num_teams(ptr @[[GLOB1]], i32 [[TMP0]], i32 [[TMP1]], i32 0)629// CHECK1-NEXT: [[TMP2:%.*]] = load i32, ptr [[N_ADDR]], align 4630// CHECK1-NEXT: store i32 [[TMP2]], ptr [[N_CASTED]], align 4631// CHECK1-NEXT: [[TMP3:%.*]] = load i64, ptr [[N_CASTED]], align 8632// CHECK1-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB1]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l71.omp_outlined, i64 [[TMP3]])633// CHECK1-NEXT: ret void634//635//636// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l71.omp_outlined637// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i64 noundef [[N:%.*]]) #[[ATTR1]] {638// CHECK1-NEXT: entry:639// CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8640// CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8641// CHECK1-NEXT: [[N_ADDR:%.*]] = alloca i64, align 8642// CHECK1-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 4643// CHECK1-NEXT: [[TMP:%.*]] = alloca i32, align 4644// CHECK1-NEXT: [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 4645// CHECK1-NEXT: [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 4646// CHECK1-NEXT: [[I:%.*]] = alloca i32, align 4647// CHECK1-NEXT: [[DOTOMP_COMB_LB:%.*]] = alloca i32, align 4648// CHECK1-NEXT: [[DOTOMP_COMB_UB:%.*]] = alloca i32, align 4649// CHECK1-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4650// CHECK1-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4651// CHECK1-NEXT: [[I3:%.*]] = alloca i32, align 4652// CHECK1-NEXT: [[N_CASTED:%.*]] = alloca i64, align 8653// CHECK1-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8654// CHECK1-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8655// CHECK1-NEXT: store i64 [[N]], ptr [[N_ADDR]], align 8656// CHECK1-NEXT: [[TMP0:%.*]] = load i32, ptr [[N_ADDR]], align 4657// CHECK1-NEXT: store i32 [[TMP0]], ptr [[DOTCAPTURE_EXPR_]], align 4658// CHECK1-NEXT: [[TMP1:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4659// CHECK1-NEXT: [[SUB:%.*]] = sub nsw i32 [[TMP1]], 0660// CHECK1-NEXT: [[DIV:%.*]] = sdiv i32 [[SUB]], 1661// CHECK1-NEXT: [[SUB2:%.*]] = sub nsw i32 [[DIV]], 1662// CHECK1-NEXT: store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 4663// CHECK1-NEXT: store i32 0, ptr [[I]], align 4664// CHECK1-NEXT: [[TMP2:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4665// CHECK1-NEXT: [[CMP:%.*]] = icmp slt i32 0, [[TMP2]]666// CHECK1-NEXT: br i1 [[CMP]], label [[OMP_PRECOND_THEN:%.*]], label [[OMP_PRECOND_END:%.*]]667// CHECK1: omp.precond.then:668// CHECK1-NEXT: store i32 0, ptr [[DOTOMP_COMB_LB]], align 4669// CHECK1-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4670// CHECK1-NEXT: store i32 [[TMP3]], ptr [[DOTOMP_COMB_UB]], align 4671// CHECK1-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 4672// CHECK1-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 4673// CHECK1-NEXT: [[TMP4:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8674// CHECK1-NEXT: [[TMP5:%.*]] = load i32, ptr [[TMP4]], align 4675// CHECK1-NEXT: call void @__kmpc_for_static_init_4(ptr @[[GLOB2:[0-9]+]], i32 [[TMP5]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_COMB_LB]], ptr [[DOTOMP_COMB_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)676// CHECK1-NEXT: [[TMP6:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4677// CHECK1-NEXT: [[TMP7:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4678// CHECK1-NEXT: [[CMP4:%.*]] = icmp sgt i32 [[TMP6]], [[TMP7]]679// CHECK1-NEXT: br i1 [[CMP4]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]680// CHECK1: cond.true:681// CHECK1-NEXT: [[TMP8:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4682// CHECK1-NEXT: br label [[COND_END:%.*]]683// CHECK1: cond.false:684// CHECK1-NEXT: [[TMP9:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4685// CHECK1-NEXT: br label [[COND_END]]686// CHECK1: cond.end:687// CHECK1-NEXT: [[COND:%.*]] = phi i32 [ [[TMP8]], [[COND_TRUE]] ], [ [[TMP9]], [[COND_FALSE]] ]688// CHECK1-NEXT: store i32 [[COND]], ptr [[DOTOMP_COMB_UB]], align 4689// CHECK1-NEXT: [[TMP10:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 4690// CHECK1-NEXT: store i32 [[TMP10]], ptr [[DOTOMP_IV]], align 4691// CHECK1-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]692// CHECK1: omp.inner.for.cond:693// CHECK1-NEXT: [[TMP11:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP17:![0-9]+]]694// CHECK1-NEXT: [[TMP12:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4, !llvm.access.group [[ACC_GRP17]]695// CHECK1-NEXT: [[CMP5:%.*]] = icmp sle i32 [[TMP11]], [[TMP12]]696// CHECK1-NEXT: br i1 [[CMP5]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]697// CHECK1: omp.inner.for.body:698// CHECK1-NEXT: [[TMP13:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 4, !llvm.access.group [[ACC_GRP17]]699// CHECK1-NEXT: [[TMP14:%.*]] = zext i32 [[TMP13]] to i64700// CHECK1-NEXT: [[TMP15:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4, !llvm.access.group [[ACC_GRP17]]701// CHECK1-NEXT: [[TMP16:%.*]] = zext i32 [[TMP15]] to i64702// CHECK1-NEXT: [[TMP17:%.*]] = load i32, ptr [[N_ADDR]], align 4, !llvm.access.group [[ACC_GRP17]]703// CHECK1-NEXT: store i32 [[TMP17]], ptr [[N_CASTED]], align 4, !llvm.access.group [[ACC_GRP17]]704// CHECK1-NEXT: [[TMP18:%.*]] = load i64, ptr [[N_CASTED]], align 8, !llvm.access.group [[ACC_GRP17]]705// CHECK1-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB1]], i32 3, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l71.omp_outlined.omp_outlined, i64 [[TMP14]], i64 [[TMP16]], i64 [[TMP18]]), !llvm.access.group [[ACC_GRP17]]706// CHECK1-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]707// CHECK1: omp.inner.for.inc:708// CHECK1-NEXT: [[TMP19:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP17]]709// CHECK1-NEXT: [[TMP20:%.*]] = load i32, ptr [[DOTOMP_STRIDE]], align 4, !llvm.access.group [[ACC_GRP17]]710// CHECK1-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP19]], [[TMP20]]711// CHECK1-NEXT: store i32 [[ADD]], ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP17]]712// CHECK1-NEXT: br label [[OMP_INNER_FOR_COND]], !llvm.loop [[LOOP18:![0-9]+]]713// CHECK1: omp.inner.for.end:714// CHECK1-NEXT: br label [[OMP_LOOP_EXIT:%.*]]715// CHECK1: omp.loop.exit:716// CHECK1-NEXT: [[TMP21:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8717// CHECK1-NEXT: [[TMP22:%.*]] = load i32, ptr [[TMP21]], align 4718// CHECK1-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB2]], i32 [[TMP22]])719// CHECK1-NEXT: [[TMP23:%.*]] = load i32, ptr [[DOTOMP_IS_LAST]], align 4720// CHECK1-NEXT: [[TMP24:%.*]] = icmp ne i32 [[TMP23]], 0721// CHECK1-NEXT: br i1 [[TMP24]], label [[DOTOMP_FINAL_THEN:%.*]], label [[DOTOMP_FINAL_DONE:%.*]]722// CHECK1: .omp.final.then:723// CHECK1-NEXT: [[TMP25:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4724// CHECK1-NEXT: [[SUB6:%.*]] = sub nsw i32 [[TMP25]], 0725// CHECK1-NEXT: [[DIV7:%.*]] = sdiv i32 [[SUB6]], 1726// CHECK1-NEXT: [[MUL:%.*]] = mul nsw i32 [[DIV7]], 1727// CHECK1-NEXT: [[ADD8:%.*]] = add nsw i32 0, [[MUL]]728// CHECK1-NEXT: store i32 [[ADD8]], ptr [[I3]], align 4729// CHECK1-NEXT: br label [[DOTOMP_FINAL_DONE]]730// CHECK1: .omp.final.done:731// CHECK1-NEXT: br label [[OMP_PRECOND_END]]732// CHECK1: omp.precond.end:733// CHECK1-NEXT: ret void734//735//736// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l71.omp_outlined.omp_outlined737// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i64 noundef [[DOTPREVIOUS_LB_:%.*]], i64 noundef [[DOTPREVIOUS_UB_:%.*]], i64 noundef [[N:%.*]]) #[[ATTR1]] {738// CHECK1-NEXT: entry:739// CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8740// CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8741// CHECK1-NEXT: [[DOTPREVIOUS_LB__ADDR:%.*]] = alloca i64, align 8742// CHECK1-NEXT: [[DOTPREVIOUS_UB__ADDR:%.*]] = alloca i64, align 8743// CHECK1-NEXT: [[N_ADDR:%.*]] = alloca i64, align 8744// CHECK1-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 4745// CHECK1-NEXT: [[TMP:%.*]] = alloca i32, align 4746// CHECK1-NEXT: [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 4747// CHECK1-NEXT: [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 4748// CHECK1-NEXT: [[I:%.*]] = alloca i32, align 4749// CHECK1-NEXT: [[DOTOMP_LB:%.*]] = alloca i32, align 4750// CHECK1-NEXT: [[DOTOMP_UB:%.*]] = alloca i32, align 4751// CHECK1-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4752// CHECK1-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4753// CHECK1-NEXT: [[I4:%.*]] = alloca i32, align 4754// CHECK1-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8755// CHECK1-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8756// CHECK1-NEXT: store i64 [[DOTPREVIOUS_LB_]], ptr [[DOTPREVIOUS_LB__ADDR]], align 8757// CHECK1-NEXT: store i64 [[DOTPREVIOUS_UB_]], ptr [[DOTPREVIOUS_UB__ADDR]], align 8758// CHECK1-NEXT: store i64 [[N]], ptr [[N_ADDR]], align 8759// CHECK1-NEXT: [[TMP0:%.*]] = load i32, ptr [[N_ADDR]], align 4760// CHECK1-NEXT: store i32 [[TMP0]], ptr [[DOTCAPTURE_EXPR_]], align 4761// CHECK1-NEXT: [[TMP1:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4762// CHECK1-NEXT: [[SUB:%.*]] = sub nsw i32 [[TMP1]], 0763// CHECK1-NEXT: [[DIV:%.*]] = sdiv i32 [[SUB]], 1764// CHECK1-NEXT: [[SUB2:%.*]] = sub nsw i32 [[DIV]], 1765// CHECK1-NEXT: store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 4766// CHECK1-NEXT: store i32 0, ptr [[I]], align 4767// CHECK1-NEXT: [[TMP2:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4768// CHECK1-NEXT: [[CMP:%.*]] = icmp slt i32 0, [[TMP2]]769// CHECK1-NEXT: br i1 [[CMP]], label [[OMP_PRECOND_THEN:%.*]], label [[OMP_PRECOND_END:%.*]]770// CHECK1: omp.precond.then:771// CHECK1-NEXT: store i32 0, ptr [[DOTOMP_LB]], align 4772// CHECK1-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4773// CHECK1-NEXT: store i32 [[TMP3]], ptr [[DOTOMP_UB]], align 4774// CHECK1-NEXT: [[TMP4:%.*]] = load i64, ptr [[DOTPREVIOUS_LB__ADDR]], align 8775// CHECK1-NEXT: [[CONV:%.*]] = trunc i64 [[TMP4]] to i32776// CHECK1-NEXT: [[TMP5:%.*]] = load i64, ptr [[DOTPREVIOUS_UB__ADDR]], align 8777// CHECK1-NEXT: [[CONV3:%.*]] = trunc i64 [[TMP5]] to i32778// CHECK1-NEXT: store i32 [[CONV]], ptr [[DOTOMP_LB]], align 4779// CHECK1-NEXT: store i32 [[CONV3]], ptr [[DOTOMP_UB]], align 4780// CHECK1-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 4781// CHECK1-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 4782// CHECK1-NEXT: [[TMP6:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8783// CHECK1-NEXT: [[TMP7:%.*]] = load i32, ptr [[TMP6]], align 4784// CHECK1-NEXT: call void @__kmpc_for_static_init_4(ptr @[[GLOB3:[0-9]+]], i32 [[TMP7]], i32 34, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)785// CHECK1-NEXT: [[TMP8:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4786// CHECK1-NEXT: [[TMP9:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4787// CHECK1-NEXT: [[CMP5:%.*]] = icmp sgt i32 [[TMP8]], [[TMP9]]788// CHECK1-NEXT: br i1 [[CMP5]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]789// CHECK1: cond.true:790// CHECK1-NEXT: [[TMP10:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4791// CHECK1-NEXT: br label [[COND_END:%.*]]792// CHECK1: cond.false:793// CHECK1-NEXT: [[TMP11:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4794// CHECK1-NEXT: br label [[COND_END]]795// CHECK1: cond.end:796// CHECK1-NEXT: [[COND:%.*]] = phi i32 [ [[TMP10]], [[COND_TRUE]] ], [ [[TMP11]], [[COND_FALSE]] ]797// CHECK1-NEXT: store i32 [[COND]], ptr [[DOTOMP_UB]], align 4798// CHECK1-NEXT: [[TMP12:%.*]] = load i32, ptr [[DOTOMP_LB]], align 4799// CHECK1-NEXT: store i32 [[TMP12]], ptr [[DOTOMP_IV]], align 4800// CHECK1-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]801// CHECK1: omp.inner.for.cond:802// CHECK1-NEXT: [[TMP13:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP21:![0-9]+]]803// CHECK1-NEXT: [[TMP14:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4, !llvm.access.group [[ACC_GRP21]]804// CHECK1-NEXT: [[CMP6:%.*]] = icmp sle i32 [[TMP13]], [[TMP14]]805// CHECK1-NEXT: br i1 [[CMP6]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]806// CHECK1: omp.inner.for.body:807// CHECK1-NEXT: [[TMP15:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP21]]808// CHECK1-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP15]], 1809// CHECK1-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]810// CHECK1-NEXT: store i32 [[ADD]], ptr [[I4]], align 4, !llvm.access.group [[ACC_GRP21]]811// CHECK1-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]812// CHECK1: omp.body.continue:813// CHECK1-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]814// CHECK1: omp.inner.for.inc:815// CHECK1-NEXT: [[TMP16:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP21]]816// CHECK1-NEXT: [[ADD7:%.*]] = add nsw i32 [[TMP16]], 1817// CHECK1-NEXT: store i32 [[ADD7]], ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP21]]818// CHECK1-NEXT: br label [[OMP_INNER_FOR_COND]], !llvm.loop [[LOOP22:![0-9]+]]819// CHECK1: omp.inner.for.end:820// CHECK1-NEXT: br label [[OMP_LOOP_EXIT:%.*]]821// CHECK1: omp.loop.exit:822// CHECK1-NEXT: [[TMP17:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8823// CHECK1-NEXT: [[TMP18:%.*]] = load i32, ptr [[TMP17]], align 4824// CHECK1-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB3]], i32 [[TMP18]])825// CHECK1-NEXT: [[TMP19:%.*]] = load i32, ptr [[DOTOMP_IS_LAST]], align 4826// CHECK1-NEXT: [[TMP20:%.*]] = icmp ne i32 [[TMP19]], 0827// CHECK1-NEXT: br i1 [[TMP20]], label [[DOTOMP_FINAL_THEN:%.*]], label [[DOTOMP_FINAL_DONE:%.*]]828// CHECK1: .omp.final.then:829// CHECK1-NEXT: [[TMP21:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4830// CHECK1-NEXT: [[SUB8:%.*]] = sub nsw i32 [[TMP21]], 0831// CHECK1-NEXT: [[DIV9:%.*]] = sdiv i32 [[SUB8]], 1832// CHECK1-NEXT: [[MUL10:%.*]] = mul nsw i32 [[DIV9]], 1833// CHECK1-NEXT: [[ADD11:%.*]] = add nsw i32 0, [[MUL10]]834// CHECK1-NEXT: store i32 [[ADD11]], ptr [[I4]], align 4835// CHECK1-NEXT: br label [[DOTOMP_FINAL_DONE]]836// CHECK1: .omp.final.done:837// CHECK1-NEXT: br label [[OMP_PRECOND_END]]838// CHECK1: omp.precond.end:839// CHECK1-NEXT: ret void840//841//842// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l75843// CHECK1-SAME: (i64 noundef [[DOTCAPTURE_EXPR_:%.*]]) #[[ATTR1]] {844// CHECK1-NEXT: entry:845// CHECK1-NEXT: [[DOTCAPTURE_EXPR__ADDR:%.*]] = alloca i64, align 8846// CHECK1-NEXT: store i64 [[DOTCAPTURE_EXPR_]], ptr [[DOTCAPTURE_EXPR__ADDR]], align 8847// CHECK1-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB1]], i32 0, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l75.omp_outlined)848// CHECK1-NEXT: ret void849//850//851// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l75.omp_outlined852// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR1]] {853// CHECK1-NEXT: entry:854// CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8855// CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8856// CHECK1-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8857// CHECK1-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8858// CHECK1-NEXT: ret void859//860//861// CHECK1-LABEL: define {{[^@]+}}@.omp_task_privates_map.862// CHECK1-SAME: (ptr noalias noundef [[TMP0:%.*]], ptr noalias noundef [[TMP1:%.*]], ptr noalias noundef [[TMP2:%.*]], ptr noalias noundef [[TMP3:%.*]]) #[[ATTR3:[0-9]+]] {863// CHECK1-NEXT: entry:864// CHECK1-NEXT: [[DOTADDR:%.*]] = alloca ptr, align 8865// CHECK1-NEXT: [[DOTADDR1:%.*]] = alloca ptr, align 8866// CHECK1-NEXT: [[DOTADDR2:%.*]] = alloca ptr, align 8867// CHECK1-NEXT: [[DOTADDR3:%.*]] = alloca ptr, align 8868// CHECK1-NEXT: store ptr [[TMP0]], ptr [[DOTADDR]], align 8869// CHECK1-NEXT: store ptr [[TMP1]], ptr [[DOTADDR1]], align 8870// CHECK1-NEXT: store ptr [[TMP2]], ptr [[DOTADDR2]], align 8871// CHECK1-NEXT: store ptr [[TMP3]], ptr [[DOTADDR3]], align 8872// CHECK1-NEXT: [[TMP4:%.*]] = load ptr, ptr [[DOTADDR]], align 8873// CHECK1-NEXT: [[TMP5:%.*]] = getelementptr inbounds nuw [[STRUCT__KMP_PRIVATES_T:%.*]], ptr [[TMP4]], i32 0, i32 0874// CHECK1-NEXT: [[TMP6:%.*]] = load ptr, ptr [[DOTADDR1]], align 8875// CHECK1-NEXT: store ptr [[TMP5]], ptr [[TMP6]], align 8876// CHECK1-NEXT: [[TMP7:%.*]] = getelementptr inbounds nuw [[STRUCT__KMP_PRIVATES_T]], ptr [[TMP4]], i32 0, i32 1877// CHECK1-NEXT: [[TMP8:%.*]] = load ptr, ptr [[DOTADDR2]], align 8878// CHECK1-NEXT: store ptr [[TMP7]], ptr [[TMP8]], align 8879// CHECK1-NEXT: [[TMP9:%.*]] = getelementptr inbounds nuw [[STRUCT__KMP_PRIVATES_T]], ptr [[TMP4]], i32 0, i32 2880// CHECK1-NEXT: [[TMP10:%.*]] = load ptr, ptr [[DOTADDR3]], align 8881// CHECK1-NEXT: store ptr [[TMP9]], ptr [[TMP10]], align 8882// CHECK1-NEXT: ret void883//884//885// CHECK1-LABEL: define {{[^@]+}}@.omp_task_entry.886// CHECK1-SAME: (i32 noundef signext [[TMP0:%.*]], ptr noalias noundef [[TMP1:%.*]]) #[[ATTR4:[0-9]+]] {887// CHECK1-NEXT: entry:888// CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR_I:%.*]] = alloca i32, align 4889// CHECK1-NEXT: [[DOTPART_ID__ADDR_I:%.*]] = alloca ptr, align 8890// CHECK1-NEXT: [[DOTPRIVATES__ADDR_I:%.*]] = alloca ptr, align 8891// CHECK1-NEXT: [[DOTCOPY_FN__ADDR_I:%.*]] = alloca ptr, align 8892// CHECK1-NEXT: [[DOTTASK_T__ADDR_I:%.*]] = alloca ptr, align 8893// CHECK1-NEXT: [[__CONTEXT_ADDR_I:%.*]] = alloca ptr, align 8894// CHECK1-NEXT: [[DOTFIRSTPRIV_PTR_ADDR_I:%.*]] = alloca ptr, align 8895// CHECK1-NEXT: [[DOTFIRSTPRIV_PTR_ADDR1_I:%.*]] = alloca ptr, align 8896// CHECK1-NEXT: [[DOTFIRSTPRIV_PTR_ADDR2_I:%.*]] = alloca ptr, align 8897// CHECK1-NEXT: [[KERNEL_ARGS_I:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8898// CHECK1-NEXT: [[DOTCAPTURE_EXPR__CASTED_I:%.*]] = alloca i64, align 8899// CHECK1-NEXT: [[DOTADDR:%.*]] = alloca i32, align 4900// CHECK1-NEXT: [[DOTADDR1:%.*]] = alloca ptr, align 8901// CHECK1-NEXT: store i32 [[TMP0]], ptr [[DOTADDR]], align 4902// CHECK1-NEXT: store ptr [[TMP1]], ptr [[DOTADDR1]], align 8903// CHECK1-NEXT: [[TMP2:%.*]] = load i32, ptr [[DOTADDR]], align 4904// CHECK1-NEXT: [[TMP3:%.*]] = load ptr, ptr [[DOTADDR1]], align 8905// CHECK1-NEXT: [[TMP4:%.*]] = getelementptr inbounds nuw [[STRUCT_KMP_TASK_T_WITH_PRIVATES:%.*]], ptr [[TMP3]], i32 0, i32 0906// CHECK1-NEXT: [[TMP5:%.*]] = getelementptr inbounds nuw [[STRUCT_KMP_TASK_T:%.*]], ptr [[TMP4]], i32 0, i32 2907// CHECK1-NEXT: [[TMP6:%.*]] = getelementptr inbounds nuw [[STRUCT_KMP_TASK_T]], ptr [[TMP4]], i32 0, i32 0908// CHECK1-NEXT: [[TMP7:%.*]] = load ptr, ptr [[TMP6]], align 8909// CHECK1-NEXT: [[TMP8:%.*]] = getelementptr inbounds nuw [[STRUCT_KMP_TASK_T_WITH_PRIVATES]], ptr [[TMP3]], i32 0, i32 1910// CHECK1-NEXT: call void @llvm.experimental.noalias.scope.decl(metadata [[META24:![0-9]+]])911// CHECK1-NEXT: call void @llvm.experimental.noalias.scope.decl(metadata [[META27:![0-9]+]])912// CHECK1-NEXT: call void @llvm.experimental.noalias.scope.decl(metadata [[META29:![0-9]+]])913// CHECK1-NEXT: call void @llvm.experimental.noalias.scope.decl(metadata [[META31:![0-9]+]])914// CHECK1-NEXT: store i32 [[TMP2]], ptr [[DOTGLOBAL_TID__ADDR_I]], align 4, !noalias [[META33:![0-9]+]]915// CHECK1-NEXT: store ptr [[TMP5]], ptr [[DOTPART_ID__ADDR_I]], align 8, !noalias [[META33]]916// CHECK1-NEXT: store ptr [[TMP8]], ptr [[DOTPRIVATES__ADDR_I]], align 8, !noalias [[META33]]917// CHECK1-NEXT: store ptr @.omp_task_privates_map., ptr [[DOTCOPY_FN__ADDR_I]], align 8, !noalias [[META33]]918// CHECK1-NEXT: store ptr [[TMP3]], ptr [[DOTTASK_T__ADDR_I]], align 8, !noalias [[META33]]919// CHECK1-NEXT: store ptr [[TMP7]], ptr [[__CONTEXT_ADDR_I]], align 8, !noalias [[META33]]920// CHECK1-NEXT: [[TMP9:%.*]] = load ptr, ptr [[__CONTEXT_ADDR_I]], align 8, !noalias [[META33]]921// CHECK1-NEXT: [[TMP10:%.*]] = load ptr, ptr [[DOTCOPY_FN__ADDR_I]], align 8, !noalias [[META33]]922// CHECK1-NEXT: [[TMP11:%.*]] = load ptr, ptr [[DOTPRIVATES__ADDR_I]], align 8, !noalias [[META33]]923// CHECK1-NEXT: call void [[TMP10]](ptr [[TMP11]], ptr [[DOTFIRSTPRIV_PTR_ADDR_I]], ptr [[DOTFIRSTPRIV_PTR_ADDR1_I]], ptr [[DOTFIRSTPRIV_PTR_ADDR2_I]]) #[[ATTR2]]924// CHECK1-NEXT: [[TMP12:%.*]] = load ptr, ptr [[DOTFIRSTPRIV_PTR_ADDR_I]], align 8, !noalias [[META33]]925// CHECK1-NEXT: [[TMP13:%.*]] = load ptr, ptr [[DOTFIRSTPRIV_PTR_ADDR1_I]], align 8, !noalias [[META33]]926// CHECK1-NEXT: [[TMP14:%.*]] = load ptr, ptr [[DOTFIRSTPRIV_PTR_ADDR2_I]], align 8, !noalias [[META33]]927// CHECK1-NEXT: [[TMP15:%.*]] = load i32, ptr [[TMP9]], align 4928// CHECK1-NEXT: store i32 3, ptr [[KERNEL_ARGS_I]], align 4, !noalias [[META33]]929// CHECK1-NEXT: [[TMP16:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS_I]], i32 0, i32 1930// CHECK1-NEXT: store i32 1, ptr [[TMP16]], align 4, !noalias [[META33]]931// CHECK1-NEXT: [[TMP17:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS_I]], i32 0, i32 2932// CHECK1-NEXT: store ptr [[TMP12]], ptr [[TMP17]], align 8, !noalias [[META33]]933// CHECK1-NEXT: [[TMP18:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS_I]], i32 0, i32 3934// CHECK1-NEXT: store ptr [[TMP13]], ptr [[TMP18]], align 8, !noalias [[META33]]935// CHECK1-NEXT: [[TMP19:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS_I]], i32 0, i32 4936// CHECK1-NEXT: store ptr [[TMP14]], ptr [[TMP19]], align 8, !noalias [[META33]]937// CHECK1-NEXT: [[TMP20:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS_I]], i32 0, i32 5938// CHECK1-NEXT: store ptr @.offload_maptypes.6, ptr [[TMP20]], align 8, !noalias [[META33]]939// CHECK1-NEXT: [[TMP21:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS_I]], i32 0, i32 6940// CHECK1-NEXT: store ptr null, ptr [[TMP21]], align 8, !noalias [[META33]]941// CHECK1-NEXT: [[TMP22:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS_I]], i32 0, i32 7942// CHECK1-NEXT: store ptr null, ptr [[TMP22]], align 8, !noalias [[META33]]943// CHECK1-NEXT: [[TMP23:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS_I]], i32 0, i32 8944// CHECK1-NEXT: store i64 0, ptr [[TMP23]], align 8, !noalias [[META33]]945// CHECK1-NEXT: [[TMP24:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS_I]], i32 0, i32 9946// CHECK1-NEXT: store i64 9, ptr [[TMP24]], align 8, !noalias [[META33]]947// CHECK1-NEXT: [[TMP25:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS_I]], i32 0, i32 10948// CHECK1-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP25]], align 4, !noalias [[META33]]949// CHECK1-NEXT: [[TMP26:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS_I]], i32 0, i32 11950// CHECK1-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP26]], align 4, !noalias [[META33]]951// CHECK1-NEXT: [[TMP27:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS_I]], i32 0, i32 12952// CHECK1-NEXT: store i32 [[TMP15]], ptr [[TMP27]], align 4, !noalias [[META33]]953// CHECK1-NEXT: [[TMP28:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 -1, i32 0, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l75.region_id, ptr [[KERNEL_ARGS_I]])954// CHECK1-NEXT: [[TMP29:%.*]] = icmp ne i32 [[TMP28]], 0955// CHECK1-NEXT: br i1 [[TMP29]], label [[OMP_OFFLOAD_FAILED_I:%.*]], label [[DOTOMP_OUTLINED__EXIT:%.*]]956// CHECK1: omp_offload.failed.i:957// CHECK1-NEXT: [[TMP30:%.*]] = load i32, ptr [[TMP9]], align 4958// CHECK1-NEXT: store i32 [[TMP30]], ptr [[DOTCAPTURE_EXPR__CASTED_I]], align 4, !noalias [[META33]]959// CHECK1-NEXT: [[TMP31:%.*]] = load i64, ptr [[DOTCAPTURE_EXPR__CASTED_I]], align 8, !noalias [[META33]]960// CHECK1-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l75(i64 [[TMP31]]) #[[ATTR2]]961// CHECK1-NEXT: br label [[DOTOMP_OUTLINED__EXIT]]962// CHECK1: .omp_outlined..exit:963// CHECK1-NEXT: ret i32 0964//965//966// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l55967// CHECK1-SAME: () #[[ATTR1]] {968// CHECK1-NEXT: entry:969// CHECK1-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB1]], i32 0, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l55.omp_outlined)970// CHECK1-NEXT: ret void971//972//973// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l55.omp_outlined974// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR1]] {975// CHECK1-NEXT: entry:976// CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8977// CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8978// CHECK1-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8979// CHECK1-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8980// CHECK1-NEXT: ret void981//982//983// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l60984// CHECK1-SAME: (i64 noundef [[A:%.*]], i64 noundef [[B:%.*]], i64 noundef [[DOTCAPTURE_EXPR_:%.*]]) #[[ATTR1]] {985// CHECK1-NEXT: entry:986// CHECK1-NEXT: [[A_ADDR:%.*]] = alloca i64, align 8987// CHECK1-NEXT: [[B_ADDR:%.*]] = alloca i64, align 8988// CHECK1-NEXT: [[DOTCAPTURE_EXPR__ADDR:%.*]] = alloca i64, align 8989// CHECK1-NEXT: [[A_CASTED:%.*]] = alloca i64, align 8990// CHECK1-NEXT: [[B_CASTED:%.*]] = alloca i64, align 8991// CHECK1-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])992// CHECK1-NEXT: store i64 [[A]], ptr [[A_ADDR]], align 8993// CHECK1-NEXT: store i64 [[B]], ptr [[B_ADDR]], align 8994// CHECK1-NEXT: store i64 [[DOTCAPTURE_EXPR_]], ptr [[DOTCAPTURE_EXPR__ADDR]], align 8995// CHECK1-NEXT: [[TMP1:%.*]] = load i16, ptr [[DOTCAPTURE_EXPR__ADDR]], align 2996// CHECK1-NEXT: [[TMP2:%.*]] = sext i16 [[TMP1]] to i32997// CHECK1-NEXT: call void @__kmpc_push_num_teams(ptr @[[GLOB1]], i32 [[TMP0]], i32 [[TMP2]], i32 0)998// CHECK1-NEXT: [[TMP3:%.*]] = load i32, ptr [[A_ADDR]], align 4999// CHECK1-NEXT: store i32 [[TMP3]], ptr [[A_CASTED]], align 41000// CHECK1-NEXT: [[TMP4:%.*]] = load i64, ptr [[A_CASTED]], align 81001// CHECK1-NEXT: [[TMP5:%.*]] = load i16, ptr [[B_ADDR]], align 21002// CHECK1-NEXT: store i16 [[TMP5]], ptr [[B_CASTED]], align 21003// CHECK1-NEXT: [[TMP6:%.*]] = load i64, ptr [[B_CASTED]], align 81004// CHECK1-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB1]], i32 2, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l60.omp_outlined, i64 [[TMP4]], i64 [[TMP6]])1005// CHECK1-NEXT: ret void1006//1007//1008// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l60.omp_outlined1009// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i64 noundef [[A:%.*]], i64 noundef [[B:%.*]]) #[[ATTR1]] {1010// CHECK1-NEXT: entry:1011// CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 81012// CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 81013// CHECK1-NEXT: [[A_ADDR:%.*]] = alloca i64, align 81014// CHECK1-NEXT: [[B_ADDR:%.*]] = alloca i64, align 81015// CHECK1-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 81016// CHECK1-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 81017// CHECK1-NEXT: store i64 [[A]], ptr [[A_ADDR]], align 81018// CHECK1-NEXT: store i64 [[B]], ptr [[B_ADDR]], align 81019// CHECK1-NEXT: [[TMP0:%.*]] = load i16, ptr [[B_ADDR]], align 21020// CHECK1-NEXT: [[CONV:%.*]] = sext i16 [[TMP0]] to i321021// CHECK1-NEXT: [[TMP1:%.*]] = load i32, ptr [[A_ADDR]], align 41022// CHECK1-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP1]], [[CONV]]1023// CHECK1-NEXT: store i32 [[ADD]], ptr [[A_ADDR]], align 41024// CHECK1-NEXT: ret void1025//1026//1027// CHECK3-LABEL: define {{[^@]+}}@_Z3bari1028// CHECK3-SAME: (i32 noundef [[N:%.*]]) #[[ATTR0:[0-9]+]] {1029// CHECK3-NEXT: entry:1030// CHECK3-NEXT: [[N_ADDR:%.*]] = alloca i32, align 41031// CHECK3-NEXT: [[A:%.*]] = alloca i32, align 41032// CHECK3-NEXT: [[S:%.*]] = alloca [[STRUCT_S1:%.*]], align 41033// CHECK3-NEXT: store i32 [[N]], ptr [[N_ADDR]], align 41034// CHECK3-NEXT: store i32 0, ptr [[A]], align 41035// CHECK3-NEXT: [[TMP0:%.*]] = load i32, ptr [[N_ADDR]], align 41036// CHECK3-NEXT: [[CALL:%.*]] = call noundef i32 @_ZN2S12r1Ei(ptr noundef nonnull align 4 dereferenceable(8) [[S]], i32 noundef [[TMP0]])1037// CHECK3-NEXT: [[TMP1:%.*]] = load i32, ptr [[A]], align 41038// CHECK3-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP1]], [[CALL]]1039// CHECK3-NEXT: store i32 [[ADD]], ptr [[A]], align 41040// CHECK3-NEXT: [[TMP2:%.*]] = load i32, ptr [[N_ADDR]], align 41041// CHECK3-NEXT: [[CALL1:%.*]] = call noundef i32 @_ZL7fstatici(i32 noundef [[TMP2]])1042// CHECK3-NEXT: [[TMP3:%.*]] = load i32, ptr [[A]], align 41043// CHECK3-NEXT: [[ADD2:%.*]] = add nsw i32 [[TMP3]], [[CALL1]]1044// CHECK3-NEXT: store i32 [[ADD2]], ptr [[A]], align 41045// CHECK3-NEXT: [[TMP4:%.*]] = load i32, ptr [[N_ADDR]], align 41046// CHECK3-NEXT: [[CALL3:%.*]] = call noundef i32 @_Z9ftemplateIiET_i(i32 noundef [[TMP4]])1047// CHECK3-NEXT: [[TMP5:%.*]] = load i32, ptr [[A]], align 41048// CHECK3-NEXT: [[ADD4:%.*]] = add nsw i32 [[TMP5]], [[CALL3]]1049// CHECK3-NEXT: store i32 [[ADD4]], ptr [[A]], align 41050// CHECK3-NEXT: [[TMP6:%.*]] = load i32, ptr [[A]], align 41051// CHECK3-NEXT: ret i32 [[TMP6]]1052//1053//1054// CHECK3-LABEL: define {{[^@]+}}@_ZN2S12r1Ei1055// CHECK3-SAME: (ptr noundef nonnull align 4 dereferenceable(8) [[THIS:%.*]], i32 noundef [[N:%.*]]) #[[ATTR0]] comdat align 2 {1056// CHECK3-NEXT: entry:1057// CHECK3-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 41058// CHECK3-NEXT: [[N_ADDR:%.*]] = alloca i32, align 41059// CHECK3-NEXT: [[B:%.*]] = alloca i32, align 41060// CHECK3-NEXT: [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 41061// CHECK3-NEXT: [[B_CASTED:%.*]] = alloca i32, align 41062// CHECK3-NEXT: [[DOTCAPTURE_EXPR__CASTED:%.*]] = alloca i32, align 41063// CHECK3-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [3 x ptr], align 41064// CHECK3-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [3 x ptr], align 41065// CHECK3-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [3 x ptr], align 41066// CHECK3-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 81067// CHECK3-NEXT: [[DOTOFFLOAD_BASEPTRS3:%.*]] = alloca [1 x ptr], align 41068// CHECK3-NEXT: [[DOTOFFLOAD_PTRS4:%.*]] = alloca [1 x ptr], align 41069// CHECK3-NEXT: [[DOTOFFLOAD_MAPPERS5:%.*]] = alloca [1 x ptr], align 41070// CHECK3-NEXT: [[KERNEL_ARGS6:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 81071// CHECK3-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 41072// CHECK3-NEXT: store i32 [[N]], ptr [[N_ADDR]], align 41073// CHECK3-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 41074// CHECK3-NEXT: store i32 1, ptr [[B]], align 41075// CHECK3-NEXT: [[TMP0:%.*]] = load i32, ptr [[N_ADDR]], align 41076// CHECK3-NEXT: [[TMP1:%.*]] = load i32, ptr [[B]], align 41077// CHECK3-NEXT: [[SUB:%.*]] = sub nsw i32 [[TMP0]], [[TMP1]]1078// CHECK3-NEXT: store i32 [[SUB]], ptr [[DOTCAPTURE_EXPR_]], align 41079// CHECK3-NEXT: [[TMP2:%.*]] = load i32, ptr [[B]], align 41080// CHECK3-NEXT: store i32 [[TMP2]], ptr [[B_CASTED]], align 41081// CHECK3-NEXT: [[TMP3:%.*]] = load i32, ptr [[B_CASTED]], align 41082// CHECK3-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 41083// CHECK3-NEXT: store i32 [[TMP4]], ptr [[DOTCAPTURE_EXPR__CASTED]], align 41084// CHECK3-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR__CASTED]], align 41085// CHECK3-NEXT: [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_S1:%.*]], ptr [[THIS1]], i32 0, i32 01086// CHECK3-NEXT: [[TMP6:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 01087// CHECK3-NEXT: store ptr [[THIS1]], ptr [[TMP6]], align 41088// CHECK3-NEXT: [[TMP7:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 01089// CHECK3-NEXT: store ptr [[A]], ptr [[TMP7]], align 41090// CHECK3-NEXT: [[TMP8:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 01091// CHECK3-NEXT: store ptr null, ptr [[TMP8]], align 41092// CHECK3-NEXT: [[TMP9:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 11093// CHECK3-NEXT: store i32 [[TMP3]], ptr [[TMP9]], align 41094// CHECK3-NEXT: [[TMP10:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 11095// CHECK3-NEXT: store i32 [[TMP3]], ptr [[TMP10]], align 41096// CHECK3-NEXT: [[TMP11:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 11097// CHECK3-NEXT: store ptr null, ptr [[TMP11]], align 41098// CHECK3-NEXT: [[TMP12:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 21099// CHECK3-NEXT: store i32 [[TMP5]], ptr [[TMP12]], align 41100// CHECK3-NEXT: [[TMP13:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 21101// CHECK3-NEXT: store i32 [[TMP5]], ptr [[TMP13]], align 41102// CHECK3-NEXT: [[TMP14:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 21103// CHECK3-NEXT: store ptr null, ptr [[TMP14]], align 41104// CHECK3-NEXT: [[TMP15:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 01105// CHECK3-NEXT: [[TMP16:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 01106// CHECK3-NEXT: [[TMP17:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 41107// CHECK3-NEXT: [[TMP18:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 01108// CHECK3-NEXT: store i32 3, ptr [[TMP18]], align 41109// CHECK3-NEXT: [[TMP19:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11110// CHECK3-NEXT: store i32 3, ptr [[TMP19]], align 41111// CHECK3-NEXT: [[TMP20:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 21112// CHECK3-NEXT: store ptr [[TMP15]], ptr [[TMP20]], align 41113// CHECK3-NEXT: [[TMP21:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 31114// CHECK3-NEXT: store ptr [[TMP16]], ptr [[TMP21]], align 41115// CHECK3-NEXT: [[TMP22:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 41116// CHECK3-NEXT: store ptr @.offload_sizes, ptr [[TMP22]], align 41117// CHECK3-NEXT: [[TMP23:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 51118// CHECK3-NEXT: store ptr @.offload_maptypes, ptr [[TMP23]], align 41119// CHECK3-NEXT: [[TMP24:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 61120// CHECK3-NEXT: store ptr null, ptr [[TMP24]], align 41121// CHECK3-NEXT: [[TMP25:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 71122// CHECK3-NEXT: store ptr null, ptr [[TMP25]], align 41123// CHECK3-NEXT: [[TMP26:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 81124// CHECK3-NEXT: store i64 0, ptr [[TMP26]], align 81125// CHECK3-NEXT: [[TMP27:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 91126// CHECK3-NEXT: store i64 4, ptr [[TMP27]], align 81127// CHECK3-NEXT: [[TMP28:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 101128// CHECK3-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP28]], align 41129// CHECK3-NEXT: [[TMP29:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 111130// CHECK3-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP29]], align 41131// CHECK3-NEXT: [[TMP30:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 121132// CHECK3-NEXT: store i32 [[TMP17]], ptr [[TMP30]], align 41133// CHECK3-NEXT: [[TMP31:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1:[0-9]+]], i64 -1, i32 0, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2S12r1Ei_l88.region_id, ptr [[KERNEL_ARGS]])1134// CHECK3-NEXT: [[TMP32:%.*]] = icmp ne i32 [[TMP31]], 01135// CHECK3-NEXT: br i1 [[TMP32]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]1136// CHECK3: omp_offload.failed:1137// CHECK3-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2S12r1Ei_l88(ptr [[THIS1]], i32 [[TMP3]], i32 [[TMP5]]) #[[ATTR2:[0-9]+]]1138// CHECK3-NEXT: br label [[OMP_OFFLOAD_CONT]]1139// CHECK3: omp_offload.cont:1140// CHECK3-NEXT: [[A2:%.*]] = getelementptr inbounds nuw [[STRUCT_S1]], ptr [[THIS1]], i32 0, i32 01141// CHECK3-NEXT: [[TMP33:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS3]], i32 0, i32 01142// CHECK3-NEXT: store ptr [[THIS1]], ptr [[TMP33]], align 41143// CHECK3-NEXT: [[TMP34:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS4]], i32 0, i32 01144// CHECK3-NEXT: store ptr [[A2]], ptr [[TMP34]], align 41145// CHECK3-NEXT: [[TMP35:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS5]], i32 0, i32 01146// CHECK3-NEXT: store ptr null, ptr [[TMP35]], align 41147// CHECK3-NEXT: [[TMP36:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS3]], i32 0, i32 01148// CHECK3-NEXT: [[TMP37:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS4]], i32 0, i32 01149// CHECK3-NEXT: [[TMP38:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 01150// CHECK3-NEXT: store i32 3, ptr [[TMP38]], align 41151// CHECK3-NEXT: [[TMP39:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 11152// CHECK3-NEXT: store i32 1, ptr [[TMP39]], align 41153// CHECK3-NEXT: [[TMP40:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 21154// CHECK3-NEXT: store ptr [[TMP36]], ptr [[TMP40]], align 41155// CHECK3-NEXT: [[TMP41:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 31156// CHECK3-NEXT: store ptr [[TMP37]], ptr [[TMP41]], align 41157// CHECK3-NEXT: [[TMP42:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 41158// CHECK3-NEXT: store ptr @.offload_sizes.1, ptr [[TMP42]], align 41159// CHECK3-NEXT: [[TMP43:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 51160// CHECK3-NEXT: store ptr @.offload_maptypes.2, ptr [[TMP43]], align 41161// CHECK3-NEXT: [[TMP44:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 61162// CHECK3-NEXT: store ptr null, ptr [[TMP44]], align 41163// CHECK3-NEXT: [[TMP45:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 71164// CHECK3-NEXT: store ptr null, ptr [[TMP45]], align 41165// CHECK3-NEXT: [[TMP46:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 81166// CHECK3-NEXT: store i64 0, ptr [[TMP46]], align 81167// CHECK3-NEXT: [[TMP47:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 91168// CHECK3-NEXT: store i64 0, ptr [[TMP47]], align 81169// CHECK3-NEXT: [[TMP48:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 101170// CHECK3-NEXT: store [3 x i32] [i32 -1, i32 0, i32 0], ptr [[TMP48]], align 41171// CHECK3-NEXT: [[TMP49:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 111172// CHECK3-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP49]], align 41173// CHECK3-NEXT: [[TMP50:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS6]], i32 0, i32 121174// CHECK3-NEXT: store i32 1024, ptr [[TMP50]], align 41175// CHECK3-NEXT: [[TMP51:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 -1, i32 -1, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2S12r1Ei_l93.region_id, ptr [[KERNEL_ARGS6]])1176// CHECK3-NEXT: [[TMP52:%.*]] = icmp ne i32 [[TMP51]], 01177// CHECK3-NEXT: br i1 [[TMP52]], label [[OMP_OFFLOAD_FAILED7:%.*]], label [[OMP_OFFLOAD_CONT8:%.*]]1178// CHECK3: omp_offload.failed7:1179// CHECK3-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2S12r1Ei_l93(ptr [[THIS1]]) #[[ATTR2]]1180// CHECK3-NEXT: br label [[OMP_OFFLOAD_CONT8]]1181// CHECK3: omp_offload.cont8:1182// CHECK3-NEXT: [[A9:%.*]] = getelementptr inbounds nuw [[STRUCT_S1]], ptr [[THIS1]], i32 0, i32 01183// CHECK3-NEXT: [[TMP53:%.*]] = load double, ptr [[A9]], align 41184// CHECK3-NEXT: [[CONV:%.*]] = fptosi double [[TMP53]] to i321185// CHECK3-NEXT: ret i32 [[CONV]]1186//1187//1188// CHECK3-LABEL: define {{[^@]+}}@_ZL7fstatici1189// CHECK3-SAME: (i32 noundef [[N:%.*]]) #[[ATTR0]] {1190// CHECK3-NEXT: entry:1191// CHECK3-NEXT: [[N_ADDR:%.*]] = alloca i32, align 41192// CHECK3-NEXT: [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 41193// CHECK3-NEXT: [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 41194// CHECK3-NEXT: [[N_CASTED:%.*]] = alloca i32, align 41195// CHECK3-NEXT: [[DOTCAPTURE_EXPR__CASTED:%.*]] = alloca i32, align 41196// CHECK3-NEXT: [[DOTCAPTURE_EXPR__CASTED2:%.*]] = alloca i32, align 41197// CHECK3-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [3 x ptr], align 41198// CHECK3-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [3 x ptr], align 41199// CHECK3-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [3 x ptr], align 41200// CHECK3-NEXT: [[TMP:%.*]] = alloca i32, align 41201// CHECK3-NEXT: [[DOTCAPTURE_EXPR_3:%.*]] = alloca i32, align 41202// CHECK3-NEXT: [[DOTCAPTURE_EXPR_4:%.*]] = alloca i32, align 41203// CHECK3-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 81204// CHECK3-NEXT: [[DOTCAPTURE_EXPR_6:%.*]] = alloca i32, align 41205// CHECK3-NEXT: [[DOTCAPTURE_EXPR__CASTED8:%.*]] = alloca i32, align 41206// CHECK3-NEXT: [[DOTOFFLOAD_BASEPTRS9:%.*]] = alloca [1 x ptr], align 41207// CHECK3-NEXT: [[DOTOFFLOAD_PTRS10:%.*]] = alloca [1 x ptr], align 41208// CHECK3-NEXT: [[DOTOFFLOAD_MAPPERS11:%.*]] = alloca [1 x ptr], align 41209// CHECK3-NEXT: [[AGG_CAPTURED:%.*]] = alloca [[STRUCT_ANON:%.*]], align 41210// CHECK3-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])1211// CHECK3-NEXT: store i32 [[N]], ptr [[N_ADDR]], align 41212// CHECK3-NEXT: [[TMP1:%.*]] = load i32, ptr [[N_ADDR]], align 41213// CHECK3-NEXT: store i32 [[TMP1]], ptr [[DOTCAPTURE_EXPR_]], align 41214// CHECK3-NEXT: [[TMP2:%.*]] = load i32, ptr [[N_ADDR]], align 41215// CHECK3-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP2]], 321216// CHECK3-NEXT: store i32 [[MUL]], ptr [[DOTCAPTURE_EXPR_1]], align 41217// CHECK3-NEXT: [[TMP3:%.*]] = load i32, ptr [[N_ADDR]], align 41218// CHECK3-NEXT: store i32 [[TMP3]], ptr [[N_CASTED]], align 41219// CHECK3-NEXT: [[TMP4:%.*]] = load i32, ptr [[N_CASTED]], align 41220// CHECK3-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 41221// CHECK3-NEXT: store i32 [[TMP5]], ptr [[DOTCAPTURE_EXPR__CASTED]], align 41222// CHECK3-NEXT: [[TMP6:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR__CASTED]], align 41223// CHECK3-NEXT: [[TMP7:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41224// CHECK3-NEXT: store i32 [[TMP7]], ptr [[DOTCAPTURE_EXPR__CASTED2]], align 41225// CHECK3-NEXT: [[TMP8:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR__CASTED2]], align 41226// CHECK3-NEXT: [[TMP9:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 01227// CHECK3-NEXT: store i32 [[TMP4]], ptr [[TMP9]], align 41228// CHECK3-NEXT: [[TMP10:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 01229// CHECK3-NEXT: store i32 [[TMP4]], ptr [[TMP10]], align 41230// CHECK3-NEXT: [[TMP11:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 01231// CHECK3-NEXT: store ptr null, ptr [[TMP11]], align 41232// CHECK3-NEXT: [[TMP12:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 11233// CHECK3-NEXT: store i32 [[TMP6]], ptr [[TMP12]], align 41234// CHECK3-NEXT: [[TMP13:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 11235// CHECK3-NEXT: store i32 [[TMP6]], ptr [[TMP13]], align 41236// CHECK3-NEXT: [[TMP14:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 11237// CHECK3-NEXT: store ptr null, ptr [[TMP14]], align 41238// CHECK3-NEXT: [[TMP15:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 21239// CHECK3-NEXT: store i32 [[TMP8]], ptr [[TMP15]], align 41240// CHECK3-NEXT: [[TMP16:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 21241// CHECK3-NEXT: store i32 [[TMP8]], ptr [[TMP16]], align 41242// CHECK3-NEXT: [[TMP17:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 21243// CHECK3-NEXT: store ptr null, ptr [[TMP17]], align 41244// CHECK3-NEXT: [[TMP18:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 01245// CHECK3-NEXT: [[TMP19:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 01246// CHECK3-NEXT: [[TMP20:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 41247// CHECK3-NEXT: [[TMP21:%.*]] = load i32, ptr [[N_ADDR]], align 41248// CHECK3-NEXT: store i32 [[TMP21]], ptr [[DOTCAPTURE_EXPR_3]], align 41249// CHECK3-NEXT: [[TMP22:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_3]], align 41250// CHECK3-NEXT: [[SUB:%.*]] = sub nsw i32 [[TMP22]], 01251// CHECK3-NEXT: [[DIV:%.*]] = sdiv i32 [[SUB]], 11252// CHECK3-NEXT: [[SUB5:%.*]] = sub nsw i32 [[DIV]], 11253// CHECK3-NEXT: store i32 [[SUB5]], ptr [[DOTCAPTURE_EXPR_4]], align 41254// CHECK3-NEXT: [[TMP23:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_4]], align 41255// CHECK3-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP23]], 11256// CHECK3-NEXT: [[TMP24:%.*]] = zext i32 [[ADD]] to i641257// CHECK3-NEXT: [[TMP25:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41258// CHECK3-NEXT: [[TMP26:%.*]] = insertvalue [3 x i32] zeroinitializer, i32 [[TMP20]], 01259// CHECK3-NEXT: [[TMP27:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 01260// CHECK3-NEXT: store i32 3, ptr [[TMP27]], align 41261// CHECK3-NEXT: [[TMP28:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11262// CHECK3-NEXT: store i32 3, ptr [[TMP28]], align 41263// CHECK3-NEXT: [[TMP29:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 21264// CHECK3-NEXT: store ptr [[TMP18]], ptr [[TMP29]], align 41265// CHECK3-NEXT: [[TMP30:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 31266// CHECK3-NEXT: store ptr [[TMP19]], ptr [[TMP30]], align 41267// CHECK3-NEXT: [[TMP31:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 41268// CHECK3-NEXT: store ptr @.offload_sizes.3, ptr [[TMP31]], align 41269// CHECK3-NEXT: [[TMP32:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 51270// CHECK3-NEXT: store ptr @.offload_maptypes.4, ptr [[TMP32]], align 41271// CHECK3-NEXT: [[TMP33:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 61272// CHECK3-NEXT: store ptr null, ptr [[TMP33]], align 41273// CHECK3-NEXT: [[TMP34:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 71274// CHECK3-NEXT: store ptr null, ptr [[TMP34]], align 41275// CHECK3-NEXT: [[TMP35:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 81276// CHECK3-NEXT: store i64 [[TMP24]], ptr [[TMP35]], align 81277// CHECK3-NEXT: [[TMP36:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 91278// CHECK3-NEXT: store i64 8, ptr [[TMP36]], align 81279// CHECK3-NEXT: [[TMP37:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 101280// CHECK3-NEXT: store [3 x i32] [[TMP26]], ptr [[TMP37]], align 41281// CHECK3-NEXT: [[TMP38:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 111282// CHECK3-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP38]], align 41283// CHECK3-NEXT: [[TMP39:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 121284// CHECK3-NEXT: store i32 [[TMP25]], ptr [[TMP39]], align 41285// CHECK3-NEXT: [[TMP40:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 -1, i32 [[TMP20]], i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l71.region_id, ptr [[KERNEL_ARGS]])1286// CHECK3-NEXT: [[TMP41:%.*]] = icmp ne i32 [[TMP40]], 01287// CHECK3-NEXT: br i1 [[TMP41]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]1288// CHECK3: omp_offload.failed:1289// CHECK3-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l71(i32 [[TMP4]], i32 [[TMP6]], i32 [[TMP8]]) #[[ATTR2]]1290// CHECK3-NEXT: br label [[OMP_OFFLOAD_CONT]]1291// CHECK3: omp_offload.cont:1292// CHECK3-NEXT: [[TMP42:%.*]] = load i32, ptr [[N_ADDR]], align 41293// CHECK3-NEXT: [[ADD7:%.*]] = add nsw i32 32, [[TMP42]]1294// CHECK3-NEXT: store i32 [[ADD7]], ptr [[DOTCAPTURE_EXPR_6]], align 41295// CHECK3-NEXT: [[TMP43:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_6]], align 41296// CHECK3-NEXT: store i32 [[TMP43]], ptr [[DOTCAPTURE_EXPR__CASTED8]], align 41297// CHECK3-NEXT: [[TMP44:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR__CASTED8]], align 41298// CHECK3-NEXT: [[TMP45:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS9]], i32 0, i32 01299// CHECK3-NEXT: store i32 [[TMP44]], ptr [[TMP45]], align 41300// CHECK3-NEXT: [[TMP46:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS10]], i32 0, i32 01301// CHECK3-NEXT: store i32 [[TMP44]], ptr [[TMP46]], align 41302// CHECK3-NEXT: [[TMP47:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS11]], i32 0, i32 01303// CHECK3-NEXT: store ptr null, ptr [[TMP47]], align 41304// CHECK3-NEXT: [[TMP48:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS9]], i32 0, i32 01305// CHECK3-NEXT: [[TMP49:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS10]], i32 0, i32 01306// CHECK3-NEXT: [[TMP50:%.*]] = getelementptr inbounds nuw [[STRUCT_ANON]], ptr [[AGG_CAPTURED]], i32 0, i32 01307// CHECK3-NEXT: [[TMP51:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_6]], align 41308// CHECK3-NEXT: store i32 [[TMP51]], ptr [[TMP50]], align 41309// CHECK3-NEXT: [[TMP52:%.*]] = call ptr @__kmpc_omp_target_task_alloc(ptr @[[GLOB1]], i32 [[TMP0]], i32 1, i32 36, i32 4, ptr @.omp_task_entry., i64 -1)1310// CHECK3-NEXT: [[TMP53:%.*]] = getelementptr inbounds nuw [[STRUCT_KMP_TASK_T_WITH_PRIVATES:%.*]], ptr [[TMP52]], i32 0, i32 01311// CHECK3-NEXT: [[TMP54:%.*]] = getelementptr inbounds nuw [[STRUCT_KMP_TASK_T:%.*]], ptr [[TMP53]], i32 0, i32 01312// CHECK3-NEXT: [[TMP55:%.*]] = load ptr, ptr [[TMP54]], align 41313// CHECK3-NEXT: call void @llvm.memcpy.p0.p0.i32(ptr align 4 [[TMP55]], ptr align 4 [[AGG_CAPTURED]], i32 4, i1 false)1314// CHECK3-NEXT: [[TMP56:%.*]] = getelementptr inbounds nuw [[STRUCT_KMP_TASK_T_WITH_PRIVATES]], ptr [[TMP52]], i32 0, i32 11315// CHECK3-NEXT: [[TMP57:%.*]] = getelementptr inbounds nuw [[STRUCT__KMP_PRIVATES_T:%.*]], ptr [[TMP56]], i32 0, i32 01316// CHECK3-NEXT: call void @llvm.memcpy.p0.p0.i32(ptr align 4 [[TMP57]], ptr align 4 @.offload_sizes.5, i32 8, i1 false)1317// CHECK3-NEXT: [[TMP58:%.*]] = getelementptr inbounds nuw [[STRUCT__KMP_PRIVATES_T]], ptr [[TMP56]], i32 0, i32 11318// CHECK3-NEXT: call void @llvm.memcpy.p0.p0.i32(ptr align 4 [[TMP58]], ptr align 4 [[TMP48]], i32 4, i1 false)1319// CHECK3-NEXT: [[TMP59:%.*]] = getelementptr inbounds nuw [[STRUCT__KMP_PRIVATES_T]], ptr [[TMP56]], i32 0, i32 21320// CHECK3-NEXT: call void @llvm.memcpy.p0.p0.i32(ptr align 4 [[TMP59]], ptr align 4 [[TMP49]], i32 4, i1 false)1321// CHECK3-NEXT: [[TMP60:%.*]] = call i32 @__kmpc_omp_task(ptr @[[GLOB1]], i32 [[TMP0]], ptr [[TMP52]])1322// CHECK3-NEXT: [[TMP61:%.*]] = load i32, ptr [[N_ADDR]], align 41323// CHECK3-NEXT: [[ADD12:%.*]] = add nsw i32 [[TMP61]], 11324// CHECK3-NEXT: ret i32 [[ADD12]]1325//1326//1327// CHECK3-LABEL: define {{[^@]+}}@_Z9ftemplateIiET_i1328// CHECK3-SAME: (i32 noundef [[N:%.*]]) #[[ATTR0]] comdat {1329// CHECK3-NEXT: entry:1330// CHECK3-NEXT: [[N_ADDR:%.*]] = alloca i32, align 41331// CHECK3-NEXT: [[A:%.*]] = alloca i32, align 41332// CHECK3-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 81333// CHECK3-NEXT: [[B:%.*]] = alloca i16, align 21334// CHECK3-NEXT: [[DOTCAPTURE_EXPR_:%.*]] = alloca i16, align 21335// CHECK3-NEXT: [[A_CASTED:%.*]] = alloca i32, align 41336// CHECK3-NEXT: [[B_CASTED:%.*]] = alloca i32, align 41337// CHECK3-NEXT: [[DOTCAPTURE_EXPR__CASTED:%.*]] = alloca i32, align 41338// CHECK3-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [3 x ptr], align 41339// CHECK3-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [3 x ptr], align 41340// CHECK3-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [3 x ptr], align 41341// CHECK3-NEXT: [[KERNEL_ARGS1:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 81342// CHECK3-NEXT: store i32 [[N]], ptr [[N_ADDR]], align 41343// CHECK3-NEXT: store i32 0, ptr [[A]], align 41344// CHECK3-NEXT: [[TMP0:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 01345// CHECK3-NEXT: store i32 3, ptr [[TMP0]], align 41346// CHECK3-NEXT: [[TMP1:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11347// CHECK3-NEXT: store i32 0, ptr [[TMP1]], align 41348// CHECK3-NEXT: [[TMP2:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 21349// CHECK3-NEXT: store ptr null, ptr [[TMP2]], align 41350// CHECK3-NEXT: [[TMP3:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 31351// CHECK3-NEXT: store ptr null, ptr [[TMP3]], align 41352// CHECK3-NEXT: [[TMP4:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 41353// CHECK3-NEXT: store ptr null, ptr [[TMP4]], align 41354// CHECK3-NEXT: [[TMP5:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 51355// CHECK3-NEXT: store ptr null, ptr [[TMP5]], align 41356// CHECK3-NEXT: [[TMP6:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 61357// CHECK3-NEXT: store ptr null, ptr [[TMP6]], align 41358// CHECK3-NEXT: [[TMP7:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 71359// CHECK3-NEXT: store ptr null, ptr [[TMP7]], align 41360// CHECK3-NEXT: [[TMP8:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 81361// CHECK3-NEXT: store i64 0, ptr [[TMP8]], align 81362// CHECK3-NEXT: [[TMP9:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 91363// CHECK3-NEXT: store i64 8, ptr [[TMP9]], align 81364// CHECK3-NEXT: [[TMP10:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 101365// CHECK3-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP10]], align 41366// CHECK3-NEXT: [[TMP11:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 111367// CHECK3-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP11]], align 41368// CHECK3-NEXT: [[TMP12:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 121369// CHECK3-NEXT: store i32 20, ptr [[TMP12]], align 41370// CHECK3-NEXT: [[TMP13:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 -1, i32 0, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l55.region_id, ptr [[KERNEL_ARGS]])1371// CHECK3-NEXT: [[TMP14:%.*]] = icmp ne i32 [[TMP13]], 01372// CHECK3-NEXT: br i1 [[TMP14]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]1373// CHECK3: omp_offload.failed:1374// CHECK3-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l55() #[[ATTR2]]1375// CHECK3-NEXT: br label [[OMP_OFFLOAD_CONT]]1376// CHECK3: omp_offload.cont:1377// CHECK3-NEXT: store i16 1, ptr [[B]], align 21378// CHECK3-NEXT: [[TMP15:%.*]] = load i16, ptr [[B]], align 21379// CHECK3-NEXT: store i16 [[TMP15]], ptr [[DOTCAPTURE_EXPR_]], align 21380// CHECK3-NEXT: [[TMP16:%.*]] = load i32, ptr [[A]], align 41381// CHECK3-NEXT: store i32 [[TMP16]], ptr [[A_CASTED]], align 41382// CHECK3-NEXT: [[TMP17:%.*]] = load i32, ptr [[A_CASTED]], align 41383// CHECK3-NEXT: [[TMP18:%.*]] = load i16, ptr [[B]], align 21384// CHECK3-NEXT: store i16 [[TMP18]], ptr [[B_CASTED]], align 21385// CHECK3-NEXT: [[TMP19:%.*]] = load i32, ptr [[B_CASTED]], align 41386// CHECK3-NEXT: [[TMP20:%.*]] = load i16, ptr [[DOTCAPTURE_EXPR_]], align 21387// CHECK3-NEXT: store i16 [[TMP20]], ptr [[DOTCAPTURE_EXPR__CASTED]], align 21388// CHECK3-NEXT: [[TMP21:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR__CASTED]], align 41389// CHECK3-NEXT: [[TMP22:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 01390// CHECK3-NEXT: store i32 [[TMP17]], ptr [[TMP22]], align 41391// CHECK3-NEXT: [[TMP23:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 01392// CHECK3-NEXT: store i32 [[TMP17]], ptr [[TMP23]], align 41393// CHECK3-NEXT: [[TMP24:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 01394// CHECK3-NEXT: store ptr null, ptr [[TMP24]], align 41395// CHECK3-NEXT: [[TMP25:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 11396// CHECK3-NEXT: store i32 [[TMP19]], ptr [[TMP25]], align 41397// CHECK3-NEXT: [[TMP26:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 11398// CHECK3-NEXT: store i32 [[TMP19]], ptr [[TMP26]], align 41399// CHECK3-NEXT: [[TMP27:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 11400// CHECK3-NEXT: store ptr null, ptr [[TMP27]], align 41401// CHECK3-NEXT: [[TMP28:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 21402// CHECK3-NEXT: store i32 [[TMP21]], ptr [[TMP28]], align 41403// CHECK3-NEXT: [[TMP29:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 21404// CHECK3-NEXT: store i32 [[TMP21]], ptr [[TMP29]], align 41405// CHECK3-NEXT: [[TMP30:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 21406// CHECK3-NEXT: store ptr null, ptr [[TMP30]], align 41407// CHECK3-NEXT: [[TMP31:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 01408// CHECK3-NEXT: [[TMP32:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 01409// CHECK3-NEXT: [[TMP33:%.*]] = load i16, ptr [[DOTCAPTURE_EXPR_]], align 21410// CHECK3-NEXT: [[TMP34:%.*]] = sext i16 [[TMP33]] to i321411// CHECK3-NEXT: [[TMP35:%.*]] = insertvalue [3 x i32] zeroinitializer, i32 [[TMP34]], 01412// CHECK3-NEXT: [[TMP36:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS1]], i32 0, i32 01413// CHECK3-NEXT: store i32 3, ptr [[TMP36]], align 41414// CHECK3-NEXT: [[TMP37:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS1]], i32 0, i32 11415// CHECK3-NEXT: store i32 3, ptr [[TMP37]], align 41416// CHECK3-NEXT: [[TMP38:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS1]], i32 0, i32 21417// CHECK3-NEXT: store ptr [[TMP31]], ptr [[TMP38]], align 41418// CHECK3-NEXT: [[TMP39:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS1]], i32 0, i32 31419// CHECK3-NEXT: store ptr [[TMP32]], ptr [[TMP39]], align 41420// CHECK3-NEXT: [[TMP40:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS1]], i32 0, i32 41421// CHECK3-NEXT: store ptr @.offload_sizes.7, ptr [[TMP40]], align 41422// CHECK3-NEXT: [[TMP41:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS1]], i32 0, i32 51423// CHECK3-NEXT: store ptr @.offload_maptypes.8, ptr [[TMP41]], align 41424// CHECK3-NEXT: [[TMP42:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS1]], i32 0, i32 61425// CHECK3-NEXT: store ptr null, ptr [[TMP42]], align 41426// CHECK3-NEXT: [[TMP43:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS1]], i32 0, i32 71427// CHECK3-NEXT: store ptr null, ptr [[TMP43]], align 41428// CHECK3-NEXT: [[TMP44:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS1]], i32 0, i32 81429// CHECK3-NEXT: store i64 0, ptr [[TMP44]], align 81430// CHECK3-NEXT: [[TMP45:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS1]], i32 0, i32 91431// CHECK3-NEXT: store i64 8, ptr [[TMP45]], align 81432// CHECK3-NEXT: [[TMP46:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS1]], i32 0, i32 101433// CHECK3-NEXT: store [3 x i32] [[TMP35]], ptr [[TMP46]], align 41434// CHECK3-NEXT: [[TMP47:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS1]], i32 0, i32 111435// CHECK3-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP47]], align 41436// CHECK3-NEXT: [[TMP48:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS1]], i32 0, i32 121437// CHECK3-NEXT: store i32 1024, ptr [[TMP48]], align 41438// CHECK3-NEXT: [[TMP49:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 -1, i32 [[TMP34]], i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l60.region_id, ptr [[KERNEL_ARGS1]])1439// CHECK3-NEXT: [[TMP50:%.*]] = icmp ne i32 [[TMP49]], 01440// CHECK3-NEXT: br i1 [[TMP50]], label [[OMP_OFFLOAD_FAILED2:%.*]], label [[OMP_OFFLOAD_CONT3:%.*]]1441// CHECK3: omp_offload.failed2:1442// CHECK3-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l60(i32 [[TMP17]], i32 [[TMP19]], i32 [[TMP21]]) #[[ATTR2]]1443// CHECK3-NEXT: br label [[OMP_OFFLOAD_CONT3]]1444// CHECK3: omp_offload.cont3:1445// CHECK3-NEXT: [[TMP51:%.*]] = load i32, ptr [[A]], align 41446// CHECK3-NEXT: ret i32 [[TMP51]]1447//1448//1449// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2S12r1Ei_l881450// CHECK3-SAME: (ptr noundef [[THIS:%.*]], i32 noundef [[B:%.*]], i32 noundef [[DOTCAPTURE_EXPR_:%.*]]) #[[ATTR1:[0-9]+]] {1451// CHECK3-NEXT: entry:1452// CHECK3-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 41453// CHECK3-NEXT: [[B_ADDR:%.*]] = alloca i32, align 41454// CHECK3-NEXT: [[DOTCAPTURE_EXPR__ADDR:%.*]] = alloca i32, align 41455// CHECK3-NEXT: [[B_CASTED:%.*]] = alloca i32, align 41456// CHECK3-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 41457// CHECK3-NEXT: store i32 [[B]], ptr [[B_ADDR]], align 41458// CHECK3-NEXT: store i32 [[DOTCAPTURE_EXPR_]], ptr [[DOTCAPTURE_EXPR__ADDR]], align 41459// CHECK3-NEXT: [[TMP0:%.*]] = load ptr, ptr [[THIS_ADDR]], align 41460// CHECK3-NEXT: [[TMP1:%.*]] = load i32, ptr [[B_ADDR]], align 41461// CHECK3-NEXT: store i32 [[TMP1]], ptr [[B_CASTED]], align 41462// CHECK3-NEXT: [[TMP2:%.*]] = load i32, ptr [[B_CASTED]], align 41463// CHECK3-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB1]], i32 2, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2S12r1Ei_l88.omp_outlined, ptr [[TMP0]], i32 [[TMP2]])1464// CHECK3-NEXT: ret void1465//1466//1467// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2S12r1Ei_l88.omp_outlined1468// CHECK3-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef [[THIS:%.*]], i32 noundef [[B:%.*]]) #[[ATTR1]] {1469// CHECK3-NEXT: entry:1470// CHECK3-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 41471// CHECK3-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 41472// CHECK3-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 41473// CHECK3-NEXT: [[B_ADDR:%.*]] = alloca i32, align 41474// CHECK3-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 41475// CHECK3-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 41476// CHECK3-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 41477// CHECK3-NEXT: store i32 [[B]], ptr [[B_ADDR]], align 41478// CHECK3-NEXT: [[TMP0:%.*]] = load ptr, ptr [[THIS_ADDR]], align 41479// CHECK3-NEXT: [[TMP1:%.*]] = load i32, ptr [[B_ADDR]], align 41480// CHECK3-NEXT: [[CONV:%.*]] = sitofp i32 [[TMP1]] to double1481// CHECK3-NEXT: [[ADD:%.*]] = fadd double [[CONV]], 1.500000e+001482// CHECK3-NEXT: [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_S1:%.*]], ptr [[TMP0]], i32 0, i32 01483// CHECK3-NEXT: store double [[ADD]], ptr [[A]], align 41484// CHECK3-NEXT: ret void1485//1486//1487// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2S12r1Ei_l931488// CHECK3-SAME: (ptr noundef [[THIS:%.*]]) #[[ATTR1]] {1489// CHECK3-NEXT: entry:1490// CHECK3-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 41491// CHECK3-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 41492// CHECK3-NEXT: [[TMP0:%.*]] = load ptr, ptr [[THIS_ADDR]], align 41493// CHECK3-NEXT: [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_S1:%.*]], ptr [[TMP0]], i32 0, i32 01494// CHECK3-NEXT: store double 2.500000e+00, ptr [[A]], align 41495// CHECK3-NEXT: ret void1496//1497//1498// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l711499// CHECK3-SAME: (i32 noundef [[N:%.*]], i32 noundef [[DOTCAPTURE_EXPR_:%.*]], i32 noundef [[DOTCAPTURE_EXPR_1:%.*]]) #[[ATTR1]] {1500// CHECK3-NEXT: entry:1501// CHECK3-NEXT: [[N_ADDR:%.*]] = alloca i32, align 41502// CHECK3-NEXT: [[DOTCAPTURE_EXPR__ADDR:%.*]] = alloca i32, align 41503// CHECK3-NEXT: [[DOTCAPTURE_EXPR__ADDR2:%.*]] = alloca i32, align 41504// CHECK3-NEXT: [[N_CASTED:%.*]] = alloca i32, align 41505// CHECK3-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])1506// CHECK3-NEXT: store i32 [[N]], ptr [[N_ADDR]], align 41507// CHECK3-NEXT: store i32 [[DOTCAPTURE_EXPR_]], ptr [[DOTCAPTURE_EXPR__ADDR]], align 41508// CHECK3-NEXT: store i32 [[DOTCAPTURE_EXPR_1]], ptr [[DOTCAPTURE_EXPR__ADDR2]], align 41509// CHECK3-NEXT: [[TMP1:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR__ADDR]], align 41510// CHECK3-NEXT: call void @__kmpc_push_num_teams(ptr @[[GLOB1]], i32 [[TMP0]], i32 [[TMP1]], i32 0)1511// CHECK3-NEXT: [[TMP2:%.*]] = load i32, ptr [[N_ADDR]], align 41512// CHECK3-NEXT: store i32 [[TMP2]], ptr [[N_CASTED]], align 41513// CHECK3-NEXT: [[TMP3:%.*]] = load i32, ptr [[N_CASTED]], align 41514// CHECK3-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB1]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l71.omp_outlined, i32 [[TMP3]])1515// CHECK3-NEXT: ret void1516//1517//1518// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l71.omp_outlined1519// CHECK3-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i32 noundef [[N:%.*]]) #[[ATTR1]] {1520// CHECK3-NEXT: entry:1521// CHECK3-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 41522// CHECK3-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 41523// CHECK3-NEXT: [[N_ADDR:%.*]] = alloca i32, align 41524// CHECK3-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 41525// CHECK3-NEXT: [[TMP:%.*]] = alloca i32, align 41526// CHECK3-NEXT: [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 41527// CHECK3-NEXT: [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 41528// CHECK3-NEXT: [[I:%.*]] = alloca i32, align 41529// CHECK3-NEXT: [[DOTOMP_COMB_LB:%.*]] = alloca i32, align 41530// CHECK3-NEXT: [[DOTOMP_COMB_UB:%.*]] = alloca i32, align 41531// CHECK3-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 41532// CHECK3-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 41533// CHECK3-NEXT: [[I3:%.*]] = alloca i32, align 41534// CHECK3-NEXT: [[N_CASTED:%.*]] = alloca i32, align 41535// CHECK3-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 41536// CHECK3-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 41537// CHECK3-NEXT: store i32 [[N]], ptr [[N_ADDR]], align 41538// CHECK3-NEXT: [[TMP0:%.*]] = load i32, ptr [[N_ADDR]], align 41539// CHECK3-NEXT: store i32 [[TMP0]], ptr [[DOTCAPTURE_EXPR_]], align 41540// CHECK3-NEXT: [[TMP1:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 41541// CHECK3-NEXT: [[SUB:%.*]] = sub nsw i32 [[TMP1]], 01542// CHECK3-NEXT: [[DIV:%.*]] = sdiv i32 [[SUB]], 11543// CHECK3-NEXT: [[SUB2:%.*]] = sub nsw i32 [[DIV]], 11544// CHECK3-NEXT: store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 41545// CHECK3-NEXT: store i32 0, ptr [[I]], align 41546// CHECK3-NEXT: [[TMP2:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 41547// CHECK3-NEXT: [[CMP:%.*]] = icmp slt i32 0, [[TMP2]]1548// CHECK3-NEXT: br i1 [[CMP]], label [[OMP_PRECOND_THEN:%.*]], label [[OMP_PRECOND_END:%.*]]1549// CHECK3: omp.precond.then:1550// CHECK3-NEXT: store i32 0, ptr [[DOTOMP_COMB_LB]], align 41551// CHECK3-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41552// CHECK3-NEXT: store i32 [[TMP3]], ptr [[DOTOMP_COMB_UB]], align 41553// CHECK3-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 41554// CHECK3-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 41555// CHECK3-NEXT: [[TMP4:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 41556// CHECK3-NEXT: [[TMP5:%.*]] = load i32, ptr [[TMP4]], align 41557// CHECK3-NEXT: call void @__kmpc_for_static_init_4(ptr @[[GLOB2:[0-9]+]], i32 [[TMP5]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_COMB_LB]], ptr [[DOTOMP_COMB_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)1558// CHECK3-NEXT: [[TMP6:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 41559// CHECK3-NEXT: [[TMP7:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41560// CHECK3-NEXT: [[CMP4:%.*]] = icmp sgt i32 [[TMP6]], [[TMP7]]1561// CHECK3-NEXT: br i1 [[CMP4]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]1562// CHECK3: cond.true:1563// CHECK3-NEXT: [[TMP8:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41564// CHECK3-NEXT: br label [[COND_END:%.*]]1565// CHECK3: cond.false:1566// CHECK3-NEXT: [[TMP9:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 41567// CHECK3-NEXT: br label [[COND_END]]1568// CHECK3: cond.end:1569// CHECK3-NEXT: [[COND:%.*]] = phi i32 [ [[TMP8]], [[COND_TRUE]] ], [ [[TMP9]], [[COND_FALSE]] ]1570// CHECK3-NEXT: store i32 [[COND]], ptr [[DOTOMP_COMB_UB]], align 41571// CHECK3-NEXT: [[TMP10:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 41572// CHECK3-NEXT: store i32 [[TMP10]], ptr [[DOTOMP_IV]], align 41573// CHECK3-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]1574// CHECK3: omp.inner.for.cond:1575// CHECK3-NEXT: [[TMP11:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP18:![0-9]+]]1576// CHECK3-NEXT: [[TMP12:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4, !llvm.access.group [[ACC_GRP18]]1577// CHECK3-NEXT: [[CMP5:%.*]] = icmp sle i32 [[TMP11]], [[TMP12]]1578// CHECK3-NEXT: br i1 [[CMP5]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]1579// CHECK3: omp.inner.for.body:1580// CHECK3-NEXT: [[TMP13:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 4, !llvm.access.group [[ACC_GRP18]]1581// CHECK3-NEXT: [[TMP14:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4, !llvm.access.group [[ACC_GRP18]]1582// CHECK3-NEXT: [[TMP15:%.*]] = load i32, ptr [[N_ADDR]], align 4, !llvm.access.group [[ACC_GRP18]]1583// CHECK3-NEXT: store i32 [[TMP15]], ptr [[N_CASTED]], align 4, !llvm.access.group [[ACC_GRP18]]1584// CHECK3-NEXT: [[TMP16:%.*]] = load i32, ptr [[N_CASTED]], align 4, !llvm.access.group [[ACC_GRP18]]1585// CHECK3-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB1]], i32 3, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l71.omp_outlined.omp_outlined, i32 [[TMP13]], i32 [[TMP14]], i32 [[TMP16]]), !llvm.access.group [[ACC_GRP18]]1586// CHECK3-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]1587// CHECK3: omp.inner.for.inc:1588// CHECK3-NEXT: [[TMP17:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP18]]1589// CHECK3-NEXT: [[TMP18:%.*]] = load i32, ptr [[DOTOMP_STRIDE]], align 4, !llvm.access.group [[ACC_GRP18]]1590// CHECK3-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP17]], [[TMP18]]1591// CHECK3-NEXT: store i32 [[ADD]], ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP18]]1592// CHECK3-NEXT: br label [[OMP_INNER_FOR_COND]], !llvm.loop [[LOOP19:![0-9]+]]1593// CHECK3: omp.inner.for.end:1594// CHECK3-NEXT: br label [[OMP_LOOP_EXIT:%.*]]1595// CHECK3: omp.loop.exit:1596// CHECK3-NEXT: [[TMP19:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 41597// CHECK3-NEXT: [[TMP20:%.*]] = load i32, ptr [[TMP19]], align 41598// CHECK3-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB2]], i32 [[TMP20]])1599// CHECK3-NEXT: [[TMP21:%.*]] = load i32, ptr [[DOTOMP_IS_LAST]], align 41600// CHECK3-NEXT: [[TMP22:%.*]] = icmp ne i32 [[TMP21]], 01601// CHECK3-NEXT: br i1 [[TMP22]], label [[DOTOMP_FINAL_THEN:%.*]], label [[DOTOMP_FINAL_DONE:%.*]]1602// CHECK3: .omp.final.then:1603// CHECK3-NEXT: [[TMP23:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 41604// CHECK3-NEXT: [[SUB6:%.*]] = sub nsw i32 [[TMP23]], 01605// CHECK3-NEXT: [[DIV7:%.*]] = sdiv i32 [[SUB6]], 11606// CHECK3-NEXT: [[MUL:%.*]] = mul nsw i32 [[DIV7]], 11607// CHECK3-NEXT: [[ADD8:%.*]] = add nsw i32 0, [[MUL]]1608// CHECK3-NEXT: store i32 [[ADD8]], ptr [[I3]], align 41609// CHECK3-NEXT: br label [[DOTOMP_FINAL_DONE]]1610// CHECK3: .omp.final.done:1611// CHECK3-NEXT: br label [[OMP_PRECOND_END]]1612// CHECK3: omp.precond.end:1613// CHECK3-NEXT: ret void1614//1615//1616// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l71.omp_outlined.omp_outlined1617// CHECK3-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i32 noundef [[DOTPREVIOUS_LB_:%.*]], i32 noundef [[DOTPREVIOUS_UB_:%.*]], i32 noundef [[N:%.*]]) #[[ATTR1]] {1618// CHECK3-NEXT: entry:1619// CHECK3-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 41620// CHECK3-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 41621// CHECK3-NEXT: [[DOTPREVIOUS_LB__ADDR:%.*]] = alloca i32, align 41622// CHECK3-NEXT: [[DOTPREVIOUS_UB__ADDR:%.*]] = alloca i32, align 41623// CHECK3-NEXT: [[N_ADDR:%.*]] = alloca i32, align 41624// CHECK3-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 41625// CHECK3-NEXT: [[TMP:%.*]] = alloca i32, align 41626// CHECK3-NEXT: [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 41627// CHECK3-NEXT: [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 41628// CHECK3-NEXT: [[I:%.*]] = alloca i32, align 41629// CHECK3-NEXT: [[DOTOMP_LB:%.*]] = alloca i32, align 41630// CHECK3-NEXT: [[DOTOMP_UB:%.*]] = alloca i32, align 41631// CHECK3-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 41632// CHECK3-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 41633// CHECK3-NEXT: [[I3:%.*]] = alloca i32, align 41634// CHECK3-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 41635// CHECK3-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 41636// CHECK3-NEXT: store i32 [[DOTPREVIOUS_LB_]], ptr [[DOTPREVIOUS_LB__ADDR]], align 41637// CHECK3-NEXT: store i32 [[DOTPREVIOUS_UB_]], ptr [[DOTPREVIOUS_UB__ADDR]], align 41638// CHECK3-NEXT: store i32 [[N]], ptr [[N_ADDR]], align 41639// CHECK3-NEXT: [[TMP0:%.*]] = load i32, ptr [[N_ADDR]], align 41640// CHECK3-NEXT: store i32 [[TMP0]], ptr [[DOTCAPTURE_EXPR_]], align 41641// CHECK3-NEXT: [[TMP1:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 41642// CHECK3-NEXT: [[SUB:%.*]] = sub nsw i32 [[TMP1]], 01643// CHECK3-NEXT: [[DIV:%.*]] = sdiv i32 [[SUB]], 11644// CHECK3-NEXT: [[SUB2:%.*]] = sub nsw i32 [[DIV]], 11645// CHECK3-NEXT: store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 41646// CHECK3-NEXT: store i32 0, ptr [[I]], align 41647// CHECK3-NEXT: [[TMP2:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 41648// CHECK3-NEXT: [[CMP:%.*]] = icmp slt i32 0, [[TMP2]]1649// CHECK3-NEXT: br i1 [[CMP]], label [[OMP_PRECOND_THEN:%.*]], label [[OMP_PRECOND_END:%.*]]1650// CHECK3: omp.precond.then:1651// CHECK3-NEXT: store i32 0, ptr [[DOTOMP_LB]], align 41652// CHECK3-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41653// CHECK3-NEXT: store i32 [[TMP3]], ptr [[DOTOMP_UB]], align 41654// CHECK3-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTPREVIOUS_LB__ADDR]], align 41655// CHECK3-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTPREVIOUS_UB__ADDR]], align 41656// CHECK3-NEXT: store i32 [[TMP4]], ptr [[DOTOMP_LB]], align 41657// CHECK3-NEXT: store i32 [[TMP5]], ptr [[DOTOMP_UB]], align 41658// CHECK3-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 41659// CHECK3-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 41660// CHECK3-NEXT: [[TMP6:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 41661// CHECK3-NEXT: [[TMP7:%.*]] = load i32, ptr [[TMP6]], align 41662// CHECK3-NEXT: call void @__kmpc_for_static_init_4(ptr @[[GLOB3:[0-9]+]], i32 [[TMP7]], i32 34, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)1663// CHECK3-NEXT: [[TMP8:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41664// CHECK3-NEXT: [[TMP9:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41665// CHECK3-NEXT: [[CMP4:%.*]] = icmp sgt i32 [[TMP8]], [[TMP9]]1666// CHECK3-NEXT: br i1 [[CMP4]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]1667// CHECK3: cond.true:1668// CHECK3-NEXT: [[TMP10:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41669// CHECK3-NEXT: br label [[COND_END:%.*]]1670// CHECK3: cond.false:1671// CHECK3-NEXT: [[TMP11:%.*]] = load i32, ptr [[DOTOMP_UB]], align 41672// CHECK3-NEXT: br label [[COND_END]]1673// CHECK3: cond.end:1674// CHECK3-NEXT: [[COND:%.*]] = phi i32 [ [[TMP10]], [[COND_TRUE]] ], [ [[TMP11]], [[COND_FALSE]] ]1675// CHECK3-NEXT: store i32 [[COND]], ptr [[DOTOMP_UB]], align 41676// CHECK3-NEXT: [[TMP12:%.*]] = load i32, ptr [[DOTOMP_LB]], align 41677// CHECK3-NEXT: store i32 [[TMP12]], ptr [[DOTOMP_IV]], align 41678// CHECK3-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]1679// CHECK3: omp.inner.for.cond:1680// CHECK3-NEXT: [[TMP13:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP22:![0-9]+]]1681// CHECK3-NEXT: [[TMP14:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4, !llvm.access.group [[ACC_GRP22]]1682// CHECK3-NEXT: [[CMP5:%.*]] = icmp sle i32 [[TMP13]], [[TMP14]]1683// CHECK3-NEXT: br i1 [[CMP5]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]1684// CHECK3: omp.inner.for.body:1685// CHECK3-NEXT: [[TMP15:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP22]]1686// CHECK3-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP15]], 11687// CHECK3-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]1688// CHECK3-NEXT: store i32 [[ADD]], ptr [[I3]], align 4, !llvm.access.group [[ACC_GRP22]]1689// CHECK3-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]1690// CHECK3: omp.body.continue:1691// CHECK3-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]1692// CHECK3: omp.inner.for.inc:1693// CHECK3-NEXT: [[TMP16:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP22]]1694// CHECK3-NEXT: [[ADD6:%.*]] = add nsw i32 [[TMP16]], 11695// CHECK3-NEXT: store i32 [[ADD6]], ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP22]]1696// CHECK3-NEXT: br label [[OMP_INNER_FOR_COND]], !llvm.loop [[LOOP23:![0-9]+]]1697// CHECK3: omp.inner.for.end:1698// CHECK3-NEXT: br label [[OMP_LOOP_EXIT:%.*]]1699// CHECK3: omp.loop.exit:1700// CHECK3-NEXT: [[TMP17:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 41701// CHECK3-NEXT: [[TMP18:%.*]] = load i32, ptr [[TMP17]], align 41702// CHECK3-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB3]], i32 [[TMP18]])1703// CHECK3-NEXT: [[TMP19:%.*]] = load i32, ptr [[DOTOMP_IS_LAST]], align 41704// CHECK3-NEXT: [[TMP20:%.*]] = icmp ne i32 [[TMP19]], 01705// CHECK3-NEXT: br i1 [[TMP20]], label [[DOTOMP_FINAL_THEN:%.*]], label [[DOTOMP_FINAL_DONE:%.*]]1706// CHECK3: .omp.final.then:1707// CHECK3-NEXT: [[TMP21:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 41708// CHECK3-NEXT: [[SUB7:%.*]] = sub nsw i32 [[TMP21]], 01709// CHECK3-NEXT: [[DIV8:%.*]] = sdiv i32 [[SUB7]], 11710// CHECK3-NEXT: [[MUL9:%.*]] = mul nsw i32 [[DIV8]], 11711// CHECK3-NEXT: [[ADD10:%.*]] = add nsw i32 0, [[MUL9]]1712// CHECK3-NEXT: store i32 [[ADD10]], ptr [[I3]], align 41713// CHECK3-NEXT: br label [[DOTOMP_FINAL_DONE]]1714// CHECK3: .omp.final.done:1715// CHECK3-NEXT: br label [[OMP_PRECOND_END]]1716// CHECK3: omp.precond.end:1717// CHECK3-NEXT: ret void1718//1719//1720// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l751721// CHECK3-SAME: (i32 noundef [[DOTCAPTURE_EXPR_:%.*]]) #[[ATTR1]] {1722// CHECK3-NEXT: entry:1723// CHECK3-NEXT: [[DOTCAPTURE_EXPR__ADDR:%.*]] = alloca i32, align 41724// CHECK3-NEXT: store i32 [[DOTCAPTURE_EXPR_]], ptr [[DOTCAPTURE_EXPR__ADDR]], align 41725// CHECK3-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB1]], i32 0, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l75.omp_outlined)1726// CHECK3-NEXT: ret void1727//1728//1729// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l75.omp_outlined1730// CHECK3-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR1]] {1731// CHECK3-NEXT: entry:1732// CHECK3-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 41733// CHECK3-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 41734// CHECK3-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 41735// CHECK3-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 41736// CHECK3-NEXT: ret void1737//1738//1739// CHECK3-LABEL: define {{[^@]+}}@.omp_task_privates_map.1740// CHECK3-SAME: (ptr noalias noundef [[TMP0:%.*]], ptr noalias noundef [[TMP1:%.*]], ptr noalias noundef [[TMP2:%.*]], ptr noalias noundef [[TMP3:%.*]]) #[[ATTR3:[0-9]+]] {1741// CHECK3-NEXT: entry:1742// CHECK3-NEXT: [[DOTADDR:%.*]] = alloca ptr, align 41743// CHECK3-NEXT: [[DOTADDR1:%.*]] = alloca ptr, align 41744// CHECK3-NEXT: [[DOTADDR2:%.*]] = alloca ptr, align 41745// CHECK3-NEXT: [[DOTADDR3:%.*]] = alloca ptr, align 41746// CHECK3-NEXT: store ptr [[TMP0]], ptr [[DOTADDR]], align 41747// CHECK3-NEXT: store ptr [[TMP1]], ptr [[DOTADDR1]], align 41748// CHECK3-NEXT: store ptr [[TMP2]], ptr [[DOTADDR2]], align 41749// CHECK3-NEXT: store ptr [[TMP3]], ptr [[DOTADDR3]], align 41750// CHECK3-NEXT: [[TMP4:%.*]] = load ptr, ptr [[DOTADDR]], align 41751// CHECK3-NEXT: [[TMP5:%.*]] = getelementptr inbounds nuw [[STRUCT__KMP_PRIVATES_T:%.*]], ptr [[TMP4]], i32 0, i32 01752// CHECK3-NEXT: [[TMP6:%.*]] = load ptr, ptr [[DOTADDR3]], align 41753// CHECK3-NEXT: store ptr [[TMP5]], ptr [[TMP6]], align 41754// CHECK3-NEXT: [[TMP7:%.*]] = getelementptr inbounds nuw [[STRUCT__KMP_PRIVATES_T]], ptr [[TMP4]], i32 0, i32 11755// CHECK3-NEXT: [[TMP8:%.*]] = load ptr, ptr [[DOTADDR1]], align 41756// CHECK3-NEXT: store ptr [[TMP7]], ptr [[TMP8]], align 41757// CHECK3-NEXT: [[TMP9:%.*]] = getelementptr inbounds nuw [[STRUCT__KMP_PRIVATES_T]], ptr [[TMP4]], i32 0, i32 21758// CHECK3-NEXT: [[TMP10:%.*]] = load ptr, ptr [[DOTADDR2]], align 41759// CHECK3-NEXT: store ptr [[TMP9]], ptr [[TMP10]], align 41760// CHECK3-NEXT: ret void1761//1762//1763// CHECK3-LABEL: define {{[^@]+}}@.omp_task_entry.1764// CHECK3-SAME: (i32 noundef [[TMP0:%.*]], ptr noalias noundef [[TMP1:%.*]]) #[[ATTR4:[0-9]+]] {1765// CHECK3-NEXT: entry:1766// CHECK3-NEXT: [[DOTGLOBAL_TID__ADDR_I:%.*]] = alloca i32, align 41767// CHECK3-NEXT: [[DOTPART_ID__ADDR_I:%.*]] = alloca ptr, align 41768// CHECK3-NEXT: [[DOTPRIVATES__ADDR_I:%.*]] = alloca ptr, align 41769// CHECK3-NEXT: [[DOTCOPY_FN__ADDR_I:%.*]] = alloca ptr, align 41770// CHECK3-NEXT: [[DOTTASK_T__ADDR_I:%.*]] = alloca ptr, align 41771// CHECK3-NEXT: [[__CONTEXT_ADDR_I:%.*]] = alloca ptr, align 41772// CHECK3-NEXT: [[DOTFIRSTPRIV_PTR_ADDR_I:%.*]] = alloca ptr, align 41773// CHECK3-NEXT: [[DOTFIRSTPRIV_PTR_ADDR1_I:%.*]] = alloca ptr, align 41774// CHECK3-NEXT: [[DOTFIRSTPRIV_PTR_ADDR2_I:%.*]] = alloca ptr, align 41775// CHECK3-NEXT: [[KERNEL_ARGS_I:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 81776// CHECK3-NEXT: [[DOTCAPTURE_EXPR__CASTED_I:%.*]] = alloca i32, align 41777// CHECK3-NEXT: [[DOTADDR:%.*]] = alloca i32, align 41778// CHECK3-NEXT: [[DOTADDR1:%.*]] = alloca ptr, align 41779// CHECK3-NEXT: store i32 [[TMP0]], ptr [[DOTADDR]], align 41780// CHECK3-NEXT: store ptr [[TMP1]], ptr [[DOTADDR1]], align 41781// CHECK3-NEXT: [[TMP2:%.*]] = load i32, ptr [[DOTADDR]], align 41782// CHECK3-NEXT: [[TMP3:%.*]] = load ptr, ptr [[DOTADDR1]], align 41783// CHECK3-NEXT: [[TMP4:%.*]] = getelementptr inbounds nuw [[STRUCT_KMP_TASK_T_WITH_PRIVATES:%.*]], ptr [[TMP3]], i32 0, i32 01784// CHECK3-NEXT: [[TMP5:%.*]] = getelementptr inbounds nuw [[STRUCT_KMP_TASK_T:%.*]], ptr [[TMP4]], i32 0, i32 21785// CHECK3-NEXT: [[TMP6:%.*]] = getelementptr inbounds nuw [[STRUCT_KMP_TASK_T]], ptr [[TMP4]], i32 0, i32 01786// CHECK3-NEXT: [[TMP7:%.*]] = load ptr, ptr [[TMP6]], align 41787// CHECK3-NEXT: [[TMP8:%.*]] = getelementptr inbounds nuw [[STRUCT_KMP_TASK_T_WITH_PRIVATES]], ptr [[TMP3]], i32 0, i32 11788// CHECK3-NEXT: call void @llvm.experimental.noalias.scope.decl(metadata [[META25:![0-9]+]])1789// CHECK3-NEXT: call void @llvm.experimental.noalias.scope.decl(metadata [[META28:![0-9]+]])1790// CHECK3-NEXT: call void @llvm.experimental.noalias.scope.decl(metadata [[META30:![0-9]+]])1791// CHECK3-NEXT: call void @llvm.experimental.noalias.scope.decl(metadata [[META32:![0-9]+]])1792// CHECK3-NEXT: store i32 [[TMP2]], ptr [[DOTGLOBAL_TID__ADDR_I]], align 4, !noalias [[META34:![0-9]+]]1793// CHECK3-NEXT: store ptr [[TMP5]], ptr [[DOTPART_ID__ADDR_I]], align 4, !noalias [[META34]]1794// CHECK3-NEXT: store ptr [[TMP8]], ptr [[DOTPRIVATES__ADDR_I]], align 4, !noalias [[META34]]1795// CHECK3-NEXT: store ptr @.omp_task_privates_map., ptr [[DOTCOPY_FN__ADDR_I]], align 4, !noalias [[META34]]1796// CHECK3-NEXT: store ptr [[TMP3]], ptr [[DOTTASK_T__ADDR_I]], align 4, !noalias [[META34]]1797// CHECK3-NEXT: store ptr [[TMP7]], ptr [[__CONTEXT_ADDR_I]], align 4, !noalias [[META34]]1798// CHECK3-NEXT: [[TMP9:%.*]] = load ptr, ptr [[__CONTEXT_ADDR_I]], align 4, !noalias [[META34]]1799// CHECK3-NEXT: [[TMP10:%.*]] = load ptr, ptr [[DOTCOPY_FN__ADDR_I]], align 4, !noalias [[META34]]1800// CHECK3-NEXT: [[TMP11:%.*]] = load ptr, ptr [[DOTPRIVATES__ADDR_I]], align 4, !noalias [[META34]]1801// CHECK3-NEXT: call void [[TMP10]](ptr [[TMP11]], ptr [[DOTFIRSTPRIV_PTR_ADDR_I]], ptr [[DOTFIRSTPRIV_PTR_ADDR1_I]], ptr [[DOTFIRSTPRIV_PTR_ADDR2_I]]) #[[ATTR2]]1802// CHECK3-NEXT: [[TMP12:%.*]] = load ptr, ptr [[DOTFIRSTPRIV_PTR_ADDR_I]], align 4, !noalias [[META34]]1803// CHECK3-NEXT: [[TMP13:%.*]] = load ptr, ptr [[DOTFIRSTPRIV_PTR_ADDR1_I]], align 4, !noalias [[META34]]1804// CHECK3-NEXT: [[TMP14:%.*]] = load ptr, ptr [[DOTFIRSTPRIV_PTR_ADDR2_I]], align 4, !noalias [[META34]]1805// CHECK3-NEXT: [[TMP15:%.*]] = load i32, ptr [[TMP9]], align 41806// CHECK3-NEXT: store i32 3, ptr [[KERNEL_ARGS_I]], align 4, !noalias [[META34]]1807// CHECK3-NEXT: [[TMP16:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS_I]], i32 0, i32 11808// CHECK3-NEXT: store i32 1, ptr [[TMP16]], align 4, !noalias [[META34]]1809// CHECK3-NEXT: [[TMP17:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS_I]], i32 0, i32 21810// CHECK3-NEXT: store ptr [[TMP12]], ptr [[TMP17]], align 4, !noalias [[META34]]1811// CHECK3-NEXT: [[TMP18:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS_I]], i32 0, i32 31812// CHECK3-NEXT: store ptr [[TMP13]], ptr [[TMP18]], align 4, !noalias [[META34]]1813// CHECK3-NEXT: [[TMP19:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS_I]], i32 0, i32 41814// CHECK3-NEXT: store ptr [[TMP14]], ptr [[TMP19]], align 4, !noalias [[META34]]1815// CHECK3-NEXT: [[TMP20:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS_I]], i32 0, i32 51816// CHECK3-NEXT: store ptr @.offload_maptypes.6, ptr [[TMP20]], align 4, !noalias [[META34]]1817// CHECK3-NEXT: [[TMP21:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS_I]], i32 0, i32 61818// CHECK3-NEXT: store ptr null, ptr [[TMP21]], align 4, !noalias [[META34]]1819// CHECK3-NEXT: [[TMP22:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS_I]], i32 0, i32 71820// CHECK3-NEXT: store ptr null, ptr [[TMP22]], align 4, !noalias [[META34]]1821// CHECK3-NEXT: [[TMP23:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS_I]], i32 0, i32 81822// CHECK3-NEXT: store i64 0, ptr [[TMP23]], align 8, !noalias [[META34]]1823// CHECK3-NEXT: [[TMP24:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS_I]], i32 0, i32 91824// CHECK3-NEXT: store i64 9, ptr [[TMP24]], align 8, !noalias [[META34]]1825// CHECK3-NEXT: [[TMP25:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS_I]], i32 0, i32 101826// CHECK3-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP25]], align 4, !noalias [[META34]]1827// CHECK3-NEXT: [[TMP26:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS_I]], i32 0, i32 111828// CHECK3-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP26]], align 4, !noalias [[META34]]1829// CHECK3-NEXT: [[TMP27:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS_I]], i32 0, i32 121830// CHECK3-NEXT: store i32 [[TMP15]], ptr [[TMP27]], align 4, !noalias [[META34]]1831// CHECK3-NEXT: [[TMP28:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 -1, i32 0, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l75.region_id, ptr [[KERNEL_ARGS_I]])1832// CHECK3-NEXT: [[TMP29:%.*]] = icmp ne i32 [[TMP28]], 01833// CHECK3-NEXT: br i1 [[TMP29]], label [[OMP_OFFLOAD_FAILED_I:%.*]], label [[DOTOMP_OUTLINED__EXIT:%.*]]1834// CHECK3: omp_offload.failed.i:1835// CHECK3-NEXT: [[TMP30:%.*]] = load i32, ptr [[TMP9]], align 41836// CHECK3-NEXT: store i32 [[TMP30]], ptr [[DOTCAPTURE_EXPR__CASTED_I]], align 4, !noalias [[META34]]1837// CHECK3-NEXT: [[TMP31:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR__CASTED_I]], align 4, !noalias [[META34]]1838// CHECK3-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l75(i32 [[TMP31]]) #[[ATTR2]]1839// CHECK3-NEXT: br label [[DOTOMP_OUTLINED__EXIT]]1840// CHECK3: .omp_outlined..exit:1841// CHECK3-NEXT: ret i32 01842//1843//1844// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l551845// CHECK3-SAME: () #[[ATTR1]] {1846// CHECK3-NEXT: entry:1847// CHECK3-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB1]], i32 0, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l55.omp_outlined)1848// CHECK3-NEXT: ret void1849//1850//1851// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l55.omp_outlined1852// CHECK3-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR1]] {1853// CHECK3-NEXT: entry:1854// CHECK3-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 41855// CHECK3-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 41856// CHECK3-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 41857// CHECK3-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 41858// CHECK3-NEXT: ret void1859//1860//1861// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l601862// CHECK3-SAME: (i32 noundef [[A:%.*]], i32 noundef [[B:%.*]], i32 noundef [[DOTCAPTURE_EXPR_:%.*]]) #[[ATTR1]] {1863// CHECK3-NEXT: entry:1864// CHECK3-NEXT: [[A_ADDR:%.*]] = alloca i32, align 41865// CHECK3-NEXT: [[B_ADDR:%.*]] = alloca i32, align 41866// CHECK3-NEXT: [[DOTCAPTURE_EXPR__ADDR:%.*]] = alloca i32, align 41867// CHECK3-NEXT: [[A_CASTED:%.*]] = alloca i32, align 41868// CHECK3-NEXT: [[B_CASTED:%.*]] = alloca i32, align 41869// CHECK3-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])1870// CHECK3-NEXT: store i32 [[A]], ptr [[A_ADDR]], align 41871// CHECK3-NEXT: store i32 [[B]], ptr [[B_ADDR]], align 41872// CHECK3-NEXT: store i32 [[DOTCAPTURE_EXPR_]], ptr [[DOTCAPTURE_EXPR__ADDR]], align 41873// CHECK3-NEXT: [[TMP1:%.*]] = load i16, ptr [[DOTCAPTURE_EXPR__ADDR]], align 21874// CHECK3-NEXT: [[TMP2:%.*]] = sext i16 [[TMP1]] to i321875// CHECK3-NEXT: call void @__kmpc_push_num_teams(ptr @[[GLOB1]], i32 [[TMP0]], i32 [[TMP2]], i32 0)1876// CHECK3-NEXT: [[TMP3:%.*]] = load i32, ptr [[A_ADDR]], align 41877// CHECK3-NEXT: store i32 [[TMP3]], ptr [[A_CASTED]], align 41878// CHECK3-NEXT: [[TMP4:%.*]] = load i32, ptr [[A_CASTED]], align 41879// CHECK3-NEXT: [[TMP5:%.*]] = load i16, ptr [[B_ADDR]], align 21880// CHECK3-NEXT: store i16 [[TMP5]], ptr [[B_CASTED]], align 21881// CHECK3-NEXT: [[TMP6:%.*]] = load i32, ptr [[B_CASTED]], align 41882// CHECK3-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB1]], i32 2, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l60.omp_outlined, i32 [[TMP4]], i32 [[TMP6]])1883// CHECK3-NEXT: ret void1884//1885//1886// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l60.omp_outlined1887// CHECK3-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i32 noundef [[A:%.*]], i32 noundef [[B:%.*]]) #[[ATTR1]] {1888// CHECK3-NEXT: entry:1889// CHECK3-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 41890// CHECK3-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 41891// CHECK3-NEXT: [[A_ADDR:%.*]] = alloca i32, align 41892// CHECK3-NEXT: [[B_ADDR:%.*]] = alloca i32, align 41893// CHECK3-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 41894// CHECK3-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 41895// CHECK3-NEXT: store i32 [[A]], ptr [[A_ADDR]], align 41896// CHECK3-NEXT: store i32 [[B]], ptr [[B_ADDR]], align 41897// CHECK3-NEXT: [[TMP0:%.*]] = load i16, ptr [[B_ADDR]], align 21898// CHECK3-NEXT: [[CONV:%.*]] = sext i16 [[TMP0]] to i321899// CHECK3-NEXT: [[TMP1:%.*]] = load i32, ptr [[A_ADDR]], align 41900// CHECK3-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP1]], [[CONV]]1901// CHECK3-NEXT: store i32 [[ADD]], ptr [[A_ADDR]], align 41902// CHECK3-NEXT: ret void1903//1904//1905// CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l711906// CHECK9-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], i64 noundef [[N:%.*]], i64 noundef [[DOTCAPTURE_EXPR_:%.*]], i64 noundef [[DOTCAPTURE_EXPR_1:%.*]]) #[[ATTR0:[0-9]+]] {1907// CHECK9-NEXT: entry:1908// CHECK9-NEXT: [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 81909// CHECK9-NEXT: [[N_ADDR:%.*]] = alloca i64, align 81910// CHECK9-NEXT: [[DOTCAPTURE_EXPR__ADDR:%.*]] = alloca i64, align 81911// CHECK9-NEXT: [[DOTCAPTURE_EXPR__ADDR2:%.*]] = alloca i64, align 81912// CHECK9-NEXT: [[N_CASTED:%.*]] = alloca i64, align 81913// CHECK9-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB3:[0-9]+]])1914// CHECK9-NEXT: store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 81915// CHECK9-NEXT: store i64 [[N]], ptr [[N_ADDR]], align 81916// CHECK9-NEXT: store i64 [[DOTCAPTURE_EXPR_]], ptr [[DOTCAPTURE_EXPR__ADDR]], align 81917// CHECK9-NEXT: store i64 [[DOTCAPTURE_EXPR_1]], ptr [[DOTCAPTURE_EXPR__ADDR2]], align 81918// CHECK9-NEXT: [[TMP1:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR__ADDR]], align 41919// CHECK9-NEXT: call void @__kmpc_push_num_teams(ptr @[[GLOB3]], i32 [[TMP0]], i32 [[TMP1]], i32 0)1920// CHECK9-NEXT: [[TMP2:%.*]] = load i32, ptr [[N_ADDR]], align 41921// CHECK9-NEXT: store i32 [[TMP2]], ptr [[N_CASTED]], align 41922// CHECK9-NEXT: [[TMP3:%.*]] = load i64, ptr [[N_CASTED]], align 81923// CHECK9-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB3]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l71.omp_outlined, i64 [[TMP3]])1924// CHECK9-NEXT: ret void1925//1926//1927// CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l71.omp_outlined1928// CHECK9-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i64 noundef [[N:%.*]]) #[[ATTR0]] {1929// CHECK9-NEXT: entry:1930// CHECK9-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 81931// CHECK9-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 81932// CHECK9-NEXT: [[N_ADDR:%.*]] = alloca i64, align 81933// CHECK9-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 41934// CHECK9-NEXT: [[TMP:%.*]] = alloca i32, align 41935// CHECK9-NEXT: [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 41936// CHECK9-NEXT: [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 41937// CHECK9-NEXT: [[I:%.*]] = alloca i32, align 41938// CHECK9-NEXT: [[DOTOMP_COMB_LB:%.*]] = alloca i32, align 41939// CHECK9-NEXT: [[DOTOMP_COMB_UB:%.*]] = alloca i32, align 41940// CHECK9-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 41941// CHECK9-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 41942// CHECK9-NEXT: [[I3:%.*]] = alloca i32, align 41943// CHECK9-NEXT: [[N_CASTED:%.*]] = alloca i64, align 81944// CHECK9-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 81945// CHECK9-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 81946// CHECK9-NEXT: store i64 [[N]], ptr [[N_ADDR]], align 81947// CHECK9-NEXT: [[TMP0:%.*]] = load i32, ptr [[N_ADDR]], align 41948// CHECK9-NEXT: store i32 [[TMP0]], ptr [[DOTCAPTURE_EXPR_]], align 41949// CHECK9-NEXT: [[TMP1:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 41950// CHECK9-NEXT: [[SUB:%.*]] = sub nsw i32 [[TMP1]], 01951// CHECK9-NEXT: [[DIV:%.*]] = sdiv i32 [[SUB]], 11952// CHECK9-NEXT: [[SUB2:%.*]] = sub nsw i32 [[DIV]], 11953// CHECK9-NEXT: store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 41954// CHECK9-NEXT: store i32 0, ptr [[I]], align 41955// CHECK9-NEXT: [[TMP2:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 41956// CHECK9-NEXT: [[CMP:%.*]] = icmp slt i32 0, [[TMP2]]1957// CHECK9-NEXT: br i1 [[CMP]], label [[OMP_PRECOND_THEN:%.*]], label [[OMP_PRECOND_END:%.*]]1958// CHECK9: omp.precond.then:1959// CHECK9-NEXT: store i32 0, ptr [[DOTOMP_COMB_LB]], align 41960// CHECK9-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41961// CHECK9-NEXT: store i32 [[TMP3]], ptr [[DOTOMP_COMB_UB]], align 41962// CHECK9-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 41963// CHECK9-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 41964// CHECK9-NEXT: [[TMP4:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 81965// CHECK9-NEXT: [[TMP5:%.*]] = load i32, ptr [[TMP4]], align 41966// CHECK9-NEXT: call void @__kmpc_for_static_init_4(ptr @[[GLOB1:[0-9]+]], i32 [[TMP5]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_COMB_LB]], ptr [[DOTOMP_COMB_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)1967// CHECK9-NEXT: [[TMP6:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 41968// CHECK9-NEXT: [[TMP7:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41969// CHECK9-NEXT: [[CMP4:%.*]] = icmp sgt i32 [[TMP6]], [[TMP7]]1970// CHECK9-NEXT: br i1 [[CMP4]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]1971// CHECK9: cond.true:1972// CHECK9-NEXT: [[TMP8:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 41973// CHECK9-NEXT: br label [[COND_END:%.*]]1974// CHECK9: cond.false:1975// CHECK9-NEXT: [[TMP9:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 41976// CHECK9-NEXT: br label [[COND_END]]1977// CHECK9: cond.end:1978// CHECK9-NEXT: [[COND:%.*]] = phi i32 [ [[TMP8]], [[COND_TRUE]] ], [ [[TMP9]], [[COND_FALSE]] ]1979// CHECK9-NEXT: store i32 [[COND]], ptr [[DOTOMP_COMB_UB]], align 41980// CHECK9-NEXT: [[TMP10:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 41981// CHECK9-NEXT: store i32 [[TMP10]], ptr [[DOTOMP_IV]], align 41982// CHECK9-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]1983// CHECK9: omp.inner.for.cond:1984// CHECK9-NEXT: [[TMP11:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP16:![0-9]+]]1985// CHECK9-NEXT: [[TMP12:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4, !llvm.access.group [[ACC_GRP16]]1986// CHECK9-NEXT: [[CMP5:%.*]] = icmp sle i32 [[TMP11]], [[TMP12]]1987// CHECK9-NEXT: br i1 [[CMP5]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]1988// CHECK9: omp.inner.for.body:1989// CHECK9-NEXT: [[TMP13:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 4, !llvm.access.group [[ACC_GRP16]]1990// CHECK9-NEXT: [[TMP14:%.*]] = zext i32 [[TMP13]] to i641991// CHECK9-NEXT: [[TMP15:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4, !llvm.access.group [[ACC_GRP16]]1992// CHECK9-NEXT: [[TMP16:%.*]] = zext i32 [[TMP15]] to i641993// CHECK9-NEXT: [[TMP17:%.*]] = load i32, ptr [[N_ADDR]], align 4, !llvm.access.group [[ACC_GRP16]]1994// CHECK9-NEXT: store i32 [[TMP17]], ptr [[N_CASTED]], align 4, !llvm.access.group [[ACC_GRP16]]1995// CHECK9-NEXT: [[TMP18:%.*]] = load i64, ptr [[N_CASTED]], align 8, !llvm.access.group [[ACC_GRP16]]1996// CHECK9-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB3]], i32 3, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l71.omp_outlined.omp_outlined, i64 [[TMP14]], i64 [[TMP16]], i64 [[TMP18]]), !llvm.access.group [[ACC_GRP16]]1997// CHECK9-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]1998// CHECK9: omp.inner.for.inc:1999// CHECK9-NEXT: [[TMP19:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP16]]2000// CHECK9-NEXT: [[TMP20:%.*]] = load i32, ptr [[DOTOMP_STRIDE]], align 4, !llvm.access.group [[ACC_GRP16]]2001// CHECK9-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP19]], [[TMP20]]2002// CHECK9-NEXT: store i32 [[ADD]], ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP16]]2003// CHECK9-NEXT: br label [[OMP_INNER_FOR_COND]], !llvm.loop [[LOOP17:![0-9]+]]2004// CHECK9: omp.inner.for.end:2005// CHECK9-NEXT: br label [[OMP_LOOP_EXIT:%.*]]2006// CHECK9: omp.loop.exit:2007// CHECK9-NEXT: [[TMP21:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 82008// CHECK9-NEXT: [[TMP22:%.*]] = load i32, ptr [[TMP21]], align 42009// CHECK9-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP22]])2010// CHECK9-NEXT: [[TMP23:%.*]] = load i32, ptr [[DOTOMP_IS_LAST]], align 42011// CHECK9-NEXT: [[TMP24:%.*]] = icmp ne i32 [[TMP23]], 02012// CHECK9-NEXT: br i1 [[TMP24]], label [[DOTOMP_FINAL_THEN:%.*]], label [[DOTOMP_FINAL_DONE:%.*]]2013// CHECK9: .omp.final.then:2014// CHECK9-NEXT: [[TMP25:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 42015// CHECK9-NEXT: [[SUB6:%.*]] = sub nsw i32 [[TMP25]], 02016// CHECK9-NEXT: [[DIV7:%.*]] = sdiv i32 [[SUB6]], 12017// CHECK9-NEXT: [[MUL:%.*]] = mul nsw i32 [[DIV7]], 12018// CHECK9-NEXT: [[ADD8:%.*]] = add nsw i32 0, [[MUL]]2019// CHECK9-NEXT: store i32 [[ADD8]], ptr [[I3]], align 42020// CHECK9-NEXT: br label [[DOTOMP_FINAL_DONE]]2021// CHECK9: .omp.final.done:2022// CHECK9-NEXT: br label [[OMP_PRECOND_END]]2023// CHECK9: omp.precond.end:2024// CHECK9-NEXT: ret void2025//2026//2027// CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l71.omp_outlined.omp_outlined2028// CHECK9-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i64 noundef [[DOTPREVIOUS_LB_:%.*]], i64 noundef [[DOTPREVIOUS_UB_:%.*]], i64 noundef [[N:%.*]]) #[[ATTR0]] {2029// CHECK9-NEXT: entry:2030// CHECK9-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 82031// CHECK9-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 82032// CHECK9-NEXT: [[DOTPREVIOUS_LB__ADDR:%.*]] = alloca i64, align 82033// CHECK9-NEXT: [[DOTPREVIOUS_UB__ADDR:%.*]] = alloca i64, align 82034// CHECK9-NEXT: [[N_ADDR:%.*]] = alloca i64, align 82035// CHECK9-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 42036// CHECK9-NEXT: [[TMP:%.*]] = alloca i32, align 42037// CHECK9-NEXT: [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 42038// CHECK9-NEXT: [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 42039// CHECK9-NEXT: [[I:%.*]] = alloca i32, align 42040// CHECK9-NEXT: [[DOTOMP_LB:%.*]] = alloca i32, align 42041// CHECK9-NEXT: [[DOTOMP_UB:%.*]] = alloca i32, align 42042// CHECK9-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 42043// CHECK9-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 42044// CHECK9-NEXT: [[I4:%.*]] = alloca i32, align 42045// CHECK9-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 82046// CHECK9-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 82047// CHECK9-NEXT: store i64 [[DOTPREVIOUS_LB_]], ptr [[DOTPREVIOUS_LB__ADDR]], align 82048// CHECK9-NEXT: store i64 [[DOTPREVIOUS_UB_]], ptr [[DOTPREVIOUS_UB__ADDR]], align 82049// CHECK9-NEXT: store i64 [[N]], ptr [[N_ADDR]], align 82050// CHECK9-NEXT: [[TMP0:%.*]] = load i32, ptr [[N_ADDR]], align 42051// CHECK9-NEXT: store i32 [[TMP0]], ptr [[DOTCAPTURE_EXPR_]], align 42052// CHECK9-NEXT: [[TMP1:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 42053// CHECK9-NEXT: [[SUB:%.*]] = sub nsw i32 [[TMP1]], 02054// CHECK9-NEXT: [[DIV:%.*]] = sdiv i32 [[SUB]], 12055// CHECK9-NEXT: [[SUB2:%.*]] = sub nsw i32 [[DIV]], 12056// CHECK9-NEXT: store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 42057// CHECK9-NEXT: store i32 0, ptr [[I]], align 42058// CHECK9-NEXT: [[TMP2:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 42059// CHECK9-NEXT: [[CMP:%.*]] = icmp slt i32 0, [[TMP2]]2060// CHECK9-NEXT: br i1 [[CMP]], label [[OMP_PRECOND_THEN:%.*]], label [[OMP_PRECOND_END:%.*]]2061// CHECK9: omp.precond.then:2062// CHECK9-NEXT: store i32 0, ptr [[DOTOMP_LB]], align 42063// CHECK9-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 42064// CHECK9-NEXT: store i32 [[TMP3]], ptr [[DOTOMP_UB]], align 42065// CHECK9-NEXT: [[TMP4:%.*]] = load i64, ptr [[DOTPREVIOUS_LB__ADDR]], align 82066// CHECK9-NEXT: [[CONV:%.*]] = trunc i64 [[TMP4]] to i322067// CHECK9-NEXT: [[TMP5:%.*]] = load i64, ptr [[DOTPREVIOUS_UB__ADDR]], align 82068// CHECK9-NEXT: [[CONV3:%.*]] = trunc i64 [[TMP5]] to i322069// CHECK9-NEXT: store i32 [[CONV]], ptr [[DOTOMP_LB]], align 42070// CHECK9-NEXT: store i32 [[CONV3]], ptr [[DOTOMP_UB]], align 42071// CHECK9-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 42072// CHECK9-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 42073// CHECK9-NEXT: [[TMP6:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 82074// CHECK9-NEXT: [[TMP7:%.*]] = load i32, ptr [[TMP6]], align 42075// CHECK9-NEXT: call void @__kmpc_for_static_init_4(ptr @[[GLOB2:[0-9]+]], i32 [[TMP7]], i32 34, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)2076// CHECK9-NEXT: [[TMP8:%.*]] = load i32, ptr [[DOTOMP_UB]], align 42077// CHECK9-NEXT: [[TMP9:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 42078// CHECK9-NEXT: [[CMP5:%.*]] = icmp sgt i32 [[TMP8]], [[TMP9]]2079// CHECK9-NEXT: br i1 [[CMP5]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]2080// CHECK9: cond.true:2081// CHECK9-NEXT: [[TMP10:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 42082// CHECK9-NEXT: br label [[COND_END:%.*]]2083// CHECK9: cond.false:2084// CHECK9-NEXT: [[TMP11:%.*]] = load i32, ptr [[DOTOMP_UB]], align 42085// CHECK9-NEXT: br label [[COND_END]]2086// CHECK9: cond.end:2087// CHECK9-NEXT: [[COND:%.*]] = phi i32 [ [[TMP10]], [[COND_TRUE]] ], [ [[TMP11]], [[COND_FALSE]] ]2088// CHECK9-NEXT: store i32 [[COND]], ptr [[DOTOMP_UB]], align 42089// CHECK9-NEXT: [[TMP12:%.*]] = load i32, ptr [[DOTOMP_LB]], align 42090// CHECK9-NEXT: store i32 [[TMP12]], ptr [[DOTOMP_IV]], align 42091// CHECK9-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]2092// CHECK9: omp.inner.for.cond:2093// CHECK9-NEXT: [[TMP13:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP20:![0-9]+]]2094// CHECK9-NEXT: [[TMP14:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4, !llvm.access.group [[ACC_GRP20]]2095// CHECK9-NEXT: [[CMP6:%.*]] = icmp sle i32 [[TMP13]], [[TMP14]]2096// CHECK9-NEXT: br i1 [[CMP6]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]2097// CHECK9: omp.inner.for.body:2098// CHECK9-NEXT: [[TMP15:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP20]]2099// CHECK9-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP15]], 12100// CHECK9-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]2101// CHECK9-NEXT: store i32 [[ADD]], ptr [[I4]], align 4, !llvm.access.group [[ACC_GRP20]]2102// CHECK9-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]2103// CHECK9: omp.body.continue:2104// CHECK9-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]2105// CHECK9: omp.inner.for.inc:2106// CHECK9-NEXT: [[TMP16:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP20]]2107// CHECK9-NEXT: [[ADD7:%.*]] = add nsw i32 [[TMP16]], 12108// CHECK9-NEXT: store i32 [[ADD7]], ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP20]]2109// CHECK9-NEXT: br label [[OMP_INNER_FOR_COND]], !llvm.loop [[LOOP21:![0-9]+]]2110// CHECK9: omp.inner.for.end:2111// CHECK9-NEXT: br label [[OMP_LOOP_EXIT:%.*]]2112// CHECK9: omp.loop.exit:2113// CHECK9-NEXT: [[TMP17:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 82114// CHECK9-NEXT: [[TMP18:%.*]] = load i32, ptr [[TMP17]], align 42115// CHECK9-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB2]], i32 [[TMP18]])2116// CHECK9-NEXT: [[TMP19:%.*]] = load i32, ptr [[DOTOMP_IS_LAST]], align 42117// CHECK9-NEXT: [[TMP20:%.*]] = icmp ne i32 [[TMP19]], 02118// CHECK9-NEXT: br i1 [[TMP20]], label [[DOTOMP_FINAL_THEN:%.*]], label [[DOTOMP_FINAL_DONE:%.*]]2119// CHECK9: .omp.final.then:2120// CHECK9-NEXT: [[TMP21:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 42121// CHECK9-NEXT: [[SUB8:%.*]] = sub nsw i32 [[TMP21]], 02122// CHECK9-NEXT: [[DIV9:%.*]] = sdiv i32 [[SUB8]], 12123// CHECK9-NEXT: [[MUL10:%.*]] = mul nsw i32 [[DIV9]], 12124// CHECK9-NEXT: [[ADD11:%.*]] = add nsw i32 0, [[MUL10]]2125// CHECK9-NEXT: store i32 [[ADD11]], ptr [[I4]], align 42126// CHECK9-NEXT: br label [[DOTOMP_FINAL_DONE]]2127// CHECK9: .omp.final.done:2128// CHECK9-NEXT: br label [[OMP_PRECOND_END]]2129// CHECK9: omp.precond.end:2130// CHECK9-NEXT: ret void2131//2132//2133// CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l752134// CHECK9-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], i64 noundef [[DOTCAPTURE_EXPR_:%.*]]) #[[ATTR0]] {2135// CHECK9-NEXT: entry:2136// CHECK9-NEXT: [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 82137// CHECK9-NEXT: [[DOTCAPTURE_EXPR__ADDR:%.*]] = alloca i64, align 82138// CHECK9-NEXT: store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 82139// CHECK9-NEXT: store i64 [[DOTCAPTURE_EXPR_]], ptr [[DOTCAPTURE_EXPR__ADDR]], align 82140// CHECK9-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB3]], i32 0, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l75.omp_outlined)2141// CHECK9-NEXT: ret void2142//2143//2144// CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l75.omp_outlined2145// CHECK9-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR0]] {2146// CHECK9-NEXT: entry:2147// CHECK9-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 82148// CHECK9-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 82149// CHECK9-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 82150// CHECK9-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 82151// CHECK9-NEXT: ret void2152//2153//2154// CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2S12r1Ei_l882155// CHECK9-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], ptr noundef [[THIS:%.*]], i64 noundef [[B:%.*]], i64 noundef [[DOTCAPTURE_EXPR_:%.*]]) #[[ATTR0]] {2156// CHECK9-NEXT: entry:2157// CHECK9-NEXT: [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 82158// CHECK9-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 82159// CHECK9-NEXT: [[B_ADDR:%.*]] = alloca i64, align 82160// CHECK9-NEXT: [[DOTCAPTURE_EXPR__ADDR:%.*]] = alloca i64, align 82161// CHECK9-NEXT: [[B_CASTED:%.*]] = alloca i64, align 82162// CHECK9-NEXT: store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 82163// CHECK9-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 82164// CHECK9-NEXT: store i64 [[B]], ptr [[B_ADDR]], align 82165// CHECK9-NEXT: store i64 [[DOTCAPTURE_EXPR_]], ptr [[DOTCAPTURE_EXPR__ADDR]], align 82166// CHECK9-NEXT: [[TMP0:%.*]] = load ptr, ptr [[THIS_ADDR]], align 82167// CHECK9-NEXT: [[TMP1:%.*]] = load i32, ptr [[B_ADDR]], align 42168// CHECK9-NEXT: store i32 [[TMP1]], ptr [[B_CASTED]], align 42169// CHECK9-NEXT: [[TMP2:%.*]] = load i64, ptr [[B_CASTED]], align 82170// CHECK9-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB3]], i32 2, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2S12r1Ei_l88.omp_outlined, ptr [[TMP0]], i64 [[TMP2]])2171// CHECK9-NEXT: ret void2172//2173//2174// CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2S12r1Ei_l88.omp_outlined2175// CHECK9-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef [[THIS:%.*]], i64 noundef [[B:%.*]]) #[[ATTR0]] {2176// CHECK9-NEXT: entry:2177// CHECK9-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 82178// CHECK9-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 82179// CHECK9-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 82180// CHECK9-NEXT: [[B_ADDR:%.*]] = alloca i64, align 82181// CHECK9-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 82182// CHECK9-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 82183// CHECK9-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 82184// CHECK9-NEXT: store i64 [[B]], ptr [[B_ADDR]], align 82185// CHECK9-NEXT: [[TMP0:%.*]] = load ptr, ptr [[THIS_ADDR]], align 82186// CHECK9-NEXT: [[TMP1:%.*]] = load i32, ptr [[B_ADDR]], align 42187// CHECK9-NEXT: [[CONV:%.*]] = sitofp i32 [[TMP1]] to double2188// CHECK9-NEXT: [[ADD:%.*]] = fadd double [[CONV]], 1.500000e+002189// CHECK9-NEXT: [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_S1:%.*]], ptr [[TMP0]], i32 0, i32 02190// CHECK9-NEXT: store double [[ADD]], ptr [[A]], align 82191// CHECK9-NEXT: ret void2192//2193//2194// CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2S12r1Ei_l932195// CHECK9-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], ptr noundef [[THIS:%.*]]) #[[ATTR0]] {2196// CHECK9-NEXT: entry:2197// CHECK9-NEXT: [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 82198// CHECK9-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 82199// CHECK9-NEXT: store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 82200// CHECK9-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 82201// CHECK9-NEXT: [[TMP0:%.*]] = load ptr, ptr [[THIS_ADDR]], align 82202// CHECK9-NEXT: [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_S1:%.*]], ptr [[TMP0]], i32 0, i32 02203// CHECK9-NEXT: store double 2.500000e+00, ptr [[A]], align 82204// CHECK9-NEXT: ret void2205//2206//2207// CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l552208// CHECK9-SAME: (ptr noalias noundef [[DYN_PTR:%.*]]) #[[ATTR0]] {2209// CHECK9-NEXT: entry:2210// CHECK9-NEXT: [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 82211// CHECK9-NEXT: store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 82212// CHECK9-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB3]], i32 0, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l55.omp_outlined)2213// CHECK9-NEXT: ret void2214//2215//2216// CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l55.omp_outlined2217// CHECK9-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR0]] {2218// CHECK9-NEXT: entry:2219// CHECK9-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 82220// CHECK9-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 82221// CHECK9-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 82222// CHECK9-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 82223// CHECK9-NEXT: ret void2224//2225//2226// CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l602227// CHECK9-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], i64 noundef [[A:%.*]], i64 noundef [[B:%.*]], i64 noundef [[DOTCAPTURE_EXPR_:%.*]]) #[[ATTR0]] {2228// CHECK9-NEXT: entry:2229// CHECK9-NEXT: [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 82230// CHECK9-NEXT: [[A_ADDR:%.*]] = alloca i64, align 82231// CHECK9-NEXT: [[B_ADDR:%.*]] = alloca i64, align 82232// CHECK9-NEXT: [[DOTCAPTURE_EXPR__ADDR:%.*]] = alloca i64, align 82233// CHECK9-NEXT: [[A_CASTED:%.*]] = alloca i64, align 82234// CHECK9-NEXT: [[B_CASTED:%.*]] = alloca i64, align 82235// CHECK9-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB3]])2236// CHECK9-NEXT: store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 82237// CHECK9-NEXT: store i64 [[A]], ptr [[A_ADDR]], align 82238// CHECK9-NEXT: store i64 [[B]], ptr [[B_ADDR]], align 82239// CHECK9-NEXT: store i64 [[DOTCAPTURE_EXPR_]], ptr [[DOTCAPTURE_EXPR__ADDR]], align 82240// CHECK9-NEXT: [[TMP1:%.*]] = load i16, ptr [[DOTCAPTURE_EXPR__ADDR]], align 22241// CHECK9-NEXT: [[TMP2:%.*]] = sext i16 [[TMP1]] to i322242// CHECK9-NEXT: call void @__kmpc_push_num_teams(ptr @[[GLOB3]], i32 [[TMP0]], i32 [[TMP2]], i32 0)2243// CHECK9-NEXT: [[TMP3:%.*]] = load i32, ptr [[A_ADDR]], align 42244// CHECK9-NEXT: store i32 [[TMP3]], ptr [[A_CASTED]], align 42245// CHECK9-NEXT: [[TMP4:%.*]] = load i64, ptr [[A_CASTED]], align 82246// CHECK9-NEXT: [[TMP5:%.*]] = load i16, ptr [[B_ADDR]], align 22247// CHECK9-NEXT: store i16 [[TMP5]], ptr [[B_CASTED]], align 22248// CHECK9-NEXT: [[TMP6:%.*]] = load i64, ptr [[B_CASTED]], align 82249// CHECK9-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB3]], i32 2, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l60.omp_outlined, i64 [[TMP4]], i64 [[TMP6]])2250// CHECK9-NEXT: ret void2251//2252//2253// CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l60.omp_outlined2254// CHECK9-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i64 noundef [[A:%.*]], i64 noundef [[B:%.*]]) #[[ATTR0]] {2255// CHECK9-NEXT: entry:2256// CHECK9-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 82257// CHECK9-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 82258// CHECK9-NEXT: [[A_ADDR:%.*]] = alloca i64, align 82259// CHECK9-NEXT: [[B_ADDR:%.*]] = alloca i64, align 82260// CHECK9-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 82261// CHECK9-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 82262// CHECK9-NEXT: store i64 [[A]], ptr [[A_ADDR]], align 82263// CHECK9-NEXT: store i64 [[B]], ptr [[B_ADDR]], align 82264// CHECK9-NEXT: [[TMP0:%.*]] = load i16, ptr [[B_ADDR]], align 22265// CHECK9-NEXT: [[CONV:%.*]] = sext i16 [[TMP0]] to i322266// CHECK9-NEXT: [[TMP1:%.*]] = load i32, ptr [[A_ADDR]], align 42267// CHECK9-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP1]], [[CONV]]2268// CHECK9-NEXT: store i32 [[ADD]], ptr [[A_ADDR]], align 42269// CHECK9-NEXT: ret void2270//2271//2272// CHECK11-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l712273// CHECK11-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], i32 noundef [[N:%.*]], i32 noundef [[DOTCAPTURE_EXPR_:%.*]], i32 noundef [[DOTCAPTURE_EXPR_1:%.*]]) #[[ATTR0:[0-9]+]] {2274// CHECK11-NEXT: entry:2275// CHECK11-NEXT: [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 42276// CHECK11-NEXT: [[N_ADDR:%.*]] = alloca i32, align 42277// CHECK11-NEXT: [[DOTCAPTURE_EXPR__ADDR:%.*]] = alloca i32, align 42278// CHECK11-NEXT: [[DOTCAPTURE_EXPR__ADDR2:%.*]] = alloca i32, align 42279// CHECK11-NEXT: [[N_CASTED:%.*]] = alloca i32, align 42280// CHECK11-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB3:[0-9]+]])2281// CHECK11-NEXT: store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 42282// CHECK11-NEXT: store i32 [[N]], ptr [[N_ADDR]], align 42283// CHECK11-NEXT: store i32 [[DOTCAPTURE_EXPR_]], ptr [[DOTCAPTURE_EXPR__ADDR]], align 42284// CHECK11-NEXT: store i32 [[DOTCAPTURE_EXPR_1]], ptr [[DOTCAPTURE_EXPR__ADDR2]], align 42285// CHECK11-NEXT: [[TMP1:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR__ADDR]], align 42286// CHECK11-NEXT: call void @__kmpc_push_num_teams(ptr @[[GLOB3]], i32 [[TMP0]], i32 [[TMP1]], i32 0)2287// CHECK11-NEXT: [[TMP2:%.*]] = load i32, ptr [[N_ADDR]], align 42288// CHECK11-NEXT: store i32 [[TMP2]], ptr [[N_CASTED]], align 42289// CHECK11-NEXT: [[TMP3:%.*]] = load i32, ptr [[N_CASTED]], align 42290// CHECK11-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB3]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l71.omp_outlined, i32 [[TMP3]])2291// CHECK11-NEXT: ret void2292//2293//2294// CHECK11-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l71.omp_outlined2295// CHECK11-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i32 noundef [[N:%.*]]) #[[ATTR0]] {2296// CHECK11-NEXT: entry:2297// CHECK11-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 42298// CHECK11-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 42299// CHECK11-NEXT: [[N_ADDR:%.*]] = alloca i32, align 42300// CHECK11-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 42301// CHECK11-NEXT: [[TMP:%.*]] = alloca i32, align 42302// CHECK11-NEXT: [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 42303// CHECK11-NEXT: [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 42304// CHECK11-NEXT: [[I:%.*]] = alloca i32, align 42305// CHECK11-NEXT: [[DOTOMP_COMB_LB:%.*]] = alloca i32, align 42306// CHECK11-NEXT: [[DOTOMP_COMB_UB:%.*]] = alloca i32, align 42307// CHECK11-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 42308// CHECK11-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 42309// CHECK11-NEXT: [[I3:%.*]] = alloca i32, align 42310// CHECK11-NEXT: [[N_CASTED:%.*]] = alloca i32, align 42311// CHECK11-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 42312// CHECK11-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 42313// CHECK11-NEXT: store i32 [[N]], ptr [[N_ADDR]], align 42314// CHECK11-NEXT: [[TMP0:%.*]] = load i32, ptr [[N_ADDR]], align 42315// CHECK11-NEXT: store i32 [[TMP0]], ptr [[DOTCAPTURE_EXPR_]], align 42316// CHECK11-NEXT: [[TMP1:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 42317// CHECK11-NEXT: [[SUB:%.*]] = sub nsw i32 [[TMP1]], 02318// CHECK11-NEXT: [[DIV:%.*]] = sdiv i32 [[SUB]], 12319// CHECK11-NEXT: [[SUB2:%.*]] = sub nsw i32 [[DIV]], 12320// CHECK11-NEXT: store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 42321// CHECK11-NEXT: store i32 0, ptr [[I]], align 42322// CHECK11-NEXT: [[TMP2:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 42323// CHECK11-NEXT: [[CMP:%.*]] = icmp slt i32 0, [[TMP2]]2324// CHECK11-NEXT: br i1 [[CMP]], label [[OMP_PRECOND_THEN:%.*]], label [[OMP_PRECOND_END:%.*]]2325// CHECK11: omp.precond.then:2326// CHECK11-NEXT: store i32 0, ptr [[DOTOMP_COMB_LB]], align 42327// CHECK11-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 42328// CHECK11-NEXT: store i32 [[TMP3]], ptr [[DOTOMP_COMB_UB]], align 42329// CHECK11-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 42330// CHECK11-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 42331// CHECK11-NEXT: [[TMP4:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 42332// CHECK11-NEXT: [[TMP5:%.*]] = load i32, ptr [[TMP4]], align 42333// CHECK11-NEXT: call void @__kmpc_for_static_init_4(ptr @[[GLOB1:[0-9]+]], i32 [[TMP5]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_COMB_LB]], ptr [[DOTOMP_COMB_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)2334// CHECK11-NEXT: [[TMP6:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 42335// CHECK11-NEXT: [[TMP7:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 42336// CHECK11-NEXT: [[CMP4:%.*]] = icmp sgt i32 [[TMP6]], [[TMP7]]2337// CHECK11-NEXT: br i1 [[CMP4]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]2338// CHECK11: cond.true:2339// CHECK11-NEXT: [[TMP8:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 42340// CHECK11-NEXT: br label [[COND_END:%.*]]2341// CHECK11: cond.false:2342// CHECK11-NEXT: [[TMP9:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 42343// CHECK11-NEXT: br label [[COND_END]]2344// CHECK11: cond.end:2345// CHECK11-NEXT: [[COND:%.*]] = phi i32 [ [[TMP8]], [[COND_TRUE]] ], [ [[TMP9]], [[COND_FALSE]] ]2346// CHECK11-NEXT: store i32 [[COND]], ptr [[DOTOMP_COMB_UB]], align 42347// CHECK11-NEXT: [[TMP10:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 42348// CHECK11-NEXT: store i32 [[TMP10]], ptr [[DOTOMP_IV]], align 42349// CHECK11-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]2350// CHECK11: omp.inner.for.cond:2351// CHECK11-NEXT: [[TMP11:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP17:![0-9]+]]2352// CHECK11-NEXT: [[TMP12:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4, !llvm.access.group [[ACC_GRP17]]2353// CHECK11-NEXT: [[CMP5:%.*]] = icmp sle i32 [[TMP11]], [[TMP12]]2354// CHECK11-NEXT: br i1 [[CMP5]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]2355// CHECK11: omp.inner.for.body:2356// CHECK11-NEXT: [[TMP13:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 4, !llvm.access.group [[ACC_GRP17]]2357// CHECK11-NEXT: [[TMP14:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4, !llvm.access.group [[ACC_GRP17]]2358// CHECK11-NEXT: [[TMP15:%.*]] = load i32, ptr [[N_ADDR]], align 4, !llvm.access.group [[ACC_GRP17]]2359// CHECK11-NEXT: store i32 [[TMP15]], ptr [[N_CASTED]], align 4, !llvm.access.group [[ACC_GRP17]]2360// CHECK11-NEXT: [[TMP16:%.*]] = load i32, ptr [[N_CASTED]], align 4, !llvm.access.group [[ACC_GRP17]]2361// CHECK11-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB3]], i32 3, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l71.omp_outlined.omp_outlined, i32 [[TMP13]], i32 [[TMP14]], i32 [[TMP16]]), !llvm.access.group [[ACC_GRP17]]2362// CHECK11-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]2363// CHECK11: omp.inner.for.inc:2364// CHECK11-NEXT: [[TMP17:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP17]]2365// CHECK11-NEXT: [[TMP18:%.*]] = load i32, ptr [[DOTOMP_STRIDE]], align 4, !llvm.access.group [[ACC_GRP17]]2366// CHECK11-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP17]], [[TMP18]]2367// CHECK11-NEXT: store i32 [[ADD]], ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP17]]2368// CHECK11-NEXT: br label [[OMP_INNER_FOR_COND]], !llvm.loop [[LOOP18:![0-9]+]]2369// CHECK11: omp.inner.for.end:2370// CHECK11-NEXT: br label [[OMP_LOOP_EXIT:%.*]]2371// CHECK11: omp.loop.exit:2372// CHECK11-NEXT: [[TMP19:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 42373// CHECK11-NEXT: [[TMP20:%.*]] = load i32, ptr [[TMP19]], align 42374// CHECK11-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP20]])2375// CHECK11-NEXT: [[TMP21:%.*]] = load i32, ptr [[DOTOMP_IS_LAST]], align 42376// CHECK11-NEXT: [[TMP22:%.*]] = icmp ne i32 [[TMP21]], 02377// CHECK11-NEXT: br i1 [[TMP22]], label [[DOTOMP_FINAL_THEN:%.*]], label [[DOTOMP_FINAL_DONE:%.*]]2378// CHECK11: .omp.final.then:2379// CHECK11-NEXT: [[TMP23:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 42380// CHECK11-NEXT: [[SUB6:%.*]] = sub nsw i32 [[TMP23]], 02381// CHECK11-NEXT: [[DIV7:%.*]] = sdiv i32 [[SUB6]], 12382// CHECK11-NEXT: [[MUL:%.*]] = mul nsw i32 [[DIV7]], 12383// CHECK11-NEXT: [[ADD8:%.*]] = add nsw i32 0, [[MUL]]2384// CHECK11-NEXT: store i32 [[ADD8]], ptr [[I3]], align 42385// CHECK11-NEXT: br label [[DOTOMP_FINAL_DONE]]2386// CHECK11: .omp.final.done:2387// CHECK11-NEXT: br label [[OMP_PRECOND_END]]2388// CHECK11: omp.precond.end:2389// CHECK11-NEXT: ret void2390//2391//2392// CHECK11-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l71.omp_outlined.omp_outlined2393// CHECK11-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i32 noundef [[DOTPREVIOUS_LB_:%.*]], i32 noundef [[DOTPREVIOUS_UB_:%.*]], i32 noundef [[N:%.*]]) #[[ATTR0]] {2394// CHECK11-NEXT: entry:2395// CHECK11-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 42396// CHECK11-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 42397// CHECK11-NEXT: [[DOTPREVIOUS_LB__ADDR:%.*]] = alloca i32, align 42398// CHECK11-NEXT: [[DOTPREVIOUS_UB__ADDR:%.*]] = alloca i32, align 42399// CHECK11-NEXT: [[N_ADDR:%.*]] = alloca i32, align 42400// CHECK11-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 42401// CHECK11-NEXT: [[TMP:%.*]] = alloca i32, align 42402// CHECK11-NEXT: [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 42403// CHECK11-NEXT: [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 42404// CHECK11-NEXT: [[I:%.*]] = alloca i32, align 42405// CHECK11-NEXT: [[DOTOMP_LB:%.*]] = alloca i32, align 42406// CHECK11-NEXT: [[DOTOMP_UB:%.*]] = alloca i32, align 42407// CHECK11-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 42408// CHECK11-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 42409// CHECK11-NEXT: [[I3:%.*]] = alloca i32, align 42410// CHECK11-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 42411// CHECK11-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 42412// CHECK11-NEXT: store i32 [[DOTPREVIOUS_LB_]], ptr [[DOTPREVIOUS_LB__ADDR]], align 42413// CHECK11-NEXT: store i32 [[DOTPREVIOUS_UB_]], ptr [[DOTPREVIOUS_UB__ADDR]], align 42414// CHECK11-NEXT: store i32 [[N]], ptr [[N_ADDR]], align 42415// CHECK11-NEXT: [[TMP0:%.*]] = load i32, ptr [[N_ADDR]], align 42416// CHECK11-NEXT: store i32 [[TMP0]], ptr [[DOTCAPTURE_EXPR_]], align 42417// CHECK11-NEXT: [[TMP1:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 42418// CHECK11-NEXT: [[SUB:%.*]] = sub nsw i32 [[TMP1]], 02419// CHECK11-NEXT: [[DIV:%.*]] = sdiv i32 [[SUB]], 12420// CHECK11-NEXT: [[SUB2:%.*]] = sub nsw i32 [[DIV]], 12421// CHECK11-NEXT: store i32 [[SUB2]], ptr [[DOTCAPTURE_EXPR_1]], align 42422// CHECK11-NEXT: store i32 0, ptr [[I]], align 42423// CHECK11-NEXT: [[TMP2:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 42424// CHECK11-NEXT: [[CMP:%.*]] = icmp slt i32 0, [[TMP2]]2425// CHECK11-NEXT: br i1 [[CMP]], label [[OMP_PRECOND_THEN:%.*]], label [[OMP_PRECOND_END:%.*]]2426// CHECK11: omp.precond.then:2427// CHECK11-NEXT: store i32 0, ptr [[DOTOMP_LB]], align 42428// CHECK11-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 42429// CHECK11-NEXT: store i32 [[TMP3]], ptr [[DOTOMP_UB]], align 42430// CHECK11-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTPREVIOUS_LB__ADDR]], align 42431// CHECK11-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTPREVIOUS_UB__ADDR]], align 42432// CHECK11-NEXT: store i32 [[TMP4]], ptr [[DOTOMP_LB]], align 42433// CHECK11-NEXT: store i32 [[TMP5]], ptr [[DOTOMP_UB]], align 42434// CHECK11-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 42435// CHECK11-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 42436// CHECK11-NEXT: [[TMP6:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 42437// CHECK11-NEXT: [[TMP7:%.*]] = load i32, ptr [[TMP6]], align 42438// CHECK11-NEXT: call void @__kmpc_for_static_init_4(ptr @[[GLOB2:[0-9]+]], i32 [[TMP7]], i32 34, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)2439// CHECK11-NEXT: [[TMP8:%.*]] = load i32, ptr [[DOTOMP_UB]], align 42440// CHECK11-NEXT: [[TMP9:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 42441// CHECK11-NEXT: [[CMP4:%.*]] = icmp sgt i32 [[TMP8]], [[TMP9]]2442// CHECK11-NEXT: br i1 [[CMP4]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]2443// CHECK11: cond.true:2444// CHECK11-NEXT: [[TMP10:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 42445// CHECK11-NEXT: br label [[COND_END:%.*]]2446// CHECK11: cond.false:2447// CHECK11-NEXT: [[TMP11:%.*]] = load i32, ptr [[DOTOMP_UB]], align 42448// CHECK11-NEXT: br label [[COND_END]]2449// CHECK11: cond.end:2450// CHECK11-NEXT: [[COND:%.*]] = phi i32 [ [[TMP10]], [[COND_TRUE]] ], [ [[TMP11]], [[COND_FALSE]] ]2451// CHECK11-NEXT: store i32 [[COND]], ptr [[DOTOMP_UB]], align 42452// CHECK11-NEXT: [[TMP12:%.*]] = load i32, ptr [[DOTOMP_LB]], align 42453// CHECK11-NEXT: store i32 [[TMP12]], ptr [[DOTOMP_IV]], align 42454// CHECK11-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]2455// CHECK11: omp.inner.for.cond:2456// CHECK11-NEXT: [[TMP13:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP21:![0-9]+]]2457// CHECK11-NEXT: [[TMP14:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4, !llvm.access.group [[ACC_GRP21]]2458// CHECK11-NEXT: [[CMP5:%.*]] = icmp sle i32 [[TMP13]], [[TMP14]]2459// CHECK11-NEXT: br i1 [[CMP5]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]2460// CHECK11: omp.inner.for.body:2461// CHECK11-NEXT: [[TMP15:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP21]]2462// CHECK11-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP15]], 12463// CHECK11-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]2464// CHECK11-NEXT: store i32 [[ADD]], ptr [[I3]], align 4, !llvm.access.group [[ACC_GRP21]]2465// CHECK11-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]2466// CHECK11: omp.body.continue:2467// CHECK11-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]2468// CHECK11: omp.inner.for.inc:2469// CHECK11-NEXT: [[TMP16:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP21]]2470// CHECK11-NEXT: [[ADD6:%.*]] = add nsw i32 [[TMP16]], 12471// CHECK11-NEXT: store i32 [[ADD6]], ptr [[DOTOMP_IV]], align 4, !llvm.access.group [[ACC_GRP21]]2472// CHECK11-NEXT: br label [[OMP_INNER_FOR_COND]], !llvm.loop [[LOOP22:![0-9]+]]2473// CHECK11: omp.inner.for.end:2474// CHECK11-NEXT: br label [[OMP_LOOP_EXIT:%.*]]2475// CHECK11: omp.loop.exit:2476// CHECK11-NEXT: [[TMP17:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 42477// CHECK11-NEXT: [[TMP18:%.*]] = load i32, ptr [[TMP17]], align 42478// CHECK11-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB2]], i32 [[TMP18]])2479// CHECK11-NEXT: [[TMP19:%.*]] = load i32, ptr [[DOTOMP_IS_LAST]], align 42480// CHECK11-NEXT: [[TMP20:%.*]] = icmp ne i32 [[TMP19]], 02481// CHECK11-NEXT: br i1 [[TMP20]], label [[DOTOMP_FINAL_THEN:%.*]], label [[DOTOMP_FINAL_DONE:%.*]]2482// CHECK11: .omp.final.then:2483// CHECK11-NEXT: [[TMP21:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 42484// CHECK11-NEXT: [[SUB7:%.*]] = sub nsw i32 [[TMP21]], 02485// CHECK11-NEXT: [[DIV8:%.*]] = sdiv i32 [[SUB7]], 12486// CHECK11-NEXT: [[MUL9:%.*]] = mul nsw i32 [[DIV8]], 12487// CHECK11-NEXT: [[ADD10:%.*]] = add nsw i32 0, [[MUL9]]2488// CHECK11-NEXT: store i32 [[ADD10]], ptr [[I3]], align 42489// CHECK11-NEXT: br label [[DOTOMP_FINAL_DONE]]2490// CHECK11: .omp.final.done:2491// CHECK11-NEXT: br label [[OMP_PRECOND_END]]2492// CHECK11: omp.precond.end:2493// CHECK11-NEXT: ret void2494//2495//2496// CHECK11-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l752497// CHECK11-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], i32 noundef [[DOTCAPTURE_EXPR_:%.*]]) #[[ATTR0]] {2498// CHECK11-NEXT: entry:2499// CHECK11-NEXT: [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 42500// CHECK11-NEXT: [[DOTCAPTURE_EXPR__ADDR:%.*]] = alloca i32, align 42501// CHECK11-NEXT: store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 42502// CHECK11-NEXT: store i32 [[DOTCAPTURE_EXPR_]], ptr [[DOTCAPTURE_EXPR__ADDR]], align 42503// CHECK11-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB3]], i32 0, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l75.omp_outlined)2504// CHECK11-NEXT: ret void2505//2506//2507// CHECK11-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZL7fstatici_l75.omp_outlined2508// CHECK11-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR0]] {2509// CHECK11-NEXT: entry:2510// CHECK11-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 42511// CHECK11-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 42512// CHECK11-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 42513// CHECK11-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 42514// CHECK11-NEXT: ret void2515//2516//2517// CHECK11-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2S12r1Ei_l882518// CHECK11-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], ptr noundef [[THIS:%.*]], i32 noundef [[B:%.*]], i32 noundef [[DOTCAPTURE_EXPR_:%.*]]) #[[ATTR0]] {2519// CHECK11-NEXT: entry:2520// CHECK11-NEXT: [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 42521// CHECK11-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 42522// CHECK11-NEXT: [[B_ADDR:%.*]] = alloca i32, align 42523// CHECK11-NEXT: [[DOTCAPTURE_EXPR__ADDR:%.*]] = alloca i32, align 42524// CHECK11-NEXT: [[B_CASTED:%.*]] = alloca i32, align 42525// CHECK11-NEXT: store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 42526// CHECK11-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 42527// CHECK11-NEXT: store i32 [[B]], ptr [[B_ADDR]], align 42528// CHECK11-NEXT: store i32 [[DOTCAPTURE_EXPR_]], ptr [[DOTCAPTURE_EXPR__ADDR]], align 42529// CHECK11-NEXT: [[TMP0:%.*]] = load ptr, ptr [[THIS_ADDR]], align 42530// CHECK11-NEXT: [[TMP1:%.*]] = load i32, ptr [[B_ADDR]], align 42531// CHECK11-NEXT: store i32 [[TMP1]], ptr [[B_CASTED]], align 42532// CHECK11-NEXT: [[TMP2:%.*]] = load i32, ptr [[B_CASTED]], align 42533// CHECK11-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB3]], i32 2, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2S12r1Ei_l88.omp_outlined, ptr [[TMP0]], i32 [[TMP2]])2534// CHECK11-NEXT: ret void2535//2536//2537// CHECK11-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2S12r1Ei_l88.omp_outlined2538// CHECK11-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef [[THIS:%.*]], i32 noundef [[B:%.*]]) #[[ATTR0]] {2539// CHECK11-NEXT: entry:2540// CHECK11-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 42541// CHECK11-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 42542// CHECK11-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 42543// CHECK11-NEXT: [[B_ADDR:%.*]] = alloca i32, align 42544// CHECK11-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 42545// CHECK11-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 42546// CHECK11-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 42547// CHECK11-NEXT: store i32 [[B]], ptr [[B_ADDR]], align 42548// CHECK11-NEXT: [[TMP0:%.*]] = load ptr, ptr [[THIS_ADDR]], align 42549// CHECK11-NEXT: [[TMP1:%.*]] = load i32, ptr [[B_ADDR]], align 42550// CHECK11-NEXT: [[CONV:%.*]] = sitofp i32 [[TMP1]] to double2551// CHECK11-NEXT: [[ADD:%.*]] = fadd double [[CONV]], 1.500000e+002552// CHECK11-NEXT: [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_S1:%.*]], ptr [[TMP0]], i32 0, i32 02553// CHECK11-NEXT: store double [[ADD]], ptr [[A]], align 42554// CHECK11-NEXT: ret void2555//2556//2557// CHECK11-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2S12r1Ei_l932558// CHECK11-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], ptr noundef [[THIS:%.*]]) #[[ATTR0]] {2559// CHECK11-NEXT: entry:2560// CHECK11-NEXT: [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 42561// CHECK11-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 42562// CHECK11-NEXT: store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 42563// CHECK11-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 42564// CHECK11-NEXT: [[TMP0:%.*]] = load ptr, ptr [[THIS_ADDR]], align 42565// CHECK11-NEXT: [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_S1:%.*]], ptr [[TMP0]], i32 0, i32 02566// CHECK11-NEXT: store double 2.500000e+00, ptr [[A]], align 42567// CHECK11-NEXT: ret void2568//2569//2570// CHECK11-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l552571// CHECK11-SAME: (ptr noalias noundef [[DYN_PTR:%.*]]) #[[ATTR0]] {2572// CHECK11-NEXT: entry:2573// CHECK11-NEXT: [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 42574// CHECK11-NEXT: store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 42575// CHECK11-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB3]], i32 0, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l55.omp_outlined)2576// CHECK11-NEXT: ret void2577//2578//2579// CHECK11-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l55.omp_outlined2580// CHECK11-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR0]] {2581// CHECK11-NEXT: entry:2582// CHECK11-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 42583// CHECK11-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 42584// CHECK11-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 42585// CHECK11-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 42586// CHECK11-NEXT: ret void2587//2588//2589// CHECK11-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l602590// CHECK11-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], i32 noundef [[A:%.*]], i32 noundef [[B:%.*]], i32 noundef [[DOTCAPTURE_EXPR_:%.*]]) #[[ATTR0]] {2591// CHECK11-NEXT: entry:2592// CHECK11-NEXT: [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 42593// CHECK11-NEXT: [[A_ADDR:%.*]] = alloca i32, align 42594// CHECK11-NEXT: [[B_ADDR:%.*]] = alloca i32, align 42595// CHECK11-NEXT: [[DOTCAPTURE_EXPR__ADDR:%.*]] = alloca i32, align 42596// CHECK11-NEXT: [[A_CASTED:%.*]] = alloca i32, align 42597// CHECK11-NEXT: [[B_CASTED:%.*]] = alloca i32, align 42598// CHECK11-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB3]])2599// CHECK11-NEXT: store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 42600// CHECK11-NEXT: store i32 [[A]], ptr [[A_ADDR]], align 42601// CHECK11-NEXT: store i32 [[B]], ptr [[B_ADDR]], align 42602// CHECK11-NEXT: store i32 [[DOTCAPTURE_EXPR_]], ptr [[DOTCAPTURE_EXPR__ADDR]], align 42603// CHECK11-NEXT: [[TMP1:%.*]] = load i16, ptr [[DOTCAPTURE_EXPR__ADDR]], align 22604// CHECK11-NEXT: [[TMP2:%.*]] = sext i16 [[TMP1]] to i322605// CHECK11-NEXT: call void @__kmpc_push_num_teams(ptr @[[GLOB3]], i32 [[TMP0]], i32 [[TMP2]], i32 0)2606// CHECK11-NEXT: [[TMP3:%.*]] = load i32, ptr [[A_ADDR]], align 42607// CHECK11-NEXT: store i32 [[TMP3]], ptr [[A_CASTED]], align 42608// CHECK11-NEXT: [[TMP4:%.*]] = load i32, ptr [[A_CASTED]], align 42609// CHECK11-NEXT: [[TMP5:%.*]] = load i16, ptr [[B_ADDR]], align 22610// CHECK11-NEXT: store i16 [[TMP5]], ptr [[B_CASTED]], align 22611// CHECK11-NEXT: [[TMP6:%.*]] = load i32, ptr [[B_CASTED]], align 42612// CHECK11-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB3]], i32 2, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l60.omp_outlined, i32 [[TMP4]], i32 [[TMP6]])2613// CHECK11-NEXT: ret void2614//2615//2616// CHECK11-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l60.omp_outlined2617// CHECK11-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i32 noundef [[A:%.*]], i32 noundef [[B:%.*]]) #[[ATTR0]] {2618// CHECK11-NEXT: entry:2619// CHECK11-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 42620// CHECK11-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 42621// CHECK11-NEXT: [[A_ADDR:%.*]] = alloca i32, align 42622// CHECK11-NEXT: [[B_ADDR:%.*]] = alloca i32, align 42623// CHECK11-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 42624// CHECK11-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 42625// CHECK11-NEXT: store i32 [[A]], ptr [[A_ADDR]], align 42626// CHECK11-NEXT: store i32 [[B]], ptr [[B_ADDR]], align 42627// CHECK11-NEXT: [[TMP0:%.*]] = load i16, ptr [[B_ADDR]], align 22628// CHECK11-NEXT: [[CONV:%.*]] = sext i16 [[TMP0]] to i322629// CHECK11-NEXT: [[TMP1:%.*]] = load i32, ptr [[A_ADDR]], align 42630// CHECK11-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP1]], [[CONV]]2631// CHECK11-NEXT: store i32 [[ADD]], ptr [[A_ADDR]], align 42632// CHECK11-NEXT: ret void2633//2634