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 // expected-no-diagnostics 3 #ifndef HEADER 4 #define HEADER 5 // Test host codegen. 6 // RUN: %clang_cc1 -no-opaque-pointers -DCK1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix=CHECK1 7 // RUN: %clang_cc1 -no-opaque-pointers -DCK1 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s 8 // RUN: %clang_cc1 -no-opaque-pointers -DCK1 -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix=CHECK1 9 // RUN: %clang_cc1 -no-opaque-pointers -DCK1 -verify -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix=CHECK3 10 // RUN: %clang_cc1 -no-opaque-pointers -DCK1 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s 11 // RUN: %clang_cc1 -no-opaque-pointers -DCK1 -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix=CHECK3 12 13 // RUN: %clang_cc1 -no-opaque-pointers -DCK1 -verify -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}" 14 // RUN: %clang_cc1 -no-opaque-pointers -DCK1 -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s 15 // RUN: %clang_cc1 -no-opaque-pointers -DCK1 -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}" 16 // RUN: %clang_cc1 -no-opaque-pointers -DCK1 -verify -fopenmp-simd -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}" 17 // RUN: %clang_cc1 -no-opaque-pointers -DCK1 -fopenmp-simd -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s 18 // RUN: %clang_cc1 -no-opaque-pointers -DCK1 -fopenmp-simd -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}" 19 #ifdef CK1 20 21 int Gbla; 22 long long Gblb; 23 int &Gblc = Gbla; 24 25 int teams_argument_global_local(int a){ 26 int comp = 1; 27 28 int la = 23; 29 float lc = 25.0; 30 31 #pragma omp target 32 #pragma omp teams 33 { 34 ++comp; 35 } 36 37 #pragma omp target 38 {{{ 39 #pragma omp teams 40 { 41 ++comp; 42 } 43 }}} 44 45 46 #pragma omp target 47 #pragma omp teams num_teams(la) 48 { 49 ++comp; 50 } 51 52 53 #pragma omp target 54 #pragma omp teams thread_limit(la) 55 { 56 ++comp; 57 } 58 59 60 61 62 #pragma omp target 63 #pragma omp teams num_teams(Gbla+a) thread_limit(Gblb+(long long)lc) 64 { 65 ++comp; 66 } 67 68 69 70 71 #pragma omp target 72 #pragma omp teams num_teams(Gblc+1) thread_limit(Gblc+2) 73 { 74 comp += Gblc; 75 } 76 77 return comp; 78 } 79 80 #endif // CK1 81 82 // Test host codegen. 83 // RUN: %clang_cc1 -no-opaque-pointers -DCK2 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix=CHECK9 84 // RUN: %clang_cc1 -no-opaque-pointers -DCK2 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s 85 // RUN: %clang_cc1 -no-opaque-pointers -DCK2 -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix=CHECK9 86 // RUN: %clang_cc1 -no-opaque-pointers -DCK2 -verify -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix=CHECK11 87 // RUN: %clang_cc1 -no-opaque-pointers -DCK2 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s 88 // RUN: %clang_cc1 -no-opaque-pointers -DCK2 -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix=CHECK11 89 90 // RUN: %clang_cc1 -no-opaque-pointers -DCK2 -verify -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}" 91 // RUN: %clang_cc1 -no-opaque-pointers -DCK2 -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s 92 // RUN: %clang_cc1 -no-opaque-pointers -DCK2 -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}" 93 // RUN: %clang_cc1 -no-opaque-pointers -DCK2 -verify -fopenmp-simd -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}" 94 // RUN: %clang_cc1 -no-opaque-pointers -DCK2 -fopenmp-simd -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s 95 // RUN: %clang_cc1 -no-opaque-pointers -DCK2 -fopenmp-simd -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}" 96 #ifdef CK2 97 98 template <typename T> 99 struct SS{ 100 T a; 101 float b; 102 }; 103 104 SS<int> Gbla; 105 SS<long long> Gblb; 106 107 int teams_template_arg(void) { 108 int comp = 1; 109 110 SS<int> la; 111 SS<long long> lb; 112 113 114 115 116 #pragma omp target 117 #pragma omp teams num_teams(Gbla.a) thread_limit((long long)la.b) 118 { 119 ++comp; 120 } 121 122 123 124 125 #pragma omp target 126 #pragma omp teams num_teams((long long)lb.b) thread_limit(Gblb.a) 127 { 128 ++comp; 129 } 130 return comp; 131 } 132 #endif // CK2 133 134 // Test host codegen. 135 // RUN: %clang_cc1 -no-opaque-pointers -DCK3 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix=CHECK17 136 // RUN: %clang_cc1 -no-opaque-pointers -DCK3 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s 137 // RUN: %clang_cc1 -no-opaque-pointers -DCK3 -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix=CHECK17 138 // RUN: %clang_cc1 -no-opaque-pointers -DCK3 -verify -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix=CHECK19 139 // RUN: %clang_cc1 -no-opaque-pointers -DCK3 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s 140 // RUN: %clang_cc1 -no-opaque-pointers -DCK3 -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix=CHECK19 141 142 // RUN: %clang_cc1 -no-opaque-pointers -DCK3 -verify -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}" 143 // RUN: %clang_cc1 -no-opaque-pointers -DCK3 -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s 144 // RUN: %clang_cc1 -no-opaque-pointers -DCK3 -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}" 145 // RUN: %clang_cc1 -no-opaque-pointers -DCK3 -verify -fopenmp-simd -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}" 146 // RUN: %clang_cc1 -no-opaque-pointers -DCK3 -fopenmp-simd -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s 147 // RUN: %clang_cc1 -no-opaque-pointers -DCK3 -fopenmp-simd -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}" 148 #ifdef CK3 149 150 151 template <typename T, int X, long long Y> 152 struct SS{ 153 T a; 154 float b; 155 156 int foo(void) { 157 int comp = 1; 158 159 160 161 #pragma omp target 162 #pragma omp teams num_teams(a) thread_limit(X) 163 { 164 ++comp; 165 } 166 167 168 169 #pragma omp target 170 #pragma omp teams num_teams(Y) thread_limit((int)b+X) 171 { 172 ++comp; 173 } 174 return comp; 175 } 176 }; 177 178 int teams_template_struct(void) { 179 SS<int, 123, 456> V; 180 return V.foo(); 181 182 } 183 #endif // CK3 184 185 // Test target codegen - host bc file has to be created first. 186 // RUN: %clang_cc1 -no-opaque-pointers -DCK4 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm-bc %s -o %t-ppc-host.bc 187 // RUN: %clang_cc1 -no-opaque-pointers -DCK4 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --check-prefix=CHECK25 188 // RUN: %clang_cc1 -no-opaque-pointers -DCK4 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o %t %s 189 // RUN: %clang_cc1 -no-opaque-pointers -DCK4 -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix=CHECK25 190 // RUN: %clang_cc1 -no-opaque-pointers -DCK4 -verify -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm-bc %s -o %t-x86-host.bc 191 // RUN: %clang_cc1 -no-opaque-pointers -DCK4 -verify -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s --check-prefix=CHECK27 192 // RUN: %clang_cc1 -no-opaque-pointers -DCK4 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o %t %s 193 // RUN: %clang_cc1 -no-opaque-pointers -DCK4 -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix=CHECK27 194 195 // RUN: %clang_cc1 -no-opaque-pointers -DCK4 -verify -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm-bc %s -o %t-ppc-host.bc 196 // RUN: %clang_cc1 -no-opaque-pointers -DCK4 -verify -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}" 197 // RUN: %clang_cc1 -no-opaque-pointers -DCK4 -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o %t %s 198 // RUN: %clang_cc1 -no-opaque-pointers -DCK4 -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -fopenmp-is-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}}" 199 // RUN: %clang_cc1 -no-opaque-pointers -DCK4 -verify -fopenmp-simd -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm-bc %s -o %t-x86-host.bc 200 // RUN: %clang_cc1 -no-opaque-pointers -DCK4 -verify -fopenmp-simd -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}" 201 // RUN: %clang_cc1 -no-opaque-pointers -DCK4 -fopenmp-simd -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o %t %s 202 // RUN: %clang_cc1 -no-opaque-pointers -DCK4 -fopenmp-simd -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -fopenmp-is-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}}" 203 204 #ifdef CK4 205 206 207 template <typename T> 208 int tmain(T argc) { 209 #pragma omp target 210 #pragma omp teams 211 argc = 0; 212 return 0; 213 } 214 215 int main (int argc, char **argv) { 216 #pragma omp target 217 #pragma omp teams 218 argc = 0; 219 return tmain(argv); 220 } 221 222 223 224 225 #endif // CK4 226 227 // Test target codegen - host bc file has to be created first. 228 // RUN: %clang_cc1 -no-opaque-pointers -DCK5 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm-bc %s -o %t-ppc-host.bc 229 // RUN: %clang_cc1 -no-opaque-pointers -DCK5 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --check-prefix=CHECK33 230 // RUN: %clang_cc1 -no-opaque-pointers -DCK5 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o %t %s 231 // RUN: %clang_cc1 -no-opaque-pointers -DCK5 -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix=CHECK33 232 // RUN: %clang_cc1 -no-opaque-pointers -DCK5 -verify -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm-bc %s -o %t-x86-host.bc 233 // RUN: %clang_cc1 -no-opaque-pointers -DCK5 -verify -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s --check-prefix=CHECK35 234 // RUN: %clang_cc1 -no-opaque-pointers -DCK5 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o %t %s 235 // RUN: %clang_cc1 -no-opaque-pointers -DCK5 -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix=CHECK35 236 237 // RUN: %clang_cc1 -no-opaque-pointers -DCK5 -verify -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm-bc %s -o %t-ppc-host.bc 238 // RUN: %clang_cc1 -no-opaque-pointers -DCK5 -verify -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}" 239 // RUN: %clang_cc1 -no-opaque-pointers -DCK5 -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o %t %s 240 // RUN: %clang_cc1 -no-opaque-pointers -DCK5 -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -fopenmp-is-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}}" 241 // RUN: %clang_cc1 -no-opaque-pointers -DCK5 -verify -fopenmp-simd -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm-bc %s -o %t-x86-host.bc 242 // RUN: %clang_cc1 -no-opaque-pointers -DCK5 -verify -fopenmp-simd -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}" 243 // RUN: %clang_cc1 -no-opaque-pointers -DCK5 -fopenmp-simd -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o %t %s 244 // RUN: %clang_cc1 -no-opaque-pointers -DCK5 -fopenmp-simd -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -fopenmp-is-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}}" 245 246 // expected-no-diagnostics 247 #ifdef CK5 248 249 250 template <typename T> 251 int tmain(T argc) { 252 int a = 10; 253 int b = 5; 254 #pragma omp target 255 #pragma omp teams num_teams(a) thread_limit(b) 256 { 257 argc = 0; 258 } 259 return 0; 260 } 261 262 int main (int argc, char **argv) { 263 int a = 20; 264 int b = 5; 265 #pragma omp target 266 #pragma omp teams num_teams(a) thread_limit(b) 267 { 268 argc = 0; 269 } 270 return tmain(argv); 271 } 272 273 274 275 #endif // CK5 276 277 // Test host codegen. 278 // RUN: %clang_cc1 -no-opaque-pointers -DCK6 -verify -fopenmp -fopenmp-version=50 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix=CHECK41 279 // RUN: %clang_cc1 -no-opaque-pointers -DCK6 -fopenmp -fopenmp-version=50 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s 280 // RUN: %clang_cc1 -no-opaque-pointers -DCK6 -fopenmp-version=50 -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix=CHECK41 281 // RUN: %clang_cc1 -no-opaque-pointers -DCK6 -verify -fopenmp -fopenmp-version=50 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix=CHECK43 282 // RUN: %clang_cc1 -no-opaque-pointers -DCK6 -fopenmp -fopenmp-version=50 -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s 283 // RUN: %clang_cc1 -no-opaque-pointers -DCK6 -fopenmp -fopenmp-version=50 -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=CHECK43 284 285 // RUN: %clang_cc1 -no-opaque-pointers -DCK6 -verify -fopenmp-version=50 -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}" 286 // RUN: %clang_cc1 -no-opaque-pointers -DCK6 -fopenmp-simd -fopenmp-version=50 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s 287 // RUN: %clang_cc1 -no-opaque-pointers -DCK6 -fopenmp-simd -fopenmp-version=50 -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}}" 288 // RUN: %clang_cc1 -no-opaque-pointers -DCK6 -verify -fopenmp-simd -fopenmp-version=50 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}" 289 // RUN: %clang_cc1 -no-opaque-pointers -DCK6 -fopenmp-simd -fopenmp-version=50 -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s 290 // RUN: %clang_cc1 -no-opaque-pointers -DCK6 -fopenmp-simd -fopenmp-version=50 -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}}" 291 #ifdef CK6 292 293 void foo() { 294 #pragma omp teams 295 ; 296 } 297 298 #endif // CK6 299 300 #endif 301 // CHECK1-LABEL: define {{[^@]+}}@_Z27teams_argument_global_locali 302 // CHECK1-SAME: (i32 noundef signext [[A:%.*]]) #[[ATTR0:[0-9]+]] { 303 // CHECK1-NEXT: entry: 304 // CHECK1-NEXT: [[A_ADDR:%.*]] = alloca i32, align 4 305 // CHECK1-NEXT: [[COMP:%.*]] = alloca i32, align 4 306 // CHECK1-NEXT: [[LA:%.*]] = alloca i32, align 4 307 // CHECK1-NEXT: [[LC:%.*]] = alloca float, align 4 308 // CHECK1-NEXT: [[COMP_CASTED:%.*]] = alloca i64, align 8 309 // CHECK1-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x i8*], align 8 310 // CHECK1-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x i8*], align 8 311 // CHECK1-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [1 x i8*], align 8 312 // CHECK1-NEXT: [[COMP_CASTED1:%.*]] = alloca i64, align 8 313 // CHECK1-NEXT: [[DOTOFFLOAD_BASEPTRS3:%.*]] = alloca [1 x i8*], align 8 314 // CHECK1-NEXT: [[DOTOFFLOAD_PTRS4:%.*]] = alloca [1 x i8*], align 8 315 // CHECK1-NEXT: [[DOTOFFLOAD_MAPPERS5:%.*]] = alloca [1 x i8*], align 8 316 // CHECK1-NEXT: [[LA_CASTED:%.*]] = alloca i64, align 8 317 // CHECK1-NEXT: [[COMP_CASTED9:%.*]] = alloca i64, align 8 318 // CHECK1-NEXT: [[DOTOFFLOAD_BASEPTRS11:%.*]] = alloca [2 x i8*], align 8 319 // CHECK1-NEXT: [[DOTOFFLOAD_PTRS12:%.*]] = alloca [2 x i8*], align 8 320 // CHECK1-NEXT: [[DOTOFFLOAD_MAPPERS13:%.*]] = alloca [2 x i8*], align 8 321 // CHECK1-NEXT: [[LA_CASTED16:%.*]] = alloca i64, align 8 322 // CHECK1-NEXT: [[COMP_CASTED18:%.*]] = alloca i64, align 8 323 // CHECK1-NEXT: [[DOTOFFLOAD_BASEPTRS20:%.*]] = alloca [2 x i8*], align 8 324 // CHECK1-NEXT: [[DOTOFFLOAD_PTRS21:%.*]] = alloca [2 x i8*], align 8 325 // CHECK1-NEXT: [[DOTOFFLOAD_MAPPERS22:%.*]] = alloca [2 x i8*], align 8 326 // CHECK1-NEXT: [[GBLA_CASTED:%.*]] = alloca i64, align 8 327 // CHECK1-NEXT: [[A_CASTED:%.*]] = alloca i64, align 8 328 // CHECK1-NEXT: [[GBLB_CASTED:%.*]] = alloca i64, align 8 329 // CHECK1-NEXT: [[LC_CASTED:%.*]] = alloca i64, align 8 330 // CHECK1-NEXT: [[COMP_CASTED28:%.*]] = alloca i64, align 8 331 // CHECK1-NEXT: [[DOTOFFLOAD_BASEPTRS30:%.*]] = alloca [5 x i8*], align 8 332 // CHECK1-NEXT: [[DOTOFFLOAD_PTRS31:%.*]] = alloca [5 x i8*], align 8 333 // CHECK1-NEXT: [[DOTOFFLOAD_MAPPERS32:%.*]] = alloca [5 x i8*], align 8 334 // CHECK1-NEXT: [[TMP:%.*]] = alloca i32*, align 8 335 // CHECK1-NEXT: [[GBLC_CASTED:%.*]] = alloca i64, align 8 336 // CHECK1-NEXT: [[COMP_CASTED38:%.*]] = alloca i64, align 8 337 // CHECK1-NEXT: [[DOTOFFLOAD_BASEPTRS40:%.*]] = alloca [2 x i8*], align 8 338 // CHECK1-NEXT: [[DOTOFFLOAD_PTRS41:%.*]] = alloca [2 x i8*], align 8 339 // CHECK1-NEXT: [[DOTOFFLOAD_MAPPERS42:%.*]] = alloca [2 x i8*], align 8 340 // CHECK1-NEXT: [[_TMP43:%.*]] = alloca i32*, align 8 341 // CHECK1-NEXT: [[_TMP45:%.*]] = alloca i32*, align 8 342 // CHECK1-NEXT: store i32 [[A]], i32* [[A_ADDR]], align 4 343 // CHECK1-NEXT: store i32 1, i32* [[COMP]], align 4 344 // CHECK1-NEXT: store i32 23, i32* [[LA]], align 4 345 // CHECK1-NEXT: store float 2.500000e+01, float* [[LC]], align 4 346 // CHECK1-NEXT: [[TMP0:%.*]] = load i32, i32* [[COMP]], align 4 347 // CHECK1-NEXT: [[CONV:%.*]] = bitcast i64* [[COMP_CASTED]] to i32* 348 // CHECK1-NEXT: store i32 [[TMP0]], i32* [[CONV]], align 4 349 // CHECK1-NEXT: [[TMP1:%.*]] = load i64, i64* [[COMP_CASTED]], align 8 350 // CHECK1-NEXT: [[TMP2:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 351 // CHECK1-NEXT: [[TMP3:%.*]] = bitcast i8** [[TMP2]] to i64* 352 // CHECK1-NEXT: store i64 [[TMP1]], i64* [[TMP3]], align 8 353 // CHECK1-NEXT: [[TMP4:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 0 354 // CHECK1-NEXT: [[TMP5:%.*]] = bitcast i8** [[TMP4]] to i64* 355 // CHECK1-NEXT: store i64 [[TMP1]], i64* [[TMP5]], align 8 356 // CHECK1-NEXT: [[TMP6:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_MAPPERS]], i64 0, i64 0 357 // CHECK1-NEXT: store i8* null, i8** [[TMP6]], align 8 358 // CHECK1-NEXT: [[TMP7:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 359 // CHECK1-NEXT: [[TMP8:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 0 360 // CHECK1-NEXT: [[TMP9:%.*]] = call i32 @__tgt_target_teams_mapper(%struct.ident_t* @[[GLOB1:[0-9]+]], i64 -1, i8* @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l31.region_id, i32 1, i8** [[TMP7]], i8** [[TMP8]], i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_sizes, i32 0, i32 0), i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_maptypes, i32 0, i32 0), i8** null, i8** null, i32 0, i32 0) 361 // CHECK1-NEXT: [[TMP10:%.*]] = icmp ne i32 [[TMP9]], 0 362 // CHECK1-NEXT: br i1 [[TMP10]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]] 363 // CHECK1: omp_offload.failed: 364 // CHECK1-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l31(i64 [[TMP1]]) #[[ATTR2:[0-9]+]] 365 // CHECK1-NEXT: br label [[OMP_OFFLOAD_CONT]] 366 // CHECK1: omp_offload.cont: 367 // CHECK1-NEXT: [[TMP11:%.*]] = load i32, i32* [[COMP]], align 4 368 // CHECK1-NEXT: [[CONV2:%.*]] = bitcast i64* [[COMP_CASTED1]] to i32* 369 // CHECK1-NEXT: store i32 [[TMP11]], i32* [[CONV2]], align 4 370 // CHECK1-NEXT: [[TMP12:%.*]] = load i64, i64* [[COMP_CASTED1]], align 8 371 // CHECK1-NEXT: [[TMP13:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_BASEPTRS3]], i32 0, i32 0 372 // CHECK1-NEXT: [[TMP14:%.*]] = bitcast i8** [[TMP13]] to i64* 373 // CHECK1-NEXT: store i64 [[TMP12]], i64* [[TMP14]], align 8 374 // CHECK1-NEXT: [[TMP15:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_PTRS4]], i32 0, i32 0 375 // CHECK1-NEXT: [[TMP16:%.*]] = bitcast i8** [[TMP15]] to i64* 376 // CHECK1-NEXT: store i64 [[TMP12]], i64* [[TMP16]], align 8 377 // CHECK1-NEXT: [[TMP17:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_MAPPERS5]], i64 0, i64 0 378 // CHECK1-NEXT: store i8* null, i8** [[TMP17]], align 8 379 // CHECK1-NEXT: [[TMP18:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_BASEPTRS3]], i32 0, i32 0 380 // CHECK1-NEXT: [[TMP19:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_PTRS4]], i32 0, i32 0 381 // CHECK1-NEXT: [[TMP20:%.*]] = call i32 @__tgt_target_teams_mapper(%struct.ident_t* @[[GLOB1]], i64 -1, i8* @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l37.region_id, i32 1, i8** [[TMP18]], i8** [[TMP19]], i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_sizes.2, i32 0, i32 0), i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_maptypes.3, i32 0, i32 0), i8** null, i8** null, i32 0, i32 0) 382 // CHECK1-NEXT: [[TMP21:%.*]] = icmp ne i32 [[TMP20]], 0 383 // CHECK1-NEXT: br i1 [[TMP21]], label [[OMP_OFFLOAD_FAILED6:%.*]], label [[OMP_OFFLOAD_CONT7:%.*]] 384 // CHECK1: omp_offload.failed6: 385 // CHECK1-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l37(i64 [[TMP12]]) #[[ATTR2]] 386 // CHECK1-NEXT: br label [[OMP_OFFLOAD_CONT7]] 387 // CHECK1: omp_offload.cont7: 388 // CHECK1-NEXT: [[TMP22:%.*]] = load i32, i32* [[LA]], align 4 389 // CHECK1-NEXT: [[CONV8:%.*]] = bitcast i64* [[LA_CASTED]] to i32* 390 // CHECK1-NEXT: store i32 [[TMP22]], i32* [[CONV8]], align 4 391 // CHECK1-NEXT: [[TMP23:%.*]] = load i64, i64* [[LA_CASTED]], align 8 392 // CHECK1-NEXT: [[TMP24:%.*]] = load i32, i32* [[COMP]], align 4 393 // CHECK1-NEXT: [[CONV10:%.*]] = bitcast i64* [[COMP_CASTED9]] to i32* 394 // CHECK1-NEXT: store i32 [[TMP24]], i32* [[CONV10]], align 4 395 // CHECK1-NEXT: [[TMP25:%.*]] = load i64, i64* [[COMP_CASTED9]], align 8 396 // CHECK1-NEXT: [[TMP26:%.*]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[DOTOFFLOAD_BASEPTRS11]], i32 0, i32 0 397 // CHECK1-NEXT: [[TMP27:%.*]] = bitcast i8** [[TMP26]] to i64* 398 // CHECK1-NEXT: store i64 [[TMP23]], i64* [[TMP27]], align 8 399 // CHECK1-NEXT: [[TMP28:%.*]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[DOTOFFLOAD_PTRS12]], i32 0, i32 0 400 // CHECK1-NEXT: [[TMP29:%.*]] = bitcast i8** [[TMP28]] to i64* 401 // CHECK1-NEXT: store i64 [[TMP23]], i64* [[TMP29]], align 8 402 // CHECK1-NEXT: [[TMP30:%.*]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[DOTOFFLOAD_MAPPERS13]], i64 0, i64 0 403 // CHECK1-NEXT: store i8* null, i8** [[TMP30]], align 8 404 // CHECK1-NEXT: [[TMP31:%.*]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[DOTOFFLOAD_BASEPTRS11]], i32 0, i32 1 405 // CHECK1-NEXT: [[TMP32:%.*]] = bitcast i8** [[TMP31]] to i64* 406 // CHECK1-NEXT: store i64 [[TMP25]], i64* [[TMP32]], align 8 407 // CHECK1-NEXT: [[TMP33:%.*]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[DOTOFFLOAD_PTRS12]], i32 0, i32 1 408 // CHECK1-NEXT: [[TMP34:%.*]] = bitcast i8** [[TMP33]] to i64* 409 // CHECK1-NEXT: store i64 [[TMP25]], i64* [[TMP34]], align 8 410 // CHECK1-NEXT: [[TMP35:%.*]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[DOTOFFLOAD_MAPPERS13]], i64 0, i64 1 411 // CHECK1-NEXT: store i8* null, i8** [[TMP35]], align 8 412 // CHECK1-NEXT: [[TMP36:%.*]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[DOTOFFLOAD_BASEPTRS11]], i32 0, i32 0 413 // CHECK1-NEXT: [[TMP37:%.*]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[DOTOFFLOAD_PTRS12]], i32 0, i32 0 414 // CHECK1-NEXT: [[TMP38:%.*]] = load i32, i32* [[LA]], align 4 415 // CHECK1-NEXT: [[TMP39:%.*]] = call i32 @__tgt_target_teams_mapper(%struct.ident_t* @[[GLOB1]], i64 -1, i8* @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l46.region_id, i32 2, i8** [[TMP36]], i8** [[TMP37]], i64* getelementptr inbounds ([2 x i64], [2 x i64]* @.offload_sizes.5, i32 0, i32 0), i64* getelementptr inbounds ([2 x i64], [2 x i64]* @.offload_maptypes.6, i32 0, i32 0), i8** null, i8** null, i32 [[TMP38]], i32 0) 416 // CHECK1-NEXT: [[TMP40:%.*]] = icmp ne i32 [[TMP39]], 0 417 // CHECK1-NEXT: br i1 [[TMP40]], label [[OMP_OFFLOAD_FAILED14:%.*]], label [[OMP_OFFLOAD_CONT15:%.*]] 418 // CHECK1: omp_offload.failed14: 419 // CHECK1-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l46(i64 [[TMP23]], i64 [[TMP25]]) #[[ATTR2]] 420 // CHECK1-NEXT: br label [[OMP_OFFLOAD_CONT15]] 421 // CHECK1: omp_offload.cont15: 422 // CHECK1-NEXT: [[TMP41:%.*]] = load i32, i32* [[LA]], align 4 423 // CHECK1-NEXT: [[CONV17:%.*]] = bitcast i64* [[LA_CASTED16]] to i32* 424 // CHECK1-NEXT: store i32 [[TMP41]], i32* [[CONV17]], align 4 425 // CHECK1-NEXT: [[TMP42:%.*]] = load i64, i64* [[LA_CASTED16]], align 8 426 // CHECK1-NEXT: [[TMP43:%.*]] = load i32, i32* [[COMP]], align 4 427 // CHECK1-NEXT: [[CONV19:%.*]] = bitcast i64* [[COMP_CASTED18]] to i32* 428 // CHECK1-NEXT: store i32 [[TMP43]], i32* [[CONV19]], align 4 429 // CHECK1-NEXT: [[TMP44:%.*]] = load i64, i64* [[COMP_CASTED18]], align 8 430 // CHECK1-NEXT: [[TMP45:%.*]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[DOTOFFLOAD_BASEPTRS20]], i32 0, i32 0 431 // CHECK1-NEXT: [[TMP46:%.*]] = bitcast i8** [[TMP45]] to i64* 432 // CHECK1-NEXT: store i64 [[TMP42]], i64* [[TMP46]], align 8 433 // CHECK1-NEXT: [[TMP47:%.*]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[DOTOFFLOAD_PTRS21]], i32 0, i32 0 434 // CHECK1-NEXT: [[TMP48:%.*]] = bitcast i8** [[TMP47]] to i64* 435 // CHECK1-NEXT: store i64 [[TMP42]], i64* [[TMP48]], align 8 436 // CHECK1-NEXT: [[TMP49:%.*]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[DOTOFFLOAD_MAPPERS22]], i64 0, i64 0 437 // CHECK1-NEXT: store i8* null, i8** [[TMP49]], align 8 438 // CHECK1-NEXT: [[TMP50:%.*]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[DOTOFFLOAD_BASEPTRS20]], i32 0, i32 1 439 // CHECK1-NEXT: [[TMP51:%.*]] = bitcast i8** [[TMP50]] to i64* 440 // CHECK1-NEXT: store i64 [[TMP44]], i64* [[TMP51]], align 8 441 // CHECK1-NEXT: [[TMP52:%.*]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[DOTOFFLOAD_PTRS21]], i32 0, i32 1 442 // CHECK1-NEXT: [[TMP53:%.*]] = bitcast i8** [[TMP52]] to i64* 443 // CHECK1-NEXT: store i64 [[TMP44]], i64* [[TMP53]], align 8 444 // CHECK1-NEXT: [[TMP54:%.*]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[DOTOFFLOAD_MAPPERS22]], i64 0, i64 1 445 // CHECK1-NEXT: store i8* null, i8** [[TMP54]], align 8 446 // CHECK1-NEXT: [[TMP55:%.*]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[DOTOFFLOAD_BASEPTRS20]], i32 0, i32 0 447 // CHECK1-NEXT: [[TMP56:%.*]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[DOTOFFLOAD_PTRS21]], i32 0, i32 0 448 // CHECK1-NEXT: [[TMP57:%.*]] = load i32, i32* [[LA]], align 4 449 // CHECK1-NEXT: [[TMP58:%.*]] = call i32 @__tgt_target_teams_mapper(%struct.ident_t* @[[GLOB1]], i64 -1, i8* @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l53.region_id, i32 2, i8** [[TMP55]], i8** [[TMP56]], i64* getelementptr inbounds ([2 x i64], [2 x i64]* @.offload_sizes.8, i32 0, i32 0), i64* getelementptr inbounds ([2 x i64], [2 x i64]* @.offload_maptypes.9, i32 0, i32 0), i8** null, i8** null, i32 0, i32 [[TMP57]]) 450 // CHECK1-NEXT: [[TMP59:%.*]] = icmp ne i32 [[TMP58]], 0 451 // CHECK1-NEXT: br i1 [[TMP59]], label [[OMP_OFFLOAD_FAILED23:%.*]], label [[OMP_OFFLOAD_CONT24:%.*]] 452 // CHECK1: omp_offload.failed23: 453 // CHECK1-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l53(i64 [[TMP42]], i64 [[TMP44]]) #[[ATTR2]] 454 // CHECK1-NEXT: br label [[OMP_OFFLOAD_CONT24]] 455 // CHECK1: omp_offload.cont24: 456 // CHECK1-NEXT: [[TMP60:%.*]] = load i32, i32* @Gbla, align 4 457 // CHECK1-NEXT: [[CONV25:%.*]] = bitcast i64* [[GBLA_CASTED]] to i32* 458 // CHECK1-NEXT: store i32 [[TMP60]], i32* [[CONV25]], align 4 459 // CHECK1-NEXT: [[TMP61:%.*]] = load i64, i64* [[GBLA_CASTED]], align 8 460 // CHECK1-NEXT: [[TMP62:%.*]] = load i32, i32* [[A_ADDR]], align 4 461 // CHECK1-NEXT: [[CONV26:%.*]] = bitcast i64* [[A_CASTED]] to i32* 462 // CHECK1-NEXT: store i32 [[TMP62]], i32* [[CONV26]], align 4 463 // CHECK1-NEXT: [[TMP63:%.*]] = load i64, i64* [[A_CASTED]], align 8 464 // CHECK1-NEXT: [[TMP64:%.*]] = load i64, i64* @Gblb, align 8 465 // CHECK1-NEXT: store i64 [[TMP64]], i64* [[GBLB_CASTED]], align 8 466 // CHECK1-NEXT: [[TMP65:%.*]] = load i64, i64* [[GBLB_CASTED]], align 8 467 // CHECK1-NEXT: [[TMP66:%.*]] = load float, float* [[LC]], align 4 468 // CHECK1-NEXT: [[CONV27:%.*]] = bitcast i64* [[LC_CASTED]] to float* 469 // CHECK1-NEXT: store float [[TMP66]], float* [[CONV27]], align 4 470 // CHECK1-NEXT: [[TMP67:%.*]] = load i64, i64* [[LC_CASTED]], align 8 471 // CHECK1-NEXT: [[TMP68:%.*]] = load i32, i32* [[COMP]], align 4 472 // CHECK1-NEXT: [[CONV29:%.*]] = bitcast i64* [[COMP_CASTED28]] to i32* 473 // CHECK1-NEXT: store i32 [[TMP68]], i32* [[CONV29]], align 4 474 // CHECK1-NEXT: [[TMP69:%.*]] = load i64, i64* [[COMP_CASTED28]], align 8 475 // CHECK1-NEXT: [[TMP70:%.*]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[DOTOFFLOAD_BASEPTRS30]], i32 0, i32 0 476 // CHECK1-NEXT: [[TMP71:%.*]] = bitcast i8** [[TMP70]] to i64* 477 // CHECK1-NEXT: store i64 [[TMP61]], i64* [[TMP71]], align 8 478 // CHECK1-NEXT: [[TMP72:%.*]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[DOTOFFLOAD_PTRS31]], i32 0, i32 0 479 // CHECK1-NEXT: [[TMP73:%.*]] = bitcast i8** [[TMP72]] to i64* 480 // CHECK1-NEXT: store i64 [[TMP61]], i64* [[TMP73]], align 8 481 // CHECK1-NEXT: [[TMP74:%.*]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[DOTOFFLOAD_MAPPERS32]], i64 0, i64 0 482 // CHECK1-NEXT: store i8* null, i8** [[TMP74]], align 8 483 // CHECK1-NEXT: [[TMP75:%.*]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[DOTOFFLOAD_BASEPTRS30]], i32 0, i32 1 484 // CHECK1-NEXT: [[TMP76:%.*]] = bitcast i8** [[TMP75]] to i64* 485 // CHECK1-NEXT: store i64 [[TMP63]], i64* [[TMP76]], align 8 486 // CHECK1-NEXT: [[TMP77:%.*]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[DOTOFFLOAD_PTRS31]], i32 0, i32 1 487 // CHECK1-NEXT: [[TMP78:%.*]] = bitcast i8** [[TMP77]] to i64* 488 // CHECK1-NEXT: store i64 [[TMP63]], i64* [[TMP78]], align 8 489 // CHECK1-NEXT: [[TMP79:%.*]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[DOTOFFLOAD_MAPPERS32]], i64 0, i64 1 490 // CHECK1-NEXT: store i8* null, i8** [[TMP79]], align 8 491 // CHECK1-NEXT: [[TMP80:%.*]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[DOTOFFLOAD_BASEPTRS30]], i32 0, i32 2 492 // CHECK1-NEXT: [[TMP81:%.*]] = bitcast i8** [[TMP80]] to i64* 493 // CHECK1-NEXT: store i64 [[TMP65]], i64* [[TMP81]], align 8 494 // CHECK1-NEXT: [[TMP82:%.*]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[DOTOFFLOAD_PTRS31]], i32 0, i32 2 495 // CHECK1-NEXT: [[TMP83:%.*]] = bitcast i8** [[TMP82]] to i64* 496 // CHECK1-NEXT: store i64 [[TMP65]], i64* [[TMP83]], align 8 497 // CHECK1-NEXT: [[TMP84:%.*]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[DOTOFFLOAD_MAPPERS32]], i64 0, i64 2 498 // CHECK1-NEXT: store i8* null, i8** [[TMP84]], align 8 499 // CHECK1-NEXT: [[TMP85:%.*]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[DOTOFFLOAD_BASEPTRS30]], i32 0, i32 3 500 // CHECK1-NEXT: [[TMP86:%.*]] = bitcast i8** [[TMP85]] to i64* 501 // CHECK1-NEXT: store i64 [[TMP67]], i64* [[TMP86]], align 8 502 // CHECK1-NEXT: [[TMP87:%.*]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[DOTOFFLOAD_PTRS31]], i32 0, i32 3 503 // CHECK1-NEXT: [[TMP88:%.*]] = bitcast i8** [[TMP87]] to i64* 504 // CHECK1-NEXT: store i64 [[TMP67]], i64* [[TMP88]], align 8 505 // CHECK1-NEXT: [[TMP89:%.*]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[DOTOFFLOAD_MAPPERS32]], i64 0, i64 3 506 // CHECK1-NEXT: store i8* null, i8** [[TMP89]], align 8 507 // CHECK1-NEXT: [[TMP90:%.*]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[DOTOFFLOAD_BASEPTRS30]], i32 0, i32 4 508 // CHECK1-NEXT: [[TMP91:%.*]] = bitcast i8** [[TMP90]] to i64* 509 // CHECK1-NEXT: store i64 [[TMP69]], i64* [[TMP91]], align 8 510 // CHECK1-NEXT: [[TMP92:%.*]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[DOTOFFLOAD_PTRS31]], i32 0, i32 4 511 // CHECK1-NEXT: [[TMP93:%.*]] = bitcast i8** [[TMP92]] to i64* 512 // CHECK1-NEXT: store i64 [[TMP69]], i64* [[TMP93]], align 8 513 // CHECK1-NEXT: [[TMP94:%.*]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[DOTOFFLOAD_MAPPERS32]], i64 0, i64 4 514 // CHECK1-NEXT: store i8* null, i8** [[TMP94]], align 8 515 // CHECK1-NEXT: [[TMP95:%.*]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[DOTOFFLOAD_BASEPTRS30]], i32 0, i32 0 516 // CHECK1-NEXT: [[TMP96:%.*]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[DOTOFFLOAD_PTRS31]], i32 0, i32 0 517 // CHECK1-NEXT: [[TMP97:%.*]] = load i32, i32* @Gbla, align 4 518 // CHECK1-NEXT: [[TMP98:%.*]] = load i32, i32* [[A_ADDR]], align 4 519 // CHECK1-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP97]], [[TMP98]] 520 // CHECK1-NEXT: [[TMP99:%.*]] = load i64, i64* @Gblb, align 8 521 // CHECK1-NEXT: [[TMP100:%.*]] = load float, float* [[LC]], align 4 522 // CHECK1-NEXT: [[CONV33:%.*]] = fptosi float [[TMP100]] to i64 523 // CHECK1-NEXT: [[ADD34:%.*]] = add nsw i64 [[TMP99]], [[CONV33]] 524 // CHECK1-NEXT: [[TMP101:%.*]] = trunc i64 [[ADD34]] to i32 525 // CHECK1-NEXT: [[TMP102:%.*]] = call i32 @__tgt_target_teams_mapper(%struct.ident_t* @[[GLOB1]], i64 -1, i8* @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l62.region_id, i32 5, i8** [[TMP95]], i8** [[TMP96]], i64* getelementptr inbounds ([5 x i64], [5 x i64]* @.offload_sizes.11, i32 0, i32 0), i64* getelementptr inbounds ([5 x i64], [5 x i64]* @.offload_maptypes.12, i32 0, i32 0), i8** null, i8** null, i32 [[ADD]], i32 [[TMP101]]) 526 // CHECK1-NEXT: [[TMP103:%.*]] = icmp ne i32 [[TMP102]], 0 527 // CHECK1-NEXT: br i1 [[TMP103]], label [[OMP_OFFLOAD_FAILED35:%.*]], label [[OMP_OFFLOAD_CONT36:%.*]] 528 // CHECK1: omp_offload.failed35: 529 // CHECK1-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l62(i64 [[TMP61]], i64 [[TMP63]], i64 [[TMP65]], i64 [[TMP67]], i64 [[TMP69]]) #[[ATTR2]] 530 // CHECK1-NEXT: br label [[OMP_OFFLOAD_CONT36]] 531 // CHECK1: omp_offload.cont36: 532 // CHECK1-NEXT: [[TMP104:%.*]] = load i32*, i32** @Gblc, align 8 533 // CHECK1-NEXT: store i32* [[TMP104]], i32** [[TMP]], align 8 534 // CHECK1-NEXT: [[TMP105:%.*]] = load i32, i32* @Gbla, align 4 535 // CHECK1-NEXT: [[CONV37:%.*]] = bitcast i64* [[GBLC_CASTED]] to i32* 536 // CHECK1-NEXT: store i32 [[TMP105]], i32* [[CONV37]], align 4 537 // CHECK1-NEXT: [[TMP106:%.*]] = load i64, i64* [[GBLC_CASTED]], align 8 538 // CHECK1-NEXT: [[TMP107:%.*]] = load i32, i32* [[COMP]], align 4 539 // CHECK1-NEXT: [[CONV39:%.*]] = bitcast i64* [[COMP_CASTED38]] to i32* 540 // CHECK1-NEXT: store i32 [[TMP107]], i32* [[CONV39]], align 4 541 // CHECK1-NEXT: [[TMP108:%.*]] = load i64, i64* [[COMP_CASTED38]], align 8 542 // CHECK1-NEXT: [[TMP109:%.*]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[DOTOFFLOAD_BASEPTRS40]], i32 0, i32 0 543 // CHECK1-NEXT: [[TMP110:%.*]] = bitcast i8** [[TMP109]] to i64* 544 // CHECK1-NEXT: store i64 [[TMP106]], i64* [[TMP110]], align 8 545 // CHECK1-NEXT: [[TMP111:%.*]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[DOTOFFLOAD_PTRS41]], i32 0, i32 0 546 // CHECK1-NEXT: [[TMP112:%.*]] = bitcast i8** [[TMP111]] to i64* 547 // CHECK1-NEXT: store i64 [[TMP106]], i64* [[TMP112]], align 8 548 // CHECK1-NEXT: [[TMP113:%.*]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[DOTOFFLOAD_MAPPERS42]], i64 0, i64 0 549 // CHECK1-NEXT: store i8* null, i8** [[TMP113]], align 8 550 // CHECK1-NEXT: [[TMP114:%.*]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[DOTOFFLOAD_BASEPTRS40]], i32 0, i32 1 551 // CHECK1-NEXT: [[TMP115:%.*]] = bitcast i8** [[TMP114]] to i64* 552 // CHECK1-NEXT: store i64 [[TMP108]], i64* [[TMP115]], align 8 553 // CHECK1-NEXT: [[TMP116:%.*]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[DOTOFFLOAD_PTRS41]], i32 0, i32 1 554 // CHECK1-NEXT: [[TMP117:%.*]] = bitcast i8** [[TMP116]] to i64* 555 // CHECK1-NEXT: store i64 [[TMP108]], i64* [[TMP117]], align 8 556 // CHECK1-NEXT: [[TMP118:%.*]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[DOTOFFLOAD_MAPPERS42]], i64 0, i64 1 557 // CHECK1-NEXT: store i8* null, i8** [[TMP118]], align 8 558 // CHECK1-NEXT: [[TMP119:%.*]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[DOTOFFLOAD_BASEPTRS40]], i32 0, i32 0 559 // CHECK1-NEXT: [[TMP120:%.*]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[DOTOFFLOAD_PTRS41]], i32 0, i32 0 560 // CHECK1-NEXT: [[TMP121:%.*]] = load i32*, i32** @Gblc, align 8 561 // CHECK1-NEXT: store i32* [[TMP121]], i32** [[_TMP43]], align 8 562 // CHECK1-NEXT: [[TMP122:%.*]] = load i32, i32* @Gbla, align 4 563 // CHECK1-NEXT: [[ADD44:%.*]] = add nsw i32 [[TMP122]], 1 564 // CHECK1-NEXT: [[TMP123:%.*]] = load i32*, i32** @Gblc, align 8 565 // CHECK1-NEXT: store i32* [[TMP123]], i32** [[_TMP45]], align 8 566 // CHECK1-NEXT: [[TMP124:%.*]] = load i32, i32* @Gbla, align 4 567 // CHECK1-NEXT: [[ADD46:%.*]] = add nsw i32 [[TMP124]], 2 568 // CHECK1-NEXT: [[TMP125:%.*]] = call i32 @__tgt_target_teams_mapper(%struct.ident_t* @[[GLOB1]], i64 -1, i8* @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l71.region_id, i32 2, i8** [[TMP119]], i8** [[TMP120]], i64* getelementptr inbounds ([2 x i64], [2 x i64]* @.offload_sizes.14, i32 0, i32 0), i64* getelementptr inbounds ([2 x i64], [2 x i64]* @.offload_maptypes.15, i32 0, i32 0), i8** null, i8** null, i32 [[ADD44]], i32 [[ADD46]]) 569 // CHECK1-NEXT: [[TMP126:%.*]] = icmp ne i32 [[TMP125]], 0 570 // CHECK1-NEXT: br i1 [[TMP126]], label [[OMP_OFFLOAD_FAILED47:%.*]], label [[OMP_OFFLOAD_CONT48:%.*]] 571 // CHECK1: omp_offload.failed47: 572 // CHECK1-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l71(i64 [[TMP106]], i64 [[TMP108]]) #[[ATTR2]] 573 // CHECK1-NEXT: br label [[OMP_OFFLOAD_CONT48]] 574 // CHECK1: omp_offload.cont48: 575 // CHECK1-NEXT: [[TMP127:%.*]] = load i32, i32* [[COMP]], align 4 576 // CHECK1-NEXT: ret i32 [[TMP127]] 577 // 578 // 579 // CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l31 580 // CHECK1-SAME: (i64 noundef [[COMP:%.*]]) #[[ATTR1:[0-9]+]] { 581 // CHECK1-NEXT: entry: 582 // CHECK1-NEXT: [[COMP_ADDR:%.*]] = alloca i64, align 8 583 // CHECK1-NEXT: store i64 [[COMP]], i64* [[COMP_ADDR]], align 8 584 // CHECK1-NEXT: [[CONV:%.*]] = bitcast i64* [[COMP_ADDR]] to i32* 585 // CHECK1-NEXT: call void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(%struct.ident_t* @[[GLOB1]], i32 1, void (i32*, i32*, ...)* bitcast (void (i32*, i32*, i32*)* @.omp_outlined. to void (i32*, i32*, ...)*), i32* [[CONV]]) 586 // CHECK1-NEXT: ret void 587 // 588 // 589 // CHECK1-LABEL: define {{[^@]+}}@.omp_outlined. 590 // CHECK1-SAME: (i32* noalias noundef [[DOTGLOBAL_TID_:%.*]], i32* noalias noundef [[DOTBOUND_TID_:%.*]], i32* noundef nonnull align 4 dereferenceable(4) [[COMP:%.*]]) #[[ATTR1]] { 591 // CHECK1-NEXT: entry: 592 // CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8 593 // CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8 594 // CHECK1-NEXT: [[COMP_ADDR:%.*]] = alloca i32*, align 8 595 // CHECK1-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 8 596 // CHECK1-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 8 597 // CHECK1-NEXT: store i32* [[COMP]], i32** [[COMP_ADDR]], align 8 598 // CHECK1-NEXT: [[TMP0:%.*]] = load i32*, i32** [[COMP_ADDR]], align 8 599 // CHECK1-NEXT: [[TMP1:%.*]] = load i32, i32* [[TMP0]], align 4 600 // CHECK1-NEXT: [[INC:%.*]] = add nsw i32 [[TMP1]], 1 601 // CHECK1-NEXT: store i32 [[INC]], i32* [[TMP0]], align 4 602 // CHECK1-NEXT: ret void 603 // 604 // 605 // CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l37 606 // CHECK1-SAME: (i64 noundef [[COMP:%.*]]) #[[ATTR1]] { 607 // CHECK1-NEXT: entry: 608 // CHECK1-NEXT: [[COMP_ADDR:%.*]] = alloca i64, align 8 609 // CHECK1-NEXT: store i64 [[COMP]], i64* [[COMP_ADDR]], align 8 610 // CHECK1-NEXT: [[CONV:%.*]] = bitcast i64* [[COMP_ADDR]] to i32* 611 // CHECK1-NEXT: call void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(%struct.ident_t* @[[GLOB1]], i32 1, void (i32*, i32*, ...)* bitcast (void (i32*, i32*, i32*)* @.omp_outlined..1 to void (i32*, i32*, ...)*), i32* [[CONV]]) 612 // CHECK1-NEXT: ret void 613 // 614 // 615 // CHECK1-LABEL: define {{[^@]+}}@.omp_outlined..1 616 // CHECK1-SAME: (i32* noalias noundef [[DOTGLOBAL_TID_:%.*]], i32* noalias noundef [[DOTBOUND_TID_:%.*]], i32* noundef nonnull align 4 dereferenceable(4) [[COMP:%.*]]) #[[ATTR1]] { 617 // CHECK1-NEXT: entry: 618 // CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8 619 // CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8 620 // CHECK1-NEXT: [[COMP_ADDR:%.*]] = alloca i32*, align 8 621 // CHECK1-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 8 622 // CHECK1-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 8 623 // CHECK1-NEXT: store i32* [[COMP]], i32** [[COMP_ADDR]], align 8 624 // CHECK1-NEXT: [[TMP0:%.*]] = load i32*, i32** [[COMP_ADDR]], align 8 625 // CHECK1-NEXT: [[TMP1:%.*]] = load i32, i32* [[TMP0]], align 4 626 // CHECK1-NEXT: [[INC:%.*]] = add nsw i32 [[TMP1]], 1 627 // CHECK1-NEXT: store i32 [[INC]], i32* [[TMP0]], align 4 628 // CHECK1-NEXT: ret void 629 // 630 // 631 // CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l46 632 // CHECK1-SAME: (i64 noundef [[LA:%.*]], i64 noundef [[COMP:%.*]]) #[[ATTR1]] { 633 // CHECK1-NEXT: entry: 634 // CHECK1-NEXT: [[LA_ADDR:%.*]] = alloca i64, align 8 635 // CHECK1-NEXT: [[COMP_ADDR:%.*]] = alloca i64, align 8 636 // CHECK1-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1]]) 637 // CHECK1-NEXT: store i64 [[LA]], i64* [[LA_ADDR]], align 8 638 // CHECK1-NEXT: store i64 [[COMP]], i64* [[COMP_ADDR]], align 8 639 // CHECK1-NEXT: [[CONV:%.*]] = bitcast i64* [[LA_ADDR]] to i32* 640 // CHECK1-NEXT: [[CONV1:%.*]] = bitcast i64* [[COMP_ADDR]] to i32* 641 // CHECK1-NEXT: [[TMP1:%.*]] = load i32, i32* [[CONV]], align 4 642 // CHECK1-NEXT: call void @__kmpc_push_num_teams(%struct.ident_t* @[[GLOB1]], i32 [[TMP0]], i32 [[TMP1]], i32 0) 643 // CHECK1-NEXT: call void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(%struct.ident_t* @[[GLOB1]], i32 1, void (i32*, i32*, ...)* bitcast (void (i32*, i32*, i32*)* @.omp_outlined..4 to void (i32*, i32*, ...)*), i32* [[CONV1]]) 644 // CHECK1-NEXT: ret void 645 // 646 // 647 // CHECK1-LABEL: define {{[^@]+}}@.omp_outlined..4 648 // CHECK1-SAME: (i32* noalias noundef [[DOTGLOBAL_TID_:%.*]], i32* noalias noundef [[DOTBOUND_TID_:%.*]], i32* noundef nonnull align 4 dereferenceable(4) [[COMP:%.*]]) #[[ATTR1]] { 649 // CHECK1-NEXT: entry: 650 // CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8 651 // CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8 652 // CHECK1-NEXT: [[COMP_ADDR:%.*]] = alloca i32*, align 8 653 // CHECK1-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 8 654 // CHECK1-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 8 655 // CHECK1-NEXT: store i32* [[COMP]], i32** [[COMP_ADDR]], align 8 656 // CHECK1-NEXT: [[TMP0:%.*]] = load i32*, i32** [[COMP_ADDR]], align 8 657 // CHECK1-NEXT: [[TMP1:%.*]] = load i32, i32* [[TMP0]], align 4 658 // CHECK1-NEXT: [[INC:%.*]] = add nsw i32 [[TMP1]], 1 659 // CHECK1-NEXT: store i32 [[INC]], i32* [[TMP0]], align 4 660 // CHECK1-NEXT: ret void 661 // 662 // 663 // CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l53 664 // CHECK1-SAME: (i64 noundef [[LA:%.*]], i64 noundef [[COMP:%.*]]) #[[ATTR1]] { 665 // CHECK1-NEXT: entry: 666 // CHECK1-NEXT: [[LA_ADDR:%.*]] = alloca i64, align 8 667 // CHECK1-NEXT: [[COMP_ADDR:%.*]] = alloca i64, align 8 668 // CHECK1-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1]]) 669 // CHECK1-NEXT: store i64 [[LA]], i64* [[LA_ADDR]], align 8 670 // CHECK1-NEXT: store i64 [[COMP]], i64* [[COMP_ADDR]], align 8 671 // CHECK1-NEXT: [[CONV:%.*]] = bitcast i64* [[LA_ADDR]] to i32* 672 // CHECK1-NEXT: [[CONV1:%.*]] = bitcast i64* [[COMP_ADDR]] to i32* 673 // CHECK1-NEXT: [[TMP1:%.*]] = load i32, i32* [[CONV]], align 4 674 // CHECK1-NEXT: call void @__kmpc_push_num_teams(%struct.ident_t* @[[GLOB1]], i32 [[TMP0]], i32 0, i32 [[TMP1]]) 675 // CHECK1-NEXT: call void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(%struct.ident_t* @[[GLOB1]], i32 1, void (i32*, i32*, ...)* bitcast (void (i32*, i32*, i32*)* @.omp_outlined..7 to void (i32*, i32*, ...)*), i32* [[CONV1]]) 676 // CHECK1-NEXT: ret void 677 // 678 // 679 // CHECK1-LABEL: define {{[^@]+}}@.omp_outlined..7 680 // CHECK1-SAME: (i32* noalias noundef [[DOTGLOBAL_TID_:%.*]], i32* noalias noundef [[DOTBOUND_TID_:%.*]], i32* noundef nonnull align 4 dereferenceable(4) [[COMP:%.*]]) #[[ATTR1]] { 681 // CHECK1-NEXT: entry: 682 // CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8 683 // CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8 684 // CHECK1-NEXT: [[COMP_ADDR:%.*]] = alloca i32*, align 8 685 // CHECK1-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 8 686 // CHECK1-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 8 687 // CHECK1-NEXT: store i32* [[COMP]], i32** [[COMP_ADDR]], align 8 688 // CHECK1-NEXT: [[TMP0:%.*]] = load i32*, i32** [[COMP_ADDR]], align 8 689 // CHECK1-NEXT: [[TMP1:%.*]] = load i32, i32* [[TMP0]], align 4 690 // CHECK1-NEXT: [[INC:%.*]] = add nsw i32 [[TMP1]], 1 691 // CHECK1-NEXT: store i32 [[INC]], i32* [[TMP0]], align 4 692 // CHECK1-NEXT: ret void 693 // 694 // 695 // CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l62 696 // CHECK1-SAME: (i64 noundef [[GBLA:%.*]], i64 noundef [[A:%.*]], i64 noundef [[GBLB:%.*]], i64 noundef [[LC:%.*]], i64 noundef [[COMP:%.*]]) #[[ATTR1]] { 697 // CHECK1-NEXT: entry: 698 // CHECK1-NEXT: [[GBLA_ADDR:%.*]] = alloca i64, align 8 699 // CHECK1-NEXT: [[A_ADDR:%.*]] = alloca i64, align 8 700 // CHECK1-NEXT: [[GBLB_ADDR:%.*]] = alloca i64, align 8 701 // CHECK1-NEXT: [[LC_ADDR:%.*]] = alloca i64, align 8 702 // CHECK1-NEXT: [[COMP_ADDR:%.*]] = alloca i64, align 8 703 // CHECK1-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1]]) 704 // CHECK1-NEXT: store i64 [[GBLA]], i64* [[GBLA_ADDR]], align 8 705 // CHECK1-NEXT: store i64 [[A]], i64* [[A_ADDR]], align 8 706 // CHECK1-NEXT: store i64 [[GBLB]], i64* [[GBLB_ADDR]], align 8 707 // CHECK1-NEXT: store i64 [[LC]], i64* [[LC_ADDR]], align 8 708 // CHECK1-NEXT: store i64 [[COMP]], i64* [[COMP_ADDR]], align 8 709 // CHECK1-NEXT: [[CONV:%.*]] = bitcast i64* [[GBLA_ADDR]] to i32* 710 // CHECK1-NEXT: [[CONV1:%.*]] = bitcast i64* [[A_ADDR]] to i32* 711 // CHECK1-NEXT: [[CONV2:%.*]] = bitcast i64* [[LC_ADDR]] to float* 712 // CHECK1-NEXT: [[CONV3:%.*]] = bitcast i64* [[COMP_ADDR]] to i32* 713 // CHECK1-NEXT: [[TMP1:%.*]] = load i32, i32* [[CONV]], align 4 714 // CHECK1-NEXT: [[TMP2:%.*]] = load i32, i32* [[CONV1]], align 4 715 // CHECK1-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP1]], [[TMP2]] 716 // CHECK1-NEXT: [[TMP3:%.*]] = load i64, i64* [[GBLB_ADDR]], align 8 717 // CHECK1-NEXT: [[TMP4:%.*]] = load float, float* [[CONV2]], align 4 718 // CHECK1-NEXT: [[CONV4:%.*]] = fptosi float [[TMP4]] to i64 719 // CHECK1-NEXT: [[ADD5:%.*]] = add nsw i64 [[TMP3]], [[CONV4]] 720 // CHECK1-NEXT: [[TMP5:%.*]] = trunc i64 [[ADD5]] to i32 721 // CHECK1-NEXT: call void @__kmpc_push_num_teams(%struct.ident_t* @[[GLOB1]], i32 [[TMP0]], i32 [[ADD]], i32 [[TMP5]]) 722 // CHECK1-NEXT: call void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(%struct.ident_t* @[[GLOB1]], i32 1, void (i32*, i32*, ...)* bitcast (void (i32*, i32*, i32*)* @.omp_outlined..10 to void (i32*, i32*, ...)*), i32* [[CONV3]]) 723 // CHECK1-NEXT: ret void 724 // 725 // 726 // CHECK1-LABEL: define {{[^@]+}}@.omp_outlined..10 727 // CHECK1-SAME: (i32* noalias noundef [[DOTGLOBAL_TID_:%.*]], i32* noalias noundef [[DOTBOUND_TID_:%.*]], i32* noundef nonnull align 4 dereferenceable(4) [[COMP:%.*]]) #[[ATTR1]] { 728 // CHECK1-NEXT: entry: 729 // CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8 730 // CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8 731 // CHECK1-NEXT: [[COMP_ADDR:%.*]] = alloca i32*, align 8 732 // CHECK1-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 8 733 // CHECK1-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 8 734 // CHECK1-NEXT: store i32* [[COMP]], i32** [[COMP_ADDR]], align 8 735 // CHECK1-NEXT: [[TMP0:%.*]] = load i32*, i32** [[COMP_ADDR]], align 8 736 // CHECK1-NEXT: [[TMP1:%.*]] = load i32, i32* [[TMP0]], align 4 737 // CHECK1-NEXT: [[INC:%.*]] = add nsw i32 [[TMP1]], 1 738 // CHECK1-NEXT: store i32 [[INC]], i32* [[TMP0]], align 4 739 // CHECK1-NEXT: ret void 740 // 741 // 742 // CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l71 743 // CHECK1-SAME: (i64 noundef [[GBLC:%.*]], i64 noundef [[COMP:%.*]]) #[[ATTR1]] { 744 // CHECK1-NEXT: entry: 745 // CHECK1-NEXT: [[GBLC_ADDR:%.*]] = alloca i64, align 8 746 // CHECK1-NEXT: [[COMP_ADDR:%.*]] = alloca i64, align 8 747 // CHECK1-NEXT: [[TMP:%.*]] = alloca i32*, align 8 748 // CHECK1-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1]]) 749 // CHECK1-NEXT: store i64 [[GBLC]], i64* [[GBLC_ADDR]], align 8 750 // CHECK1-NEXT: store i64 [[COMP]], i64* [[COMP_ADDR]], align 8 751 // CHECK1-NEXT: [[CONV:%.*]] = bitcast i64* [[GBLC_ADDR]] to i32* 752 // CHECK1-NEXT: [[CONV1:%.*]] = bitcast i64* [[COMP_ADDR]] to i32* 753 // CHECK1-NEXT: store i32* [[CONV]], i32** [[TMP]], align 8 754 // CHECK1-NEXT: [[TMP1:%.*]] = load i32, i32* @Gbla, align 4 755 // CHECK1-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP1]], 1 756 // CHECK1-NEXT: [[TMP2:%.*]] = load i32, i32* @Gbla, align 4 757 // CHECK1-NEXT: [[ADD2:%.*]] = add nsw i32 [[TMP2]], 2 758 // CHECK1-NEXT: call void @__kmpc_push_num_teams(%struct.ident_t* @[[GLOB1]], i32 [[TMP0]], i32 [[ADD]], i32 [[ADD2]]) 759 // CHECK1-NEXT: [[TMP3:%.*]] = load i32*, i32** [[TMP]], align 8 760 // CHECK1-NEXT: call void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(%struct.ident_t* @[[GLOB1]], i32 2, void (i32*, i32*, ...)* bitcast (void (i32*, i32*, i32*, i32*)* @.omp_outlined..13 to void (i32*, i32*, ...)*), i32* [[CONV1]], i32* [[TMP3]]) 761 // CHECK1-NEXT: ret void 762 // 763 // 764 // CHECK1-LABEL: define {{[^@]+}}@.omp_outlined..13 765 // CHECK1-SAME: (i32* noalias noundef [[DOTGLOBAL_TID_:%.*]], i32* noalias noundef [[DOTBOUND_TID_:%.*]], i32* noundef nonnull align 4 dereferenceable(4) [[COMP:%.*]], i32* noundef nonnull align 4 dereferenceable(4) [[GBLC:%.*]]) #[[ATTR1]] { 766 // CHECK1-NEXT: entry: 767 // CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8 768 // CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8 769 // CHECK1-NEXT: [[COMP_ADDR:%.*]] = alloca i32*, align 8 770 // CHECK1-NEXT: [[GBLC_ADDR:%.*]] = alloca i32*, align 8 771 // CHECK1-NEXT: [[TMP:%.*]] = alloca i32*, align 8 772 // CHECK1-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 8 773 // CHECK1-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 8 774 // CHECK1-NEXT: store i32* [[COMP]], i32** [[COMP_ADDR]], align 8 775 // CHECK1-NEXT: store i32* [[GBLC]], i32** [[GBLC_ADDR]], align 8 776 // CHECK1-NEXT: [[TMP0:%.*]] = load i32*, i32** [[COMP_ADDR]], align 8 777 // CHECK1-NEXT: [[TMP1:%.*]] = load i32*, i32** [[GBLC_ADDR]], align 8 778 // CHECK1-NEXT: store i32* [[TMP1]], i32** [[TMP]], align 8 779 // CHECK1-NEXT: [[TMP2:%.*]] = load i32, i32* @Gbla, align 4 780 // CHECK1-NEXT: [[TMP3:%.*]] = load i32, i32* [[TMP0]], align 4 781 // CHECK1-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP3]], [[TMP2]] 782 // CHECK1-NEXT: store i32 [[ADD]], i32* [[TMP0]], align 4 783 // CHECK1-NEXT: ret void 784 // 785 // 786 // CHECK1-LABEL: define {{[^@]+}}@.omp_offloading.requires_reg 787 // CHECK1-SAME: () #[[ATTR3:[0-9]+]] { 788 // CHECK1-NEXT: entry: 789 // CHECK1-NEXT: call void @__tgt_register_requires(i64 1) 790 // CHECK1-NEXT: ret void 791 // 792 // 793 // CHECK3-LABEL: define {{[^@]+}}@_Z27teams_argument_global_locali 794 // CHECK3-SAME: (i32 noundef [[A:%.*]]) #[[ATTR0:[0-9]+]] { 795 // CHECK3-NEXT: entry: 796 // CHECK3-NEXT: [[A_ADDR:%.*]] = alloca i32, align 4 797 // CHECK3-NEXT: [[COMP:%.*]] = alloca i32, align 4 798 // CHECK3-NEXT: [[LA:%.*]] = alloca i32, align 4 799 // CHECK3-NEXT: [[LC:%.*]] = alloca float, align 4 800 // CHECK3-NEXT: [[COMP_CASTED:%.*]] = alloca i32, align 4 801 // CHECK3-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x i8*], align 4 802 // CHECK3-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x i8*], align 4 803 // CHECK3-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [1 x i8*], align 4 804 // CHECK3-NEXT: [[COMP_CASTED1:%.*]] = alloca i32, align 4 805 // CHECK3-NEXT: [[DOTOFFLOAD_BASEPTRS2:%.*]] = alloca [1 x i8*], align 4 806 // CHECK3-NEXT: [[DOTOFFLOAD_PTRS3:%.*]] = alloca [1 x i8*], align 4 807 // CHECK3-NEXT: [[DOTOFFLOAD_MAPPERS4:%.*]] = alloca [1 x i8*], align 4 808 // CHECK3-NEXT: [[LA_CASTED:%.*]] = alloca i32, align 4 809 // CHECK3-NEXT: [[COMP_CASTED7:%.*]] = alloca i32, align 4 810 // CHECK3-NEXT: [[DOTOFFLOAD_BASEPTRS8:%.*]] = alloca [2 x i8*], align 4 811 // CHECK3-NEXT: [[DOTOFFLOAD_PTRS9:%.*]] = alloca [2 x i8*], align 4 812 // CHECK3-NEXT: [[DOTOFFLOAD_MAPPERS10:%.*]] = alloca [2 x i8*], align 4 813 // CHECK3-NEXT: [[LA_CASTED13:%.*]] = alloca i32, align 4 814 // CHECK3-NEXT: [[COMP_CASTED14:%.*]] = alloca i32, align 4 815 // CHECK3-NEXT: [[DOTOFFLOAD_BASEPTRS15:%.*]] = alloca [2 x i8*], align 4 816 // CHECK3-NEXT: [[DOTOFFLOAD_PTRS16:%.*]] = alloca [2 x i8*], align 4 817 // CHECK3-NEXT: [[DOTOFFLOAD_MAPPERS17:%.*]] = alloca [2 x i8*], align 4 818 // CHECK3-NEXT: [[GBLA_CASTED:%.*]] = alloca i32, align 4 819 // CHECK3-NEXT: [[A_CASTED:%.*]] = alloca i32, align 4 820 // CHECK3-NEXT: [[LC_CASTED:%.*]] = alloca i32, align 4 821 // CHECK3-NEXT: [[COMP_CASTED20:%.*]] = alloca i32, align 4 822 // CHECK3-NEXT: [[DOTOFFLOAD_BASEPTRS21:%.*]] = alloca [5 x i8*], align 4 823 // CHECK3-NEXT: [[DOTOFFLOAD_PTRS22:%.*]] = alloca [5 x i8*], align 4 824 // CHECK3-NEXT: [[DOTOFFLOAD_MAPPERS23:%.*]] = alloca [5 x i8*], align 4 825 // CHECK3-NEXT: [[TMP:%.*]] = alloca i32*, align 4 826 // CHECK3-NEXT: [[GBLC_CASTED:%.*]] = alloca i32, align 4 827 // CHECK3-NEXT: [[COMP_CASTED28:%.*]] = alloca i32, align 4 828 // CHECK3-NEXT: [[DOTOFFLOAD_BASEPTRS29:%.*]] = alloca [2 x i8*], align 4 829 // CHECK3-NEXT: [[DOTOFFLOAD_PTRS30:%.*]] = alloca [2 x i8*], align 4 830 // CHECK3-NEXT: [[DOTOFFLOAD_MAPPERS31:%.*]] = alloca [2 x i8*], align 4 831 // CHECK3-NEXT: [[_TMP32:%.*]] = alloca i32*, align 4 832 // CHECK3-NEXT: [[_TMP34:%.*]] = alloca i32*, align 4 833 // CHECK3-NEXT: store i32 [[A]], i32* [[A_ADDR]], align 4 834 // CHECK3-NEXT: store i32 1, i32* [[COMP]], align 4 835 // CHECK3-NEXT: store i32 23, i32* [[LA]], align 4 836 // CHECK3-NEXT: store float 2.500000e+01, float* [[LC]], align 4 837 // CHECK3-NEXT: [[TMP0:%.*]] = load i32, i32* [[COMP]], align 4 838 // CHECK3-NEXT: store i32 [[TMP0]], i32* [[COMP_CASTED]], align 4 839 // CHECK3-NEXT: [[TMP1:%.*]] = load i32, i32* [[COMP_CASTED]], align 4 840 // CHECK3-NEXT: [[TMP2:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 841 // CHECK3-NEXT: [[TMP3:%.*]] = bitcast i8** [[TMP2]] to i32* 842 // CHECK3-NEXT: store i32 [[TMP1]], i32* [[TMP3]], align 4 843 // CHECK3-NEXT: [[TMP4:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 0 844 // CHECK3-NEXT: [[TMP5:%.*]] = bitcast i8** [[TMP4]] to i32* 845 // CHECK3-NEXT: store i32 [[TMP1]], i32* [[TMP5]], align 4 846 // CHECK3-NEXT: [[TMP6:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_MAPPERS]], i32 0, i32 0 847 // CHECK3-NEXT: store i8* null, i8** [[TMP6]], align 4 848 // CHECK3-NEXT: [[TMP7:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 849 // CHECK3-NEXT: [[TMP8:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 0 850 // CHECK3-NEXT: [[TMP9:%.*]] = call i32 @__tgt_target_teams_mapper(%struct.ident_t* @[[GLOB1:[0-9]+]], i64 -1, i8* @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l31.region_id, i32 1, i8** [[TMP7]], i8** [[TMP8]], i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_sizes, i32 0, i32 0), i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_maptypes, i32 0, i32 0), i8** null, i8** null, i32 0, i32 0) 851 // CHECK3-NEXT: [[TMP10:%.*]] = icmp ne i32 [[TMP9]], 0 852 // CHECK3-NEXT: br i1 [[TMP10]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]] 853 // CHECK3: omp_offload.failed: 854 // CHECK3-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l31(i32 [[TMP1]]) #[[ATTR2:[0-9]+]] 855 // CHECK3-NEXT: br label [[OMP_OFFLOAD_CONT]] 856 // CHECK3: omp_offload.cont: 857 // CHECK3-NEXT: [[TMP11:%.*]] = load i32, i32* [[COMP]], align 4 858 // CHECK3-NEXT: store i32 [[TMP11]], i32* [[COMP_CASTED1]], align 4 859 // CHECK3-NEXT: [[TMP12:%.*]] = load i32, i32* [[COMP_CASTED1]], align 4 860 // CHECK3-NEXT: [[TMP13:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_BASEPTRS2]], i32 0, i32 0 861 // CHECK3-NEXT: [[TMP14:%.*]] = bitcast i8** [[TMP13]] to i32* 862 // CHECK3-NEXT: store i32 [[TMP12]], i32* [[TMP14]], align 4 863 // CHECK3-NEXT: [[TMP15:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_PTRS3]], i32 0, i32 0 864 // CHECK3-NEXT: [[TMP16:%.*]] = bitcast i8** [[TMP15]] to i32* 865 // CHECK3-NEXT: store i32 [[TMP12]], i32* [[TMP16]], align 4 866 // CHECK3-NEXT: [[TMP17:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_MAPPERS4]], i32 0, i32 0 867 // CHECK3-NEXT: store i8* null, i8** [[TMP17]], align 4 868 // CHECK3-NEXT: [[TMP18:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_BASEPTRS2]], i32 0, i32 0 869 // CHECK3-NEXT: [[TMP19:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_PTRS3]], i32 0, i32 0 870 // CHECK3-NEXT: [[TMP20:%.*]] = call i32 @__tgt_target_teams_mapper(%struct.ident_t* @[[GLOB1]], i64 -1, i8* @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l37.region_id, i32 1, i8** [[TMP18]], i8** [[TMP19]], i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_sizes.2, i32 0, i32 0), i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_maptypes.3, i32 0, i32 0), i8** null, i8** null, i32 0, i32 0) 871 // CHECK3-NEXT: [[TMP21:%.*]] = icmp ne i32 [[TMP20]], 0 872 // CHECK3-NEXT: br i1 [[TMP21]], label [[OMP_OFFLOAD_FAILED5:%.*]], label [[OMP_OFFLOAD_CONT6:%.*]] 873 // CHECK3: omp_offload.failed5: 874 // CHECK3-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l37(i32 [[TMP12]]) #[[ATTR2]] 875 // CHECK3-NEXT: br label [[OMP_OFFLOAD_CONT6]] 876 // CHECK3: omp_offload.cont6: 877 // CHECK3-NEXT: [[TMP22:%.*]] = load i32, i32* [[LA]], align 4 878 // CHECK3-NEXT: store i32 [[TMP22]], i32* [[LA_CASTED]], align 4 879 // CHECK3-NEXT: [[TMP23:%.*]] = load i32, i32* [[LA_CASTED]], align 4 880 // CHECK3-NEXT: [[TMP24:%.*]] = load i32, i32* [[COMP]], align 4 881 // CHECK3-NEXT: store i32 [[TMP24]], i32* [[COMP_CASTED7]], align 4 882 // CHECK3-NEXT: [[TMP25:%.*]] = load i32, i32* [[COMP_CASTED7]], align 4 883 // CHECK3-NEXT: [[TMP26:%.*]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[DOTOFFLOAD_BASEPTRS8]], i32 0, i32 0 884 // CHECK3-NEXT: [[TMP27:%.*]] = bitcast i8** [[TMP26]] to i32* 885 // CHECK3-NEXT: store i32 [[TMP23]], i32* [[TMP27]], align 4 886 // CHECK3-NEXT: [[TMP28:%.*]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[DOTOFFLOAD_PTRS9]], i32 0, i32 0 887 // CHECK3-NEXT: [[TMP29:%.*]] = bitcast i8** [[TMP28]] to i32* 888 // CHECK3-NEXT: store i32 [[TMP23]], i32* [[TMP29]], align 4 889 // CHECK3-NEXT: [[TMP30:%.*]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[DOTOFFLOAD_MAPPERS10]], i32 0, i32 0 890 // CHECK3-NEXT: store i8* null, i8** [[TMP30]], align 4 891 // CHECK3-NEXT: [[TMP31:%.*]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[DOTOFFLOAD_BASEPTRS8]], i32 0, i32 1 892 // CHECK3-NEXT: [[TMP32:%.*]] = bitcast i8** [[TMP31]] to i32* 893 // CHECK3-NEXT: store i32 [[TMP25]], i32* [[TMP32]], align 4 894 // CHECK3-NEXT: [[TMP33:%.*]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[DOTOFFLOAD_PTRS9]], i32 0, i32 1 895 // CHECK3-NEXT: [[TMP34:%.*]] = bitcast i8** [[TMP33]] to i32* 896 // CHECK3-NEXT: store i32 [[TMP25]], i32* [[TMP34]], align 4 897 // CHECK3-NEXT: [[TMP35:%.*]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[DOTOFFLOAD_MAPPERS10]], i32 0, i32 1 898 // CHECK3-NEXT: store i8* null, i8** [[TMP35]], align 4 899 // CHECK3-NEXT: [[TMP36:%.*]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[DOTOFFLOAD_BASEPTRS8]], i32 0, i32 0 900 // CHECK3-NEXT: [[TMP37:%.*]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[DOTOFFLOAD_PTRS9]], i32 0, i32 0 901 // CHECK3-NEXT: [[TMP38:%.*]] = load i32, i32* [[LA]], align 4 902 // CHECK3-NEXT: [[TMP39:%.*]] = call i32 @__tgt_target_teams_mapper(%struct.ident_t* @[[GLOB1]], i64 -1, i8* @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l46.region_id, i32 2, i8** [[TMP36]], i8** [[TMP37]], i64* getelementptr inbounds ([2 x i64], [2 x i64]* @.offload_sizes.5, i32 0, i32 0), i64* getelementptr inbounds ([2 x i64], [2 x i64]* @.offload_maptypes.6, i32 0, i32 0), i8** null, i8** null, i32 [[TMP38]], i32 0) 903 // CHECK3-NEXT: [[TMP40:%.*]] = icmp ne i32 [[TMP39]], 0 904 // CHECK3-NEXT: br i1 [[TMP40]], label [[OMP_OFFLOAD_FAILED11:%.*]], label [[OMP_OFFLOAD_CONT12:%.*]] 905 // CHECK3: omp_offload.failed11: 906 // CHECK3-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l46(i32 [[TMP23]], i32 [[TMP25]]) #[[ATTR2]] 907 // CHECK3-NEXT: br label [[OMP_OFFLOAD_CONT12]] 908 // CHECK3: omp_offload.cont12: 909 // CHECK3-NEXT: [[TMP41:%.*]] = load i32, i32* [[LA]], align 4 910 // CHECK3-NEXT: store i32 [[TMP41]], i32* [[LA_CASTED13]], align 4 911 // CHECK3-NEXT: [[TMP42:%.*]] = load i32, i32* [[LA_CASTED13]], align 4 912 // CHECK3-NEXT: [[TMP43:%.*]] = load i32, i32* [[COMP]], align 4 913 // CHECK3-NEXT: store i32 [[TMP43]], i32* [[COMP_CASTED14]], align 4 914 // CHECK3-NEXT: [[TMP44:%.*]] = load i32, i32* [[COMP_CASTED14]], align 4 915 // CHECK3-NEXT: [[TMP45:%.*]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[DOTOFFLOAD_BASEPTRS15]], i32 0, i32 0 916 // CHECK3-NEXT: [[TMP46:%.*]] = bitcast i8** [[TMP45]] to i32* 917 // CHECK3-NEXT: store i32 [[TMP42]], i32* [[TMP46]], align 4 918 // CHECK3-NEXT: [[TMP47:%.*]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[DOTOFFLOAD_PTRS16]], i32 0, i32 0 919 // CHECK3-NEXT: [[TMP48:%.*]] = bitcast i8** [[TMP47]] to i32* 920 // CHECK3-NEXT: store i32 [[TMP42]], i32* [[TMP48]], align 4 921 // CHECK3-NEXT: [[TMP49:%.*]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[DOTOFFLOAD_MAPPERS17]], i32 0, i32 0 922 // CHECK3-NEXT: store i8* null, i8** [[TMP49]], align 4 923 // CHECK3-NEXT: [[TMP50:%.*]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[DOTOFFLOAD_BASEPTRS15]], i32 0, i32 1 924 // CHECK3-NEXT: [[TMP51:%.*]] = bitcast i8** [[TMP50]] to i32* 925 // CHECK3-NEXT: store i32 [[TMP44]], i32* [[TMP51]], align 4 926 // CHECK3-NEXT: [[TMP52:%.*]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[DOTOFFLOAD_PTRS16]], i32 0, i32 1 927 // CHECK3-NEXT: [[TMP53:%.*]] = bitcast i8** [[TMP52]] to i32* 928 // CHECK3-NEXT: store i32 [[TMP44]], i32* [[TMP53]], align 4 929 // CHECK3-NEXT: [[TMP54:%.*]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[DOTOFFLOAD_MAPPERS17]], i32 0, i32 1 930 // CHECK3-NEXT: store i8* null, i8** [[TMP54]], align 4 931 // CHECK3-NEXT: [[TMP55:%.*]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[DOTOFFLOAD_BASEPTRS15]], i32 0, i32 0 932 // CHECK3-NEXT: [[TMP56:%.*]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[DOTOFFLOAD_PTRS16]], i32 0, i32 0 933 // CHECK3-NEXT: [[TMP57:%.*]] = load i32, i32* [[LA]], align 4 934 // CHECK3-NEXT: [[TMP58:%.*]] = call i32 @__tgt_target_teams_mapper(%struct.ident_t* @[[GLOB1]], i64 -1, i8* @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l53.region_id, i32 2, i8** [[TMP55]], i8** [[TMP56]], i64* getelementptr inbounds ([2 x i64], [2 x i64]* @.offload_sizes.8, i32 0, i32 0), i64* getelementptr inbounds ([2 x i64], [2 x i64]* @.offload_maptypes.9, i32 0, i32 0), i8** null, i8** null, i32 0, i32 [[TMP57]]) 935 // CHECK3-NEXT: [[TMP59:%.*]] = icmp ne i32 [[TMP58]], 0 936 // CHECK3-NEXT: br i1 [[TMP59]], label [[OMP_OFFLOAD_FAILED18:%.*]], label [[OMP_OFFLOAD_CONT19:%.*]] 937 // CHECK3: omp_offload.failed18: 938 // CHECK3-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l53(i32 [[TMP42]], i32 [[TMP44]]) #[[ATTR2]] 939 // CHECK3-NEXT: br label [[OMP_OFFLOAD_CONT19]] 940 // CHECK3: omp_offload.cont19: 941 // CHECK3-NEXT: [[TMP60:%.*]] = load i32, i32* @Gbla, align 4 942 // CHECK3-NEXT: store i32 [[TMP60]], i32* [[GBLA_CASTED]], align 4 943 // CHECK3-NEXT: [[TMP61:%.*]] = load i32, i32* [[GBLA_CASTED]], align 4 944 // CHECK3-NEXT: [[TMP62:%.*]] = load i32, i32* [[A_ADDR]], align 4 945 // CHECK3-NEXT: store i32 [[TMP62]], i32* [[A_CASTED]], align 4 946 // CHECK3-NEXT: [[TMP63:%.*]] = load i32, i32* [[A_CASTED]], align 4 947 // CHECK3-NEXT: [[TMP64:%.*]] = load float, float* [[LC]], align 4 948 // CHECK3-NEXT: [[CONV:%.*]] = bitcast i32* [[LC_CASTED]] to float* 949 // CHECK3-NEXT: store float [[TMP64]], float* [[CONV]], align 4 950 // CHECK3-NEXT: [[TMP65:%.*]] = load i32, i32* [[LC_CASTED]], align 4 951 // CHECK3-NEXT: [[TMP66:%.*]] = load i32, i32* [[COMP]], align 4 952 // CHECK3-NEXT: store i32 [[TMP66]], i32* [[COMP_CASTED20]], align 4 953 // CHECK3-NEXT: [[TMP67:%.*]] = load i32, i32* [[COMP_CASTED20]], align 4 954 // CHECK3-NEXT: [[TMP68:%.*]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[DOTOFFLOAD_BASEPTRS21]], i32 0, i32 0 955 // CHECK3-NEXT: [[TMP69:%.*]] = bitcast i8** [[TMP68]] to i32* 956 // CHECK3-NEXT: store i32 [[TMP61]], i32* [[TMP69]], align 4 957 // CHECK3-NEXT: [[TMP70:%.*]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[DOTOFFLOAD_PTRS22]], i32 0, i32 0 958 // CHECK3-NEXT: [[TMP71:%.*]] = bitcast i8** [[TMP70]] to i32* 959 // CHECK3-NEXT: store i32 [[TMP61]], i32* [[TMP71]], align 4 960 // CHECK3-NEXT: [[TMP72:%.*]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[DOTOFFLOAD_MAPPERS23]], i32 0, i32 0 961 // CHECK3-NEXT: store i8* null, i8** [[TMP72]], align 4 962 // CHECK3-NEXT: [[TMP73:%.*]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[DOTOFFLOAD_BASEPTRS21]], i32 0, i32 1 963 // CHECK3-NEXT: [[TMP74:%.*]] = bitcast i8** [[TMP73]] to i32* 964 // CHECK3-NEXT: store i32 [[TMP63]], i32* [[TMP74]], align 4 965 // CHECK3-NEXT: [[TMP75:%.*]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[DOTOFFLOAD_PTRS22]], i32 0, i32 1 966 // CHECK3-NEXT: [[TMP76:%.*]] = bitcast i8** [[TMP75]] to i32* 967 // CHECK3-NEXT: store i32 [[TMP63]], i32* [[TMP76]], align 4 968 // CHECK3-NEXT: [[TMP77:%.*]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[DOTOFFLOAD_MAPPERS23]], i32 0, i32 1 969 // CHECK3-NEXT: store i8* null, i8** [[TMP77]], align 4 970 // CHECK3-NEXT: [[TMP78:%.*]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[DOTOFFLOAD_BASEPTRS21]], i32 0, i32 2 971 // CHECK3-NEXT: [[TMP79:%.*]] = bitcast i8** [[TMP78]] to i64** 972 // CHECK3-NEXT: store i64* @Gblb, i64** [[TMP79]], align 4 973 // CHECK3-NEXT: [[TMP80:%.*]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[DOTOFFLOAD_PTRS22]], i32 0, i32 2 974 // CHECK3-NEXT: [[TMP81:%.*]] = bitcast i8** [[TMP80]] to i64** 975 // CHECK3-NEXT: store i64* @Gblb, i64** [[TMP81]], align 4 976 // CHECK3-NEXT: [[TMP82:%.*]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[DOTOFFLOAD_MAPPERS23]], i32 0, i32 2 977 // CHECK3-NEXT: store i8* null, i8** [[TMP82]], align 4 978 // CHECK3-NEXT: [[TMP83:%.*]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[DOTOFFLOAD_BASEPTRS21]], i32 0, i32 3 979 // CHECK3-NEXT: [[TMP84:%.*]] = bitcast i8** [[TMP83]] to i32* 980 // CHECK3-NEXT: store i32 [[TMP65]], i32* [[TMP84]], align 4 981 // CHECK3-NEXT: [[TMP85:%.*]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[DOTOFFLOAD_PTRS22]], i32 0, i32 3 982 // CHECK3-NEXT: [[TMP86:%.*]] = bitcast i8** [[TMP85]] to i32* 983 // CHECK3-NEXT: store i32 [[TMP65]], i32* [[TMP86]], align 4 984 // CHECK3-NEXT: [[TMP87:%.*]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[DOTOFFLOAD_MAPPERS23]], i32 0, i32 3 985 // CHECK3-NEXT: store i8* null, i8** [[TMP87]], align 4 986 // CHECK3-NEXT: [[TMP88:%.*]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[DOTOFFLOAD_BASEPTRS21]], i32 0, i32 4 987 // CHECK3-NEXT: [[TMP89:%.*]] = bitcast i8** [[TMP88]] to i32* 988 // CHECK3-NEXT: store i32 [[TMP67]], i32* [[TMP89]], align 4 989 // CHECK3-NEXT: [[TMP90:%.*]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[DOTOFFLOAD_PTRS22]], i32 0, i32 4 990 // CHECK3-NEXT: [[TMP91:%.*]] = bitcast i8** [[TMP90]] to i32* 991 // CHECK3-NEXT: store i32 [[TMP67]], i32* [[TMP91]], align 4 992 // CHECK3-NEXT: [[TMP92:%.*]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[DOTOFFLOAD_MAPPERS23]], i32 0, i32 4 993 // CHECK3-NEXT: store i8* null, i8** [[TMP92]], align 4 994 // CHECK3-NEXT: [[TMP93:%.*]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[DOTOFFLOAD_BASEPTRS21]], i32 0, i32 0 995 // CHECK3-NEXT: [[TMP94:%.*]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[DOTOFFLOAD_PTRS22]], i32 0, i32 0 996 // CHECK3-NEXT: [[TMP95:%.*]] = load i32, i32* @Gbla, align 4 997 // CHECK3-NEXT: [[TMP96:%.*]] = load i32, i32* [[A_ADDR]], align 4 998 // CHECK3-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP95]], [[TMP96]] 999 // CHECK3-NEXT: [[TMP97:%.*]] = load i64, i64* @Gblb, align 8 1000 // CHECK3-NEXT: [[TMP98:%.*]] = load float, float* [[LC]], align 4 1001 // CHECK3-NEXT: [[CONV24:%.*]] = fptosi float [[TMP98]] to i64 1002 // CHECK3-NEXT: [[ADD25:%.*]] = add nsw i64 [[TMP97]], [[CONV24]] 1003 // CHECK3-NEXT: [[TMP99:%.*]] = trunc i64 [[ADD25]] to i32 1004 // CHECK3-NEXT: [[TMP100:%.*]] = call i32 @__tgt_target_teams_mapper(%struct.ident_t* @[[GLOB1]], i64 -1, i8* @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l62.region_id, i32 5, i8** [[TMP93]], i8** [[TMP94]], i64* getelementptr inbounds ([5 x i64], [5 x i64]* @.offload_sizes.11, i32 0, i32 0), i64* getelementptr inbounds ([5 x i64], [5 x i64]* @.offload_maptypes.12, i32 0, i32 0), i8** null, i8** null, i32 [[ADD]], i32 [[TMP99]]) 1005 // CHECK3-NEXT: [[TMP101:%.*]] = icmp ne i32 [[TMP100]], 0 1006 // CHECK3-NEXT: br i1 [[TMP101]], label [[OMP_OFFLOAD_FAILED26:%.*]], label [[OMP_OFFLOAD_CONT27:%.*]] 1007 // CHECK3: omp_offload.failed26: 1008 // CHECK3-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l62(i32 [[TMP61]], i32 [[TMP63]], i64* @Gblb, i32 [[TMP65]], i32 [[TMP67]]) #[[ATTR2]] 1009 // CHECK3-NEXT: br label [[OMP_OFFLOAD_CONT27]] 1010 // CHECK3: omp_offload.cont27: 1011 // CHECK3-NEXT: [[TMP102:%.*]] = load i32*, i32** @Gblc, align 4 1012 // CHECK3-NEXT: store i32* [[TMP102]], i32** [[TMP]], align 4 1013 // CHECK3-NEXT: [[TMP103:%.*]] = load i32, i32* @Gbla, align 4 1014 // CHECK3-NEXT: store i32 [[TMP103]], i32* [[GBLC_CASTED]], align 4 1015 // CHECK3-NEXT: [[TMP104:%.*]] = load i32, i32* [[GBLC_CASTED]], align 4 1016 // CHECK3-NEXT: [[TMP105:%.*]] = load i32, i32* [[COMP]], align 4 1017 // CHECK3-NEXT: store i32 [[TMP105]], i32* [[COMP_CASTED28]], align 4 1018 // CHECK3-NEXT: [[TMP106:%.*]] = load i32, i32* [[COMP_CASTED28]], align 4 1019 // CHECK3-NEXT: [[TMP107:%.*]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[DOTOFFLOAD_BASEPTRS29]], i32 0, i32 0 1020 // CHECK3-NEXT: [[TMP108:%.*]] = bitcast i8** [[TMP107]] to i32* 1021 // CHECK3-NEXT: store i32 [[TMP104]], i32* [[TMP108]], align 4 1022 // CHECK3-NEXT: [[TMP109:%.*]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[DOTOFFLOAD_PTRS30]], i32 0, i32 0 1023 // CHECK3-NEXT: [[TMP110:%.*]] = bitcast i8** [[TMP109]] to i32* 1024 // CHECK3-NEXT: store i32 [[TMP104]], i32* [[TMP110]], align 4 1025 // CHECK3-NEXT: [[TMP111:%.*]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[DOTOFFLOAD_MAPPERS31]], i32 0, i32 0 1026 // CHECK3-NEXT: store i8* null, i8** [[TMP111]], align 4 1027 // CHECK3-NEXT: [[TMP112:%.*]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[DOTOFFLOAD_BASEPTRS29]], i32 0, i32 1 1028 // CHECK3-NEXT: [[TMP113:%.*]] = bitcast i8** [[TMP112]] to i32* 1029 // CHECK3-NEXT: store i32 [[TMP106]], i32* [[TMP113]], align 4 1030 // CHECK3-NEXT: [[TMP114:%.*]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[DOTOFFLOAD_PTRS30]], i32 0, i32 1 1031 // CHECK3-NEXT: [[TMP115:%.*]] = bitcast i8** [[TMP114]] to i32* 1032 // CHECK3-NEXT: store i32 [[TMP106]], i32* [[TMP115]], align 4 1033 // CHECK3-NEXT: [[TMP116:%.*]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[DOTOFFLOAD_MAPPERS31]], i32 0, i32 1 1034 // CHECK3-NEXT: store i8* null, i8** [[TMP116]], align 4 1035 // CHECK3-NEXT: [[TMP117:%.*]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[DOTOFFLOAD_BASEPTRS29]], i32 0, i32 0 1036 // CHECK3-NEXT: [[TMP118:%.*]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[DOTOFFLOAD_PTRS30]], i32 0, i32 0 1037 // CHECK3-NEXT: [[TMP119:%.*]] = load i32*, i32** @Gblc, align 4 1038 // CHECK3-NEXT: store i32* [[TMP119]], i32** [[_TMP32]], align 4 1039 // CHECK3-NEXT: [[TMP120:%.*]] = load i32, i32* @Gbla, align 4 1040 // CHECK3-NEXT: [[ADD33:%.*]] = add nsw i32 [[TMP120]], 1 1041 // CHECK3-NEXT: [[TMP121:%.*]] = load i32*, i32** @Gblc, align 4 1042 // CHECK3-NEXT: store i32* [[TMP121]], i32** [[_TMP34]], align 4 1043 // CHECK3-NEXT: [[TMP122:%.*]] = load i32, i32* @Gbla, align 4 1044 // CHECK3-NEXT: [[ADD35:%.*]] = add nsw i32 [[TMP122]], 2 1045 // CHECK3-NEXT: [[TMP123:%.*]] = call i32 @__tgt_target_teams_mapper(%struct.ident_t* @[[GLOB1]], i64 -1, i8* @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l71.region_id, i32 2, i8** [[TMP117]], i8** [[TMP118]], i64* getelementptr inbounds ([2 x i64], [2 x i64]* @.offload_sizes.14, i32 0, i32 0), i64* getelementptr inbounds ([2 x i64], [2 x i64]* @.offload_maptypes.15, i32 0, i32 0), i8** null, i8** null, i32 [[ADD33]], i32 [[ADD35]]) 1046 // CHECK3-NEXT: [[TMP124:%.*]] = icmp ne i32 [[TMP123]], 0 1047 // CHECK3-NEXT: br i1 [[TMP124]], label [[OMP_OFFLOAD_FAILED36:%.*]], label [[OMP_OFFLOAD_CONT37:%.*]] 1048 // CHECK3: omp_offload.failed36: 1049 // CHECK3-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l71(i32 [[TMP104]], i32 [[TMP106]]) #[[ATTR2]] 1050 // CHECK3-NEXT: br label [[OMP_OFFLOAD_CONT37]] 1051 // CHECK3: omp_offload.cont37: 1052 // CHECK3-NEXT: [[TMP125:%.*]] = load i32, i32* [[COMP]], align 4 1053 // CHECK3-NEXT: ret i32 [[TMP125]] 1054 // 1055 // 1056 // CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l31 1057 // CHECK3-SAME: (i32 noundef [[COMP:%.*]]) #[[ATTR1:[0-9]+]] { 1058 // CHECK3-NEXT: entry: 1059 // CHECK3-NEXT: [[COMP_ADDR:%.*]] = alloca i32, align 4 1060 // CHECK3-NEXT: store i32 [[COMP]], i32* [[COMP_ADDR]], align 4 1061 // CHECK3-NEXT: call void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(%struct.ident_t* @[[GLOB1]], i32 1, void (i32*, i32*, ...)* bitcast (void (i32*, i32*, i32*)* @.omp_outlined. to void (i32*, i32*, ...)*), i32* [[COMP_ADDR]]) 1062 // CHECK3-NEXT: ret void 1063 // 1064 // 1065 // CHECK3-LABEL: define {{[^@]+}}@.omp_outlined. 1066 // CHECK3-SAME: (i32* noalias noundef [[DOTGLOBAL_TID_:%.*]], i32* noalias noundef [[DOTBOUND_TID_:%.*]], i32* noundef nonnull align 4 dereferenceable(4) [[COMP:%.*]]) #[[ATTR1]] { 1067 // CHECK3-NEXT: entry: 1068 // CHECK3-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 4 1069 // CHECK3-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 4 1070 // CHECK3-NEXT: [[COMP_ADDR:%.*]] = alloca i32*, align 4 1071 // CHECK3-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 4 1072 // CHECK3-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 4 1073 // CHECK3-NEXT: store i32* [[COMP]], i32** [[COMP_ADDR]], align 4 1074 // CHECK3-NEXT: [[TMP0:%.*]] = load i32*, i32** [[COMP_ADDR]], align 4 1075 // CHECK3-NEXT: [[TMP1:%.*]] = load i32, i32* [[TMP0]], align 4 1076 // CHECK3-NEXT: [[INC:%.*]] = add nsw i32 [[TMP1]], 1 1077 // CHECK3-NEXT: store i32 [[INC]], i32* [[TMP0]], align 4 1078 // CHECK3-NEXT: ret void 1079 // 1080 // 1081 // CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l37 1082 // CHECK3-SAME: (i32 noundef [[COMP:%.*]]) #[[ATTR1]] { 1083 // CHECK3-NEXT: entry: 1084 // CHECK3-NEXT: [[COMP_ADDR:%.*]] = alloca i32, align 4 1085 // CHECK3-NEXT: store i32 [[COMP]], i32* [[COMP_ADDR]], align 4 1086 // CHECK3-NEXT: call void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(%struct.ident_t* @[[GLOB1]], i32 1, void (i32*, i32*, ...)* bitcast (void (i32*, i32*, i32*)* @.omp_outlined..1 to void (i32*, i32*, ...)*), i32* [[COMP_ADDR]]) 1087 // CHECK3-NEXT: ret void 1088 // 1089 // 1090 // CHECK3-LABEL: define {{[^@]+}}@.omp_outlined..1 1091 // CHECK3-SAME: (i32* noalias noundef [[DOTGLOBAL_TID_:%.*]], i32* noalias noundef [[DOTBOUND_TID_:%.*]], i32* noundef nonnull align 4 dereferenceable(4) [[COMP:%.*]]) #[[ATTR1]] { 1092 // CHECK3-NEXT: entry: 1093 // CHECK3-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 4 1094 // CHECK3-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 4 1095 // CHECK3-NEXT: [[COMP_ADDR:%.*]] = alloca i32*, align 4 1096 // CHECK3-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 4 1097 // CHECK3-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 4 1098 // CHECK3-NEXT: store i32* [[COMP]], i32** [[COMP_ADDR]], align 4 1099 // CHECK3-NEXT: [[TMP0:%.*]] = load i32*, i32** [[COMP_ADDR]], align 4 1100 // CHECK3-NEXT: [[TMP1:%.*]] = load i32, i32* [[TMP0]], align 4 1101 // CHECK3-NEXT: [[INC:%.*]] = add nsw i32 [[TMP1]], 1 1102 // CHECK3-NEXT: store i32 [[INC]], i32* [[TMP0]], align 4 1103 // CHECK3-NEXT: ret void 1104 // 1105 // 1106 // CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l46 1107 // CHECK3-SAME: (i32 noundef [[LA:%.*]], i32 noundef [[COMP:%.*]]) #[[ATTR1]] { 1108 // CHECK3-NEXT: entry: 1109 // CHECK3-NEXT: [[LA_ADDR:%.*]] = alloca i32, align 4 1110 // CHECK3-NEXT: [[COMP_ADDR:%.*]] = alloca i32, align 4 1111 // CHECK3-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1]]) 1112 // CHECK3-NEXT: store i32 [[LA]], i32* [[LA_ADDR]], align 4 1113 // CHECK3-NEXT: store i32 [[COMP]], i32* [[COMP_ADDR]], align 4 1114 // CHECK3-NEXT: [[TMP1:%.*]] = load i32, i32* [[LA_ADDR]], align 4 1115 // CHECK3-NEXT: call void @__kmpc_push_num_teams(%struct.ident_t* @[[GLOB1]], i32 [[TMP0]], i32 [[TMP1]], i32 0) 1116 // CHECK3-NEXT: call void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(%struct.ident_t* @[[GLOB1]], i32 1, void (i32*, i32*, ...)* bitcast (void (i32*, i32*, i32*)* @.omp_outlined..4 to void (i32*, i32*, ...)*), i32* [[COMP_ADDR]]) 1117 // CHECK3-NEXT: ret void 1118 // 1119 // 1120 // CHECK3-LABEL: define {{[^@]+}}@.omp_outlined..4 1121 // CHECK3-SAME: (i32* noalias noundef [[DOTGLOBAL_TID_:%.*]], i32* noalias noundef [[DOTBOUND_TID_:%.*]], i32* noundef nonnull align 4 dereferenceable(4) [[COMP:%.*]]) #[[ATTR1]] { 1122 // CHECK3-NEXT: entry: 1123 // CHECK3-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 4 1124 // CHECK3-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 4 1125 // CHECK3-NEXT: [[COMP_ADDR:%.*]] = alloca i32*, align 4 1126 // CHECK3-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 4 1127 // CHECK3-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 4 1128 // CHECK3-NEXT: store i32* [[COMP]], i32** [[COMP_ADDR]], align 4 1129 // CHECK3-NEXT: [[TMP0:%.*]] = load i32*, i32** [[COMP_ADDR]], align 4 1130 // CHECK3-NEXT: [[TMP1:%.*]] = load i32, i32* [[TMP0]], align 4 1131 // CHECK3-NEXT: [[INC:%.*]] = add nsw i32 [[TMP1]], 1 1132 // CHECK3-NEXT: store i32 [[INC]], i32* [[TMP0]], align 4 1133 // CHECK3-NEXT: ret void 1134 // 1135 // 1136 // CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l53 1137 // CHECK3-SAME: (i32 noundef [[LA:%.*]], i32 noundef [[COMP:%.*]]) #[[ATTR1]] { 1138 // CHECK3-NEXT: entry: 1139 // CHECK3-NEXT: [[LA_ADDR:%.*]] = alloca i32, align 4 1140 // CHECK3-NEXT: [[COMP_ADDR:%.*]] = alloca i32, align 4 1141 // CHECK3-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1]]) 1142 // CHECK3-NEXT: store i32 [[LA]], i32* [[LA_ADDR]], align 4 1143 // CHECK3-NEXT: store i32 [[COMP]], i32* [[COMP_ADDR]], align 4 1144 // CHECK3-NEXT: [[TMP1:%.*]] = load i32, i32* [[LA_ADDR]], align 4 1145 // CHECK3-NEXT: call void @__kmpc_push_num_teams(%struct.ident_t* @[[GLOB1]], i32 [[TMP0]], i32 0, i32 [[TMP1]]) 1146 // CHECK3-NEXT: call void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(%struct.ident_t* @[[GLOB1]], i32 1, void (i32*, i32*, ...)* bitcast (void (i32*, i32*, i32*)* @.omp_outlined..7 to void (i32*, i32*, ...)*), i32* [[COMP_ADDR]]) 1147 // CHECK3-NEXT: ret void 1148 // 1149 // 1150 // CHECK3-LABEL: define {{[^@]+}}@.omp_outlined..7 1151 // CHECK3-SAME: (i32* noalias noundef [[DOTGLOBAL_TID_:%.*]], i32* noalias noundef [[DOTBOUND_TID_:%.*]], i32* noundef nonnull align 4 dereferenceable(4) [[COMP:%.*]]) #[[ATTR1]] { 1152 // CHECK3-NEXT: entry: 1153 // CHECK3-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 4 1154 // CHECK3-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 4 1155 // CHECK3-NEXT: [[COMP_ADDR:%.*]] = alloca i32*, align 4 1156 // CHECK3-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 4 1157 // CHECK3-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 4 1158 // CHECK3-NEXT: store i32* [[COMP]], i32** [[COMP_ADDR]], align 4 1159 // CHECK3-NEXT: [[TMP0:%.*]] = load i32*, i32** [[COMP_ADDR]], align 4 1160 // CHECK3-NEXT: [[TMP1:%.*]] = load i32, i32* [[TMP0]], align 4 1161 // CHECK3-NEXT: [[INC:%.*]] = add nsw i32 [[TMP1]], 1 1162 // CHECK3-NEXT: store i32 [[INC]], i32* [[TMP0]], align 4 1163 // CHECK3-NEXT: ret void 1164 // 1165 // 1166 // CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l62 1167 // CHECK3-SAME: (i32 noundef [[GBLA:%.*]], i32 noundef [[A:%.*]], i64* noundef nonnull align 4 dereferenceable(8) [[GBLB:%.*]], i32 noundef [[LC:%.*]], i32 noundef [[COMP:%.*]]) #[[ATTR1]] { 1168 // CHECK3-NEXT: entry: 1169 // CHECK3-NEXT: [[GBLA_ADDR:%.*]] = alloca i32, align 4 1170 // CHECK3-NEXT: [[A_ADDR:%.*]] = alloca i32, align 4 1171 // CHECK3-NEXT: [[GBLB_ADDR:%.*]] = alloca i64*, align 4 1172 // CHECK3-NEXT: [[LC_ADDR:%.*]] = alloca i32, align 4 1173 // CHECK3-NEXT: [[COMP_ADDR:%.*]] = alloca i32, align 4 1174 // CHECK3-NEXT: [[GBLB1:%.*]] = alloca i64, align 8 1175 // CHECK3-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1]]) 1176 // CHECK3-NEXT: store i32 [[GBLA]], i32* [[GBLA_ADDR]], align 4 1177 // CHECK3-NEXT: store i32 [[A]], i32* [[A_ADDR]], align 4 1178 // CHECK3-NEXT: store i64* [[GBLB]], i64** [[GBLB_ADDR]], align 4 1179 // CHECK3-NEXT: store i32 [[LC]], i32* [[LC_ADDR]], align 4 1180 // CHECK3-NEXT: store i32 [[COMP]], i32* [[COMP_ADDR]], align 4 1181 // CHECK3-NEXT: [[TMP1:%.*]] = load i64*, i64** [[GBLB_ADDR]], align 4 1182 // CHECK3-NEXT: [[CONV:%.*]] = bitcast i32* [[LC_ADDR]] to float* 1183 // CHECK3-NEXT: [[TMP2:%.*]] = load i64, i64* [[TMP1]], align 8 1184 // CHECK3-NEXT: store i64 [[TMP2]], i64* [[GBLB1]], align 8 1185 // CHECK3-NEXT: [[TMP3:%.*]] = load i32, i32* [[GBLA_ADDR]], align 4 1186 // CHECK3-NEXT: [[TMP4:%.*]] = load i32, i32* [[A_ADDR]], align 4 1187 // CHECK3-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP3]], [[TMP4]] 1188 // CHECK3-NEXT: [[TMP5:%.*]] = load i64, i64* [[GBLB1]], align 8 1189 // CHECK3-NEXT: [[TMP6:%.*]] = load float, float* [[CONV]], align 4 1190 // CHECK3-NEXT: [[CONV2:%.*]] = fptosi float [[TMP6]] to i64 1191 // CHECK3-NEXT: [[ADD3:%.*]] = add nsw i64 [[TMP5]], [[CONV2]] 1192 // CHECK3-NEXT: [[TMP7:%.*]] = trunc i64 [[ADD3]] to i32 1193 // CHECK3-NEXT: call void @__kmpc_push_num_teams(%struct.ident_t* @[[GLOB1]], i32 [[TMP0]], i32 [[ADD]], i32 [[TMP7]]) 1194 // CHECK3-NEXT: call void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(%struct.ident_t* @[[GLOB1]], i32 1, void (i32*, i32*, ...)* bitcast (void (i32*, i32*, i32*)* @.omp_outlined..10 to void (i32*, i32*, ...)*), i32* [[COMP_ADDR]]) 1195 // CHECK3-NEXT: ret void 1196 // 1197 // 1198 // CHECK3-LABEL: define {{[^@]+}}@.omp_outlined..10 1199 // CHECK3-SAME: (i32* noalias noundef [[DOTGLOBAL_TID_:%.*]], i32* noalias noundef [[DOTBOUND_TID_:%.*]], i32* noundef nonnull align 4 dereferenceable(4) [[COMP:%.*]]) #[[ATTR1]] { 1200 // CHECK3-NEXT: entry: 1201 // CHECK3-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 4 1202 // CHECK3-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 4 1203 // CHECK3-NEXT: [[COMP_ADDR:%.*]] = alloca i32*, align 4 1204 // CHECK3-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 4 1205 // CHECK3-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 4 1206 // CHECK3-NEXT: store i32* [[COMP]], i32** [[COMP_ADDR]], align 4 1207 // CHECK3-NEXT: [[TMP0:%.*]] = load i32*, i32** [[COMP_ADDR]], align 4 1208 // CHECK3-NEXT: [[TMP1:%.*]] = load i32, i32* [[TMP0]], align 4 1209 // CHECK3-NEXT: [[INC:%.*]] = add nsw i32 [[TMP1]], 1 1210 // CHECK3-NEXT: store i32 [[INC]], i32* [[TMP0]], align 4 1211 // CHECK3-NEXT: ret void 1212 // 1213 // 1214 // CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l71 1215 // CHECK3-SAME: (i32 noundef [[GBLC:%.*]], i32 noundef [[COMP:%.*]]) #[[ATTR1]] { 1216 // CHECK3-NEXT: entry: 1217 // CHECK3-NEXT: [[GBLC_ADDR:%.*]] = alloca i32, align 4 1218 // CHECK3-NEXT: [[COMP_ADDR:%.*]] = alloca i32, align 4 1219 // CHECK3-NEXT: [[TMP:%.*]] = alloca i32*, align 4 1220 // CHECK3-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1]]) 1221 // CHECK3-NEXT: store i32 [[GBLC]], i32* [[GBLC_ADDR]], align 4 1222 // CHECK3-NEXT: store i32 [[COMP]], i32* [[COMP_ADDR]], align 4 1223 // CHECK3-NEXT: store i32* [[GBLC_ADDR]], i32** [[TMP]], align 4 1224 // CHECK3-NEXT: [[TMP1:%.*]] = load i32, i32* @Gbla, align 4 1225 // CHECK3-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP1]], 1 1226 // CHECK3-NEXT: [[TMP2:%.*]] = load i32, i32* @Gbla, align 4 1227 // CHECK3-NEXT: [[ADD1:%.*]] = add nsw i32 [[TMP2]], 2 1228 // CHECK3-NEXT: call void @__kmpc_push_num_teams(%struct.ident_t* @[[GLOB1]], i32 [[TMP0]], i32 [[ADD]], i32 [[ADD1]]) 1229 // CHECK3-NEXT: [[TMP3:%.*]] = load i32*, i32** [[TMP]], align 4 1230 // CHECK3-NEXT: call void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(%struct.ident_t* @[[GLOB1]], i32 2, void (i32*, i32*, ...)* bitcast (void (i32*, i32*, i32*, i32*)* @.omp_outlined..13 to void (i32*, i32*, ...)*), i32* [[COMP_ADDR]], i32* [[TMP3]]) 1231 // CHECK3-NEXT: ret void 1232 // 1233 // 1234 // CHECK3-LABEL: define {{[^@]+}}@.omp_outlined..13 1235 // CHECK3-SAME: (i32* noalias noundef [[DOTGLOBAL_TID_:%.*]], i32* noalias noundef [[DOTBOUND_TID_:%.*]], i32* noundef nonnull align 4 dereferenceable(4) [[COMP:%.*]], i32* noundef nonnull align 4 dereferenceable(4) [[GBLC:%.*]]) #[[ATTR1]] { 1236 // CHECK3-NEXT: entry: 1237 // CHECK3-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 4 1238 // CHECK3-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 4 1239 // CHECK3-NEXT: [[COMP_ADDR:%.*]] = alloca i32*, align 4 1240 // CHECK3-NEXT: [[GBLC_ADDR:%.*]] = alloca i32*, align 4 1241 // CHECK3-NEXT: [[TMP:%.*]] = alloca i32*, align 4 1242 // CHECK3-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 4 1243 // CHECK3-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 4 1244 // CHECK3-NEXT: store i32* [[COMP]], i32** [[COMP_ADDR]], align 4 1245 // CHECK3-NEXT: store i32* [[GBLC]], i32** [[GBLC_ADDR]], align 4 1246 // CHECK3-NEXT: [[TMP0:%.*]] = load i32*, i32** [[COMP_ADDR]], align 4 1247 // CHECK3-NEXT: [[TMP1:%.*]] = load i32*, i32** [[GBLC_ADDR]], align 4 1248 // CHECK3-NEXT: store i32* [[TMP1]], i32** [[TMP]], align 4 1249 // CHECK3-NEXT: [[TMP2:%.*]] = load i32, i32* @Gbla, align 4 1250 // CHECK3-NEXT: [[TMP3:%.*]] = load i32, i32* [[TMP0]], align 4 1251 // CHECK3-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP3]], [[TMP2]] 1252 // CHECK3-NEXT: store i32 [[ADD]], i32* [[TMP0]], align 4 1253 // CHECK3-NEXT: ret void 1254 // 1255 // 1256 // CHECK3-LABEL: define {{[^@]+}}@.omp_offloading.requires_reg 1257 // CHECK3-SAME: () #[[ATTR3:[0-9]+]] { 1258 // CHECK3-NEXT: entry: 1259 // CHECK3-NEXT: call void @__tgt_register_requires(i64 1) 1260 // CHECK3-NEXT: ret void 1261 // 1262 // 1263 // CHECK9-LABEL: define {{[^@]+}}@_Z18teams_template_argv 1264 // CHECK9-SAME: () #[[ATTR0:[0-9]+]] { 1265 // CHECK9-NEXT: entry: 1266 // CHECK9-NEXT: [[COMP:%.*]] = alloca i32, align 4 1267 // CHECK9-NEXT: [[LA:%.*]] = alloca [[STRUCT_SS:%.*]], align 4 1268 // CHECK9-NEXT: [[LB:%.*]] = alloca [[STRUCT_SS_0:%.*]], align 8 1269 // CHECK9-NEXT: [[COMP_CASTED:%.*]] = alloca i64, align 8 1270 // CHECK9-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [3 x i8*], align 8 1271 // CHECK9-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [3 x i8*], align 8 1272 // CHECK9-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [3 x i8*], align 8 1273 // CHECK9-NEXT: [[COMP_CASTED2:%.*]] = alloca i64, align 8 1274 // CHECK9-NEXT: [[DOTOFFLOAD_BASEPTRS4:%.*]] = alloca [3 x i8*], align 8 1275 // CHECK9-NEXT: [[DOTOFFLOAD_PTRS5:%.*]] = alloca [3 x i8*], align 8 1276 // CHECK9-NEXT: [[DOTOFFLOAD_MAPPERS6:%.*]] = alloca [3 x i8*], align 8 1277 // CHECK9-NEXT: store i32 1, i32* [[COMP]], align 4 1278 // CHECK9-NEXT: [[TMP0:%.*]] = load i32, i32* [[COMP]], align 4 1279 // CHECK9-NEXT: [[CONV:%.*]] = bitcast i64* [[COMP_CASTED]] to i32* 1280 // CHECK9-NEXT: store i32 [[TMP0]], i32* [[CONV]], align 4 1281 // CHECK9-NEXT: [[TMP1:%.*]] = load i64, i64* [[COMP_CASTED]], align 8 1282 // CHECK9-NEXT: [[TMP2:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 1283 // CHECK9-NEXT: [[TMP3:%.*]] = bitcast i8** [[TMP2]] to %struct.SS** 1284 // CHECK9-NEXT: store %struct.SS* @Gbla, %struct.SS** [[TMP3]], align 8 1285 // CHECK9-NEXT: [[TMP4:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 0 1286 // CHECK9-NEXT: [[TMP5:%.*]] = bitcast i8** [[TMP4]] to %struct.SS** 1287 // CHECK9-NEXT: store %struct.SS* @Gbla, %struct.SS** [[TMP5]], align 8 1288 // CHECK9-NEXT: [[TMP6:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_MAPPERS]], i64 0, i64 0 1289 // CHECK9-NEXT: store i8* null, i8** [[TMP6]], align 8 1290 // CHECK9-NEXT: [[TMP7:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 1 1291 // CHECK9-NEXT: [[TMP8:%.*]] = bitcast i8** [[TMP7]] to %struct.SS** 1292 // CHECK9-NEXT: store %struct.SS* [[LA]], %struct.SS** [[TMP8]], align 8 1293 // CHECK9-NEXT: [[TMP9:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 1 1294 // CHECK9-NEXT: [[TMP10:%.*]] = bitcast i8** [[TMP9]] to %struct.SS** 1295 // CHECK9-NEXT: store %struct.SS* [[LA]], %struct.SS** [[TMP10]], align 8 1296 // CHECK9-NEXT: [[TMP11:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_MAPPERS]], i64 0, i64 1 1297 // CHECK9-NEXT: store i8* null, i8** [[TMP11]], align 8 1298 // CHECK9-NEXT: [[TMP12:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 2 1299 // CHECK9-NEXT: [[TMP13:%.*]] = bitcast i8** [[TMP12]] to i64* 1300 // CHECK9-NEXT: store i64 [[TMP1]], i64* [[TMP13]], align 8 1301 // CHECK9-NEXT: [[TMP14:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 2 1302 // CHECK9-NEXT: [[TMP15:%.*]] = bitcast i8** [[TMP14]] to i64* 1303 // CHECK9-NEXT: store i64 [[TMP1]], i64* [[TMP15]], align 8 1304 // CHECK9-NEXT: [[TMP16:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_MAPPERS]], i64 0, i64 2 1305 // CHECK9-NEXT: store i8* null, i8** [[TMP16]], align 8 1306 // CHECK9-NEXT: [[TMP17:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 1307 // CHECK9-NEXT: [[TMP18:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 0 1308 // CHECK9-NEXT: [[TMP19:%.*]] = load i32, i32* getelementptr inbounds ([[STRUCT_SS]], %struct.SS* @Gbla, i32 0, i32 0), align 4 1309 // CHECK9-NEXT: [[B:%.*]] = getelementptr inbounds [[STRUCT_SS]], %struct.SS* [[LA]], i32 0, i32 1 1310 // CHECK9-NEXT: [[TMP20:%.*]] = load float, float* [[B]], align 4 1311 // CHECK9-NEXT: [[CONV1:%.*]] = fptosi float [[TMP20]] to i64 1312 // CHECK9-NEXT: [[TMP21:%.*]] = trunc i64 [[CONV1]] to i32 1313 // CHECK9-NEXT: [[TMP22:%.*]] = call i32 @__tgt_target_teams_mapper(%struct.ident_t* @[[GLOB1:[0-9]+]], i64 -1, i8* @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z18teams_template_argv_l116.region_id, i32 3, i8** [[TMP17]], i8** [[TMP18]], i64* getelementptr inbounds ([3 x i64], [3 x i64]* @.offload_sizes, i32 0, i32 0), i64* getelementptr inbounds ([3 x i64], [3 x i64]* @.offload_maptypes, i32 0, i32 0), i8** null, i8** null, i32 [[TMP19]], i32 [[TMP21]]) 1314 // CHECK9-NEXT: [[TMP23:%.*]] = icmp ne i32 [[TMP22]], 0 1315 // CHECK9-NEXT: br i1 [[TMP23]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]] 1316 // CHECK9: omp_offload.failed: 1317 // CHECK9-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z18teams_template_argv_l116(%struct.SS* @Gbla, %struct.SS* [[LA]], i64 [[TMP1]]) #[[ATTR2:[0-9]+]] 1318 // CHECK9-NEXT: br label [[OMP_OFFLOAD_CONT]] 1319 // CHECK9: omp_offload.cont: 1320 // CHECK9-NEXT: [[TMP24:%.*]] = load i32, i32* [[COMP]], align 4 1321 // CHECK9-NEXT: [[CONV3:%.*]] = bitcast i64* [[COMP_CASTED2]] to i32* 1322 // CHECK9-NEXT: store i32 [[TMP24]], i32* [[CONV3]], align 4 1323 // CHECK9-NEXT: [[TMP25:%.*]] = load i64, i64* [[COMP_CASTED2]], align 8 1324 // CHECK9-NEXT: [[TMP26:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_BASEPTRS4]], i32 0, i32 0 1325 // CHECK9-NEXT: [[TMP27:%.*]] = bitcast i8** [[TMP26]] to %struct.SS.0** 1326 // CHECK9-NEXT: store %struct.SS.0* [[LB]], %struct.SS.0** [[TMP27]], align 8 1327 // CHECK9-NEXT: [[TMP28:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_PTRS5]], i32 0, i32 0 1328 // CHECK9-NEXT: [[TMP29:%.*]] = bitcast i8** [[TMP28]] to %struct.SS.0** 1329 // CHECK9-NEXT: store %struct.SS.0* [[LB]], %struct.SS.0** [[TMP29]], align 8 1330 // CHECK9-NEXT: [[TMP30:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_MAPPERS6]], i64 0, i64 0 1331 // CHECK9-NEXT: store i8* null, i8** [[TMP30]], align 8 1332 // CHECK9-NEXT: [[TMP31:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_BASEPTRS4]], i32 0, i32 1 1333 // CHECK9-NEXT: [[TMP32:%.*]] = bitcast i8** [[TMP31]] to %struct.SS.0** 1334 // CHECK9-NEXT: store %struct.SS.0* @Gblb, %struct.SS.0** [[TMP32]], align 8 1335 // CHECK9-NEXT: [[TMP33:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_PTRS5]], i32 0, i32 1 1336 // CHECK9-NEXT: [[TMP34:%.*]] = bitcast i8** [[TMP33]] to %struct.SS.0** 1337 // CHECK9-NEXT: store %struct.SS.0* @Gblb, %struct.SS.0** [[TMP34]], align 8 1338 // CHECK9-NEXT: [[TMP35:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_MAPPERS6]], i64 0, i64 1 1339 // CHECK9-NEXT: store i8* null, i8** [[TMP35]], align 8 1340 // CHECK9-NEXT: [[TMP36:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_BASEPTRS4]], i32 0, i32 2 1341 // CHECK9-NEXT: [[TMP37:%.*]] = bitcast i8** [[TMP36]] to i64* 1342 // CHECK9-NEXT: store i64 [[TMP25]], i64* [[TMP37]], align 8 1343 // CHECK9-NEXT: [[TMP38:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_PTRS5]], i32 0, i32 2 1344 // CHECK9-NEXT: [[TMP39:%.*]] = bitcast i8** [[TMP38]] to i64* 1345 // CHECK9-NEXT: store i64 [[TMP25]], i64* [[TMP39]], align 8 1346 // CHECK9-NEXT: [[TMP40:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_MAPPERS6]], i64 0, i64 2 1347 // CHECK9-NEXT: store i8* null, i8** [[TMP40]], align 8 1348 // CHECK9-NEXT: [[TMP41:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_BASEPTRS4]], i32 0, i32 0 1349 // CHECK9-NEXT: [[TMP42:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_PTRS5]], i32 0, i32 0 1350 // CHECK9-NEXT: [[B7:%.*]] = getelementptr inbounds [[STRUCT_SS_0]], %struct.SS.0* [[LB]], i32 0, i32 1 1351 // CHECK9-NEXT: [[TMP43:%.*]] = load float, float* [[B7]], align 8 1352 // CHECK9-NEXT: [[CONV8:%.*]] = fptosi float [[TMP43]] to i64 1353 // CHECK9-NEXT: [[TMP44:%.*]] = trunc i64 [[CONV8]] to i32 1354 // CHECK9-NEXT: [[TMP45:%.*]] = load i64, i64* getelementptr inbounds ([[STRUCT_SS_0]], %struct.SS.0* @Gblb, i32 0, i32 0), align 8 1355 // CHECK9-NEXT: [[TMP46:%.*]] = trunc i64 [[TMP45]] to i32 1356 // CHECK9-NEXT: [[TMP47:%.*]] = call i32 @__tgt_target_teams_mapper(%struct.ident_t* @[[GLOB1]], i64 -1, i8* @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z18teams_template_argv_l125.region_id, i32 3, i8** [[TMP41]], i8** [[TMP42]], i64* getelementptr inbounds ([3 x i64], [3 x i64]* @.offload_sizes.2, i32 0, i32 0), i64* getelementptr inbounds ([3 x i64], [3 x i64]* @.offload_maptypes.3, i32 0, i32 0), i8** null, i8** null, i32 [[TMP44]], i32 [[TMP46]]) 1357 // CHECK9-NEXT: [[TMP48:%.*]] = icmp ne i32 [[TMP47]], 0 1358 // CHECK9-NEXT: br i1 [[TMP48]], label [[OMP_OFFLOAD_FAILED9:%.*]], label [[OMP_OFFLOAD_CONT10:%.*]] 1359 // CHECK9: omp_offload.failed9: 1360 // CHECK9-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z18teams_template_argv_l125(%struct.SS.0* [[LB]], %struct.SS.0* @Gblb, i64 [[TMP25]]) #[[ATTR2]] 1361 // CHECK9-NEXT: br label [[OMP_OFFLOAD_CONT10]] 1362 // CHECK9: omp_offload.cont10: 1363 // CHECK9-NEXT: [[TMP49:%.*]] = load i32, i32* [[COMP]], align 4 1364 // CHECK9-NEXT: ret i32 [[TMP49]] 1365 // 1366 // 1367 // CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z18teams_template_argv_l116 1368 // CHECK9-SAME: (%struct.SS* noundef nonnull align 4 dereferenceable(8) [[GBLA:%.*]], %struct.SS* noundef nonnull align 4 dereferenceable(8) [[LA:%.*]], i64 noundef [[COMP:%.*]]) #[[ATTR1:[0-9]+]] { 1369 // CHECK9-NEXT: entry: 1370 // CHECK9-NEXT: [[GBLA_ADDR:%.*]] = alloca %struct.SS*, align 8 1371 // CHECK9-NEXT: [[LA_ADDR:%.*]] = alloca %struct.SS*, align 8 1372 // CHECK9-NEXT: [[COMP_ADDR:%.*]] = alloca i64, align 8 1373 // CHECK9-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1]]) 1374 // CHECK9-NEXT: store %struct.SS* [[GBLA]], %struct.SS** [[GBLA_ADDR]], align 8 1375 // CHECK9-NEXT: store %struct.SS* [[LA]], %struct.SS** [[LA_ADDR]], align 8 1376 // CHECK9-NEXT: store i64 [[COMP]], i64* [[COMP_ADDR]], align 8 1377 // CHECK9-NEXT: [[TMP1:%.*]] = load %struct.SS*, %struct.SS** [[GBLA_ADDR]], align 8 1378 // CHECK9-NEXT: [[TMP2:%.*]] = load %struct.SS*, %struct.SS** [[LA_ADDR]], align 8 1379 // CHECK9-NEXT: [[CONV:%.*]] = bitcast i64* [[COMP_ADDR]] to i32* 1380 // CHECK9-NEXT: [[A:%.*]] = getelementptr inbounds [[STRUCT_SS:%.*]], %struct.SS* [[TMP1]], i32 0, i32 0 1381 // CHECK9-NEXT: [[TMP3:%.*]] = load i32, i32* [[A]], align 4 1382 // CHECK9-NEXT: [[B:%.*]] = getelementptr inbounds [[STRUCT_SS]], %struct.SS* [[TMP2]], i32 0, i32 1 1383 // CHECK9-NEXT: [[TMP4:%.*]] = load float, float* [[B]], align 4 1384 // CHECK9-NEXT: [[CONV1:%.*]] = fptosi float [[TMP4]] to i64 1385 // CHECK9-NEXT: [[TMP5:%.*]] = trunc i64 [[CONV1]] to i32 1386 // CHECK9-NEXT: call void @__kmpc_push_num_teams(%struct.ident_t* @[[GLOB1]], i32 [[TMP0]], i32 [[TMP3]], i32 [[TMP5]]) 1387 // CHECK9-NEXT: call void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(%struct.ident_t* @[[GLOB1]], i32 1, void (i32*, i32*, ...)* bitcast (void (i32*, i32*, i32*)* @.omp_outlined. to void (i32*, i32*, ...)*), i32* [[CONV]]) 1388 // CHECK9-NEXT: ret void 1389 // 1390 // 1391 // CHECK9-LABEL: define {{[^@]+}}@.omp_outlined. 1392 // CHECK9-SAME: (i32* noalias noundef [[DOTGLOBAL_TID_:%.*]], i32* noalias noundef [[DOTBOUND_TID_:%.*]], i32* noundef nonnull align 4 dereferenceable(4) [[COMP:%.*]]) #[[ATTR1]] { 1393 // CHECK9-NEXT: entry: 1394 // CHECK9-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8 1395 // CHECK9-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8 1396 // CHECK9-NEXT: [[COMP_ADDR:%.*]] = alloca i32*, align 8 1397 // CHECK9-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 8 1398 // CHECK9-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 8 1399 // CHECK9-NEXT: store i32* [[COMP]], i32** [[COMP_ADDR]], align 8 1400 // CHECK9-NEXT: [[TMP0:%.*]] = load i32*, i32** [[COMP_ADDR]], align 8 1401 // CHECK9-NEXT: [[TMP1:%.*]] = load i32, i32* [[TMP0]], align 4 1402 // CHECK9-NEXT: [[INC:%.*]] = add nsw i32 [[TMP1]], 1 1403 // CHECK9-NEXT: store i32 [[INC]], i32* [[TMP0]], align 4 1404 // CHECK9-NEXT: ret void 1405 // 1406 // 1407 // CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z18teams_template_argv_l125 1408 // CHECK9-SAME: (%struct.SS.0* noundef nonnull align 8 dereferenceable(16) [[LB:%.*]], %struct.SS.0* noundef nonnull align 8 dereferenceable(16) [[GBLB:%.*]], i64 noundef [[COMP:%.*]]) #[[ATTR1]] { 1409 // CHECK9-NEXT: entry: 1410 // CHECK9-NEXT: [[LB_ADDR:%.*]] = alloca %struct.SS.0*, align 8 1411 // CHECK9-NEXT: [[GBLB_ADDR:%.*]] = alloca %struct.SS.0*, align 8 1412 // CHECK9-NEXT: [[COMP_ADDR:%.*]] = alloca i64, align 8 1413 // CHECK9-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1]]) 1414 // CHECK9-NEXT: store %struct.SS.0* [[LB]], %struct.SS.0** [[LB_ADDR]], align 8 1415 // CHECK9-NEXT: store %struct.SS.0* [[GBLB]], %struct.SS.0** [[GBLB_ADDR]], align 8 1416 // CHECK9-NEXT: store i64 [[COMP]], i64* [[COMP_ADDR]], align 8 1417 // CHECK9-NEXT: [[TMP1:%.*]] = load %struct.SS.0*, %struct.SS.0** [[LB_ADDR]], align 8 1418 // CHECK9-NEXT: [[TMP2:%.*]] = load %struct.SS.0*, %struct.SS.0** [[GBLB_ADDR]], align 8 1419 // CHECK9-NEXT: [[CONV:%.*]] = bitcast i64* [[COMP_ADDR]] to i32* 1420 // CHECK9-NEXT: [[B:%.*]] = getelementptr inbounds [[STRUCT_SS_0:%.*]], %struct.SS.0* [[TMP1]], i32 0, i32 1 1421 // CHECK9-NEXT: [[TMP3:%.*]] = load float, float* [[B]], align 8 1422 // CHECK9-NEXT: [[CONV1:%.*]] = fptosi float [[TMP3]] to i64 1423 // CHECK9-NEXT: [[TMP4:%.*]] = trunc i64 [[CONV1]] to i32 1424 // CHECK9-NEXT: [[A:%.*]] = getelementptr inbounds [[STRUCT_SS_0]], %struct.SS.0* [[TMP2]], i32 0, i32 0 1425 // CHECK9-NEXT: [[TMP5:%.*]] = load i64, i64* [[A]], align 8 1426 // CHECK9-NEXT: [[TMP6:%.*]] = trunc i64 [[TMP5]] to i32 1427 // CHECK9-NEXT: call void @__kmpc_push_num_teams(%struct.ident_t* @[[GLOB1]], i32 [[TMP0]], i32 [[TMP4]], i32 [[TMP6]]) 1428 // CHECK9-NEXT: call void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(%struct.ident_t* @[[GLOB1]], i32 1, void (i32*, i32*, ...)* bitcast (void (i32*, i32*, i32*)* @.omp_outlined..1 to void (i32*, i32*, ...)*), i32* [[CONV]]) 1429 // CHECK9-NEXT: ret void 1430 // 1431 // 1432 // CHECK9-LABEL: define {{[^@]+}}@.omp_outlined..1 1433 // CHECK9-SAME: (i32* noalias noundef [[DOTGLOBAL_TID_:%.*]], i32* noalias noundef [[DOTBOUND_TID_:%.*]], i32* noundef nonnull align 4 dereferenceable(4) [[COMP:%.*]]) #[[ATTR1]] { 1434 // CHECK9-NEXT: entry: 1435 // CHECK9-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8 1436 // CHECK9-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8 1437 // CHECK9-NEXT: [[COMP_ADDR:%.*]] = alloca i32*, align 8 1438 // CHECK9-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 8 1439 // CHECK9-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 8 1440 // CHECK9-NEXT: store i32* [[COMP]], i32** [[COMP_ADDR]], align 8 1441 // CHECK9-NEXT: [[TMP0:%.*]] = load i32*, i32** [[COMP_ADDR]], align 8 1442 // CHECK9-NEXT: [[TMP1:%.*]] = load i32, i32* [[TMP0]], align 4 1443 // CHECK9-NEXT: [[INC:%.*]] = add nsw i32 [[TMP1]], 1 1444 // CHECK9-NEXT: store i32 [[INC]], i32* [[TMP0]], align 4 1445 // CHECK9-NEXT: ret void 1446 // 1447 // 1448 // CHECK9-LABEL: define {{[^@]+}}@.omp_offloading.requires_reg 1449 // CHECK9-SAME: () #[[ATTR3:[0-9]+]] { 1450 // CHECK9-NEXT: entry: 1451 // CHECK9-NEXT: call void @__tgt_register_requires(i64 1) 1452 // CHECK9-NEXT: ret void 1453 // 1454 // 1455 // CHECK11-LABEL: define {{[^@]+}}@_Z18teams_template_argv 1456 // CHECK11-SAME: () #[[ATTR0:[0-9]+]] { 1457 // CHECK11-NEXT: entry: 1458 // CHECK11-NEXT: [[COMP:%.*]] = alloca i32, align 4 1459 // CHECK11-NEXT: [[LA:%.*]] = alloca [[STRUCT_SS:%.*]], align 4 1460 // CHECK11-NEXT: [[LB:%.*]] = alloca [[STRUCT_SS_0:%.*]], align 4 1461 // CHECK11-NEXT: [[COMP_CASTED:%.*]] = alloca i32, align 4 1462 // CHECK11-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [3 x i8*], align 4 1463 // CHECK11-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [3 x i8*], align 4 1464 // CHECK11-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [3 x i8*], align 4 1465 // CHECK11-NEXT: [[COMP_CASTED1:%.*]] = alloca i32, align 4 1466 // CHECK11-NEXT: [[DOTOFFLOAD_BASEPTRS2:%.*]] = alloca [3 x i8*], align 4 1467 // CHECK11-NEXT: [[DOTOFFLOAD_PTRS3:%.*]] = alloca [3 x i8*], align 4 1468 // CHECK11-NEXT: [[DOTOFFLOAD_MAPPERS4:%.*]] = alloca [3 x i8*], align 4 1469 // CHECK11-NEXT: store i32 1, i32* [[COMP]], align 4 1470 // CHECK11-NEXT: [[TMP0:%.*]] = load i32, i32* [[COMP]], align 4 1471 // CHECK11-NEXT: store i32 [[TMP0]], i32* [[COMP_CASTED]], align 4 1472 // CHECK11-NEXT: [[TMP1:%.*]] = load i32, i32* [[COMP_CASTED]], align 4 1473 // CHECK11-NEXT: [[TMP2:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 1474 // CHECK11-NEXT: [[TMP3:%.*]] = bitcast i8** [[TMP2]] to %struct.SS** 1475 // CHECK11-NEXT: store %struct.SS* @Gbla, %struct.SS** [[TMP3]], align 4 1476 // CHECK11-NEXT: [[TMP4:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 0 1477 // CHECK11-NEXT: [[TMP5:%.*]] = bitcast i8** [[TMP4]] to %struct.SS** 1478 // CHECK11-NEXT: store %struct.SS* @Gbla, %struct.SS** [[TMP5]], align 4 1479 // CHECK11-NEXT: [[TMP6:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_MAPPERS]], i32 0, i32 0 1480 // CHECK11-NEXT: store i8* null, i8** [[TMP6]], align 4 1481 // CHECK11-NEXT: [[TMP7:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 1 1482 // CHECK11-NEXT: [[TMP8:%.*]] = bitcast i8** [[TMP7]] to %struct.SS** 1483 // CHECK11-NEXT: store %struct.SS* [[LA]], %struct.SS** [[TMP8]], align 4 1484 // CHECK11-NEXT: [[TMP9:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 1 1485 // CHECK11-NEXT: [[TMP10:%.*]] = bitcast i8** [[TMP9]] to %struct.SS** 1486 // CHECK11-NEXT: store %struct.SS* [[LA]], %struct.SS** [[TMP10]], align 4 1487 // CHECK11-NEXT: [[TMP11:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_MAPPERS]], i32 0, i32 1 1488 // CHECK11-NEXT: store i8* null, i8** [[TMP11]], align 4 1489 // CHECK11-NEXT: [[TMP12:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 2 1490 // CHECK11-NEXT: [[TMP13:%.*]] = bitcast i8** [[TMP12]] to i32* 1491 // CHECK11-NEXT: store i32 [[TMP1]], i32* [[TMP13]], align 4 1492 // CHECK11-NEXT: [[TMP14:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 2 1493 // CHECK11-NEXT: [[TMP15:%.*]] = bitcast i8** [[TMP14]] to i32* 1494 // CHECK11-NEXT: store i32 [[TMP1]], i32* [[TMP15]], align 4 1495 // CHECK11-NEXT: [[TMP16:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_MAPPERS]], i32 0, i32 2 1496 // CHECK11-NEXT: store i8* null, i8** [[TMP16]], align 4 1497 // CHECK11-NEXT: [[TMP17:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 1498 // CHECK11-NEXT: [[TMP18:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 0 1499 // CHECK11-NEXT: [[TMP19:%.*]] = load i32, i32* getelementptr inbounds ([[STRUCT_SS]], %struct.SS* @Gbla, i32 0, i32 0), align 4 1500 // CHECK11-NEXT: [[B:%.*]] = getelementptr inbounds [[STRUCT_SS]], %struct.SS* [[LA]], i32 0, i32 1 1501 // CHECK11-NEXT: [[TMP20:%.*]] = load float, float* [[B]], align 4 1502 // CHECK11-NEXT: [[CONV:%.*]] = fptosi float [[TMP20]] to i64 1503 // CHECK11-NEXT: [[TMP21:%.*]] = trunc i64 [[CONV]] to i32 1504 // CHECK11-NEXT: [[TMP22:%.*]] = call i32 @__tgt_target_teams_mapper(%struct.ident_t* @[[GLOB1:[0-9]+]], i64 -1, i8* @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z18teams_template_argv_l116.region_id, i32 3, i8** [[TMP17]], i8** [[TMP18]], i64* getelementptr inbounds ([3 x i64], [3 x i64]* @.offload_sizes, i32 0, i32 0), i64* getelementptr inbounds ([3 x i64], [3 x i64]* @.offload_maptypes, i32 0, i32 0), i8** null, i8** null, i32 [[TMP19]], i32 [[TMP21]]) 1505 // CHECK11-NEXT: [[TMP23:%.*]] = icmp ne i32 [[TMP22]], 0 1506 // CHECK11-NEXT: br i1 [[TMP23]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]] 1507 // CHECK11: omp_offload.failed: 1508 // CHECK11-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z18teams_template_argv_l116(%struct.SS* @Gbla, %struct.SS* [[LA]], i32 [[TMP1]]) #[[ATTR2:[0-9]+]] 1509 // CHECK11-NEXT: br label [[OMP_OFFLOAD_CONT]] 1510 // CHECK11: omp_offload.cont: 1511 // CHECK11-NEXT: [[TMP24:%.*]] = load i32, i32* [[COMP]], align 4 1512 // CHECK11-NEXT: store i32 [[TMP24]], i32* [[COMP_CASTED1]], align 4 1513 // CHECK11-NEXT: [[TMP25:%.*]] = load i32, i32* [[COMP_CASTED1]], align 4 1514 // CHECK11-NEXT: [[TMP26:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_BASEPTRS2]], i32 0, i32 0 1515 // CHECK11-NEXT: [[TMP27:%.*]] = bitcast i8** [[TMP26]] to %struct.SS.0** 1516 // CHECK11-NEXT: store %struct.SS.0* [[LB]], %struct.SS.0** [[TMP27]], align 4 1517 // CHECK11-NEXT: [[TMP28:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_PTRS3]], i32 0, i32 0 1518 // CHECK11-NEXT: [[TMP29:%.*]] = bitcast i8** [[TMP28]] to %struct.SS.0** 1519 // CHECK11-NEXT: store %struct.SS.0* [[LB]], %struct.SS.0** [[TMP29]], align 4 1520 // CHECK11-NEXT: [[TMP30:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_MAPPERS4]], i32 0, i32 0 1521 // CHECK11-NEXT: store i8* null, i8** [[TMP30]], align 4 1522 // CHECK11-NEXT: [[TMP31:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_BASEPTRS2]], i32 0, i32 1 1523 // CHECK11-NEXT: [[TMP32:%.*]] = bitcast i8** [[TMP31]] to %struct.SS.0** 1524 // CHECK11-NEXT: store %struct.SS.0* @Gblb, %struct.SS.0** [[TMP32]], align 4 1525 // CHECK11-NEXT: [[TMP33:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_PTRS3]], i32 0, i32 1 1526 // CHECK11-NEXT: [[TMP34:%.*]] = bitcast i8** [[TMP33]] to %struct.SS.0** 1527 // CHECK11-NEXT: store %struct.SS.0* @Gblb, %struct.SS.0** [[TMP34]], align 4 1528 // CHECK11-NEXT: [[TMP35:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_MAPPERS4]], i32 0, i32 1 1529 // CHECK11-NEXT: store i8* null, i8** [[TMP35]], align 4 1530 // CHECK11-NEXT: [[TMP36:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_BASEPTRS2]], i32 0, i32 2 1531 // CHECK11-NEXT: [[TMP37:%.*]] = bitcast i8** [[TMP36]] to i32* 1532 // CHECK11-NEXT: store i32 [[TMP25]], i32* [[TMP37]], align 4 1533 // CHECK11-NEXT: [[TMP38:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_PTRS3]], i32 0, i32 2 1534 // CHECK11-NEXT: [[TMP39:%.*]] = bitcast i8** [[TMP38]] to i32* 1535 // CHECK11-NEXT: store i32 [[TMP25]], i32* [[TMP39]], align 4 1536 // CHECK11-NEXT: [[TMP40:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_MAPPERS4]], i32 0, i32 2 1537 // CHECK11-NEXT: store i8* null, i8** [[TMP40]], align 4 1538 // CHECK11-NEXT: [[TMP41:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_BASEPTRS2]], i32 0, i32 0 1539 // CHECK11-NEXT: [[TMP42:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_PTRS3]], i32 0, i32 0 1540 // CHECK11-NEXT: [[B5:%.*]] = getelementptr inbounds [[STRUCT_SS_0]], %struct.SS.0* [[LB]], i32 0, i32 1 1541 // CHECK11-NEXT: [[TMP43:%.*]] = load float, float* [[B5]], align 4 1542 // CHECK11-NEXT: [[CONV6:%.*]] = fptosi float [[TMP43]] to i64 1543 // CHECK11-NEXT: [[TMP44:%.*]] = trunc i64 [[CONV6]] to i32 1544 // CHECK11-NEXT: [[TMP45:%.*]] = load i64, i64* getelementptr inbounds ([[STRUCT_SS_0]], %struct.SS.0* @Gblb, i32 0, i32 0), align 4 1545 // CHECK11-NEXT: [[TMP46:%.*]] = trunc i64 [[TMP45]] to i32 1546 // CHECK11-NEXT: [[TMP47:%.*]] = call i32 @__tgt_target_teams_mapper(%struct.ident_t* @[[GLOB1]], i64 -1, i8* @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z18teams_template_argv_l125.region_id, i32 3, i8** [[TMP41]], i8** [[TMP42]], i64* getelementptr inbounds ([3 x i64], [3 x i64]* @.offload_sizes.2, i32 0, i32 0), i64* getelementptr inbounds ([3 x i64], [3 x i64]* @.offload_maptypes.3, i32 0, i32 0), i8** null, i8** null, i32 [[TMP44]], i32 [[TMP46]]) 1547 // CHECK11-NEXT: [[TMP48:%.*]] = icmp ne i32 [[TMP47]], 0 1548 // CHECK11-NEXT: br i1 [[TMP48]], label [[OMP_OFFLOAD_FAILED7:%.*]], label [[OMP_OFFLOAD_CONT8:%.*]] 1549 // CHECK11: omp_offload.failed7: 1550 // CHECK11-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z18teams_template_argv_l125(%struct.SS.0* [[LB]], %struct.SS.0* @Gblb, i32 [[TMP25]]) #[[ATTR2]] 1551 // CHECK11-NEXT: br label [[OMP_OFFLOAD_CONT8]] 1552 // CHECK11: omp_offload.cont8: 1553 // CHECK11-NEXT: [[TMP49:%.*]] = load i32, i32* [[COMP]], align 4 1554 // CHECK11-NEXT: ret i32 [[TMP49]] 1555 // 1556 // 1557 // CHECK11-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z18teams_template_argv_l116 1558 // CHECK11-SAME: (%struct.SS* noundef nonnull align 4 dereferenceable(8) [[GBLA:%.*]], %struct.SS* noundef nonnull align 4 dereferenceable(8) [[LA:%.*]], i32 noundef [[COMP:%.*]]) #[[ATTR1:[0-9]+]] { 1559 // CHECK11-NEXT: entry: 1560 // CHECK11-NEXT: [[GBLA_ADDR:%.*]] = alloca %struct.SS*, align 4 1561 // CHECK11-NEXT: [[LA_ADDR:%.*]] = alloca %struct.SS*, align 4 1562 // CHECK11-NEXT: [[COMP_ADDR:%.*]] = alloca i32, align 4 1563 // CHECK11-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1]]) 1564 // CHECK11-NEXT: store %struct.SS* [[GBLA]], %struct.SS** [[GBLA_ADDR]], align 4 1565 // CHECK11-NEXT: store %struct.SS* [[LA]], %struct.SS** [[LA_ADDR]], align 4 1566 // CHECK11-NEXT: store i32 [[COMP]], i32* [[COMP_ADDR]], align 4 1567 // CHECK11-NEXT: [[TMP1:%.*]] = load %struct.SS*, %struct.SS** [[GBLA_ADDR]], align 4 1568 // CHECK11-NEXT: [[TMP2:%.*]] = load %struct.SS*, %struct.SS** [[LA_ADDR]], align 4 1569 // CHECK11-NEXT: [[A:%.*]] = getelementptr inbounds [[STRUCT_SS:%.*]], %struct.SS* [[TMP1]], i32 0, i32 0 1570 // CHECK11-NEXT: [[TMP3:%.*]] = load i32, i32* [[A]], align 4 1571 // CHECK11-NEXT: [[B:%.*]] = getelementptr inbounds [[STRUCT_SS]], %struct.SS* [[TMP2]], i32 0, i32 1 1572 // CHECK11-NEXT: [[TMP4:%.*]] = load float, float* [[B]], align 4 1573 // CHECK11-NEXT: [[CONV:%.*]] = fptosi float [[TMP4]] to i64 1574 // CHECK11-NEXT: [[TMP5:%.*]] = trunc i64 [[CONV]] to i32 1575 // CHECK11-NEXT: call void @__kmpc_push_num_teams(%struct.ident_t* @[[GLOB1]], i32 [[TMP0]], i32 [[TMP3]], i32 [[TMP5]]) 1576 // CHECK11-NEXT: call void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(%struct.ident_t* @[[GLOB1]], i32 1, void (i32*, i32*, ...)* bitcast (void (i32*, i32*, i32*)* @.omp_outlined. to void (i32*, i32*, ...)*), i32* [[COMP_ADDR]]) 1577 // CHECK11-NEXT: ret void 1578 // 1579 // 1580 // CHECK11-LABEL: define {{[^@]+}}@.omp_outlined. 1581 // CHECK11-SAME: (i32* noalias noundef [[DOTGLOBAL_TID_:%.*]], i32* noalias noundef [[DOTBOUND_TID_:%.*]], i32* noundef nonnull align 4 dereferenceable(4) [[COMP:%.*]]) #[[ATTR1]] { 1582 // CHECK11-NEXT: entry: 1583 // CHECK11-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 4 1584 // CHECK11-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 4 1585 // CHECK11-NEXT: [[COMP_ADDR:%.*]] = alloca i32*, align 4 1586 // CHECK11-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 4 1587 // CHECK11-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 4 1588 // CHECK11-NEXT: store i32* [[COMP]], i32** [[COMP_ADDR]], align 4 1589 // CHECK11-NEXT: [[TMP0:%.*]] = load i32*, i32** [[COMP_ADDR]], align 4 1590 // CHECK11-NEXT: [[TMP1:%.*]] = load i32, i32* [[TMP0]], align 4 1591 // CHECK11-NEXT: [[INC:%.*]] = add nsw i32 [[TMP1]], 1 1592 // CHECK11-NEXT: store i32 [[INC]], i32* [[TMP0]], align 4 1593 // CHECK11-NEXT: ret void 1594 // 1595 // 1596 // CHECK11-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z18teams_template_argv_l125 1597 // CHECK11-SAME: (%struct.SS.0* noundef nonnull align 4 dereferenceable(12) [[LB:%.*]], %struct.SS.0* noundef nonnull align 4 dereferenceable(12) [[GBLB:%.*]], i32 noundef [[COMP:%.*]]) #[[ATTR1]] { 1598 // CHECK11-NEXT: entry: 1599 // CHECK11-NEXT: [[LB_ADDR:%.*]] = alloca %struct.SS.0*, align 4 1600 // CHECK11-NEXT: [[GBLB_ADDR:%.*]] = alloca %struct.SS.0*, align 4 1601 // CHECK11-NEXT: [[COMP_ADDR:%.*]] = alloca i32, align 4 1602 // CHECK11-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1]]) 1603 // CHECK11-NEXT: store %struct.SS.0* [[LB]], %struct.SS.0** [[LB_ADDR]], align 4 1604 // CHECK11-NEXT: store %struct.SS.0* [[GBLB]], %struct.SS.0** [[GBLB_ADDR]], align 4 1605 // CHECK11-NEXT: store i32 [[COMP]], i32* [[COMP_ADDR]], align 4 1606 // CHECK11-NEXT: [[TMP1:%.*]] = load %struct.SS.0*, %struct.SS.0** [[LB_ADDR]], align 4 1607 // CHECK11-NEXT: [[TMP2:%.*]] = load %struct.SS.0*, %struct.SS.0** [[GBLB_ADDR]], align 4 1608 // CHECK11-NEXT: [[B:%.*]] = getelementptr inbounds [[STRUCT_SS_0:%.*]], %struct.SS.0* [[TMP1]], i32 0, i32 1 1609 // CHECK11-NEXT: [[TMP3:%.*]] = load float, float* [[B]], align 4 1610 // CHECK11-NEXT: [[CONV:%.*]] = fptosi float [[TMP3]] to i64 1611 // CHECK11-NEXT: [[TMP4:%.*]] = trunc i64 [[CONV]] to i32 1612 // CHECK11-NEXT: [[A:%.*]] = getelementptr inbounds [[STRUCT_SS_0]], %struct.SS.0* [[TMP2]], i32 0, i32 0 1613 // CHECK11-NEXT: [[TMP5:%.*]] = load i64, i64* [[A]], align 4 1614 // CHECK11-NEXT: [[TMP6:%.*]] = trunc i64 [[TMP5]] to i32 1615 // CHECK11-NEXT: call void @__kmpc_push_num_teams(%struct.ident_t* @[[GLOB1]], i32 [[TMP0]], i32 [[TMP4]], i32 [[TMP6]]) 1616 // CHECK11-NEXT: call void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(%struct.ident_t* @[[GLOB1]], i32 1, void (i32*, i32*, ...)* bitcast (void (i32*, i32*, i32*)* @.omp_outlined..1 to void (i32*, i32*, ...)*), i32* [[COMP_ADDR]]) 1617 // CHECK11-NEXT: ret void 1618 // 1619 // 1620 // CHECK11-LABEL: define {{[^@]+}}@.omp_outlined..1 1621 // CHECK11-SAME: (i32* noalias noundef [[DOTGLOBAL_TID_:%.*]], i32* noalias noundef [[DOTBOUND_TID_:%.*]], i32* noundef nonnull align 4 dereferenceable(4) [[COMP:%.*]]) #[[ATTR1]] { 1622 // CHECK11-NEXT: entry: 1623 // CHECK11-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 4 1624 // CHECK11-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 4 1625 // CHECK11-NEXT: [[COMP_ADDR:%.*]] = alloca i32*, align 4 1626 // CHECK11-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 4 1627 // CHECK11-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 4 1628 // CHECK11-NEXT: store i32* [[COMP]], i32** [[COMP_ADDR]], align 4 1629 // CHECK11-NEXT: [[TMP0:%.*]] = load i32*, i32** [[COMP_ADDR]], align 4 1630 // CHECK11-NEXT: [[TMP1:%.*]] = load i32, i32* [[TMP0]], align 4 1631 // CHECK11-NEXT: [[INC:%.*]] = add nsw i32 [[TMP1]], 1 1632 // CHECK11-NEXT: store i32 [[INC]], i32* [[TMP0]], align 4 1633 // CHECK11-NEXT: ret void 1634 // 1635 // 1636 // CHECK11-LABEL: define {{[^@]+}}@.omp_offloading.requires_reg 1637 // CHECK11-SAME: () #[[ATTR3:[0-9]+]] { 1638 // CHECK11-NEXT: entry: 1639 // CHECK11-NEXT: call void @__tgt_register_requires(i64 1) 1640 // CHECK11-NEXT: ret void 1641 // 1642 // 1643 // CHECK17-LABEL: define {{[^@]+}}@_Z21teams_template_structv 1644 // CHECK17-SAME: () #[[ATTR0:[0-9]+]] { 1645 // CHECK17-NEXT: entry: 1646 // CHECK17-NEXT: [[V:%.*]] = alloca [[STRUCT_SS:%.*]], align 4 1647 // CHECK17-NEXT: [[CALL:%.*]] = call noundef signext i32 @_ZN2SSIiLi123ELx456EE3fooEv(%struct.SS* noundef nonnull align 4 dereferenceable(8) [[V]]) 1648 // CHECK17-NEXT: ret i32 [[CALL]] 1649 // 1650 // 1651 // CHECK17-LABEL: define {{[^@]+}}@_ZN2SSIiLi123ELx456EE3fooEv 1652 // CHECK17-SAME: (%struct.SS* noundef nonnull align 4 dereferenceable(8) [[THIS:%.*]]) #[[ATTR0]] comdat align 2 { 1653 // CHECK17-NEXT: entry: 1654 // CHECK17-NEXT: [[THIS_ADDR:%.*]] = alloca %struct.SS*, align 8 1655 // CHECK17-NEXT: [[COMP:%.*]] = alloca i32, align 4 1656 // CHECK17-NEXT: [[COMP_CASTED:%.*]] = alloca i64, align 8 1657 // CHECK17-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [2 x i8*], align 8 1658 // CHECK17-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [2 x i8*], align 8 1659 // CHECK17-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [2 x i8*], align 8 1660 // CHECK17-NEXT: [[COMP_CASTED3:%.*]] = alloca i64, align 8 1661 // CHECK17-NEXT: [[DOTOFFLOAD_BASEPTRS5:%.*]] = alloca [2 x i8*], align 8 1662 // CHECK17-NEXT: [[DOTOFFLOAD_PTRS6:%.*]] = alloca [2 x i8*], align 8 1663 // CHECK17-NEXT: [[DOTOFFLOAD_MAPPERS7:%.*]] = alloca [2 x i8*], align 8 1664 // CHECK17-NEXT: store %struct.SS* [[THIS]], %struct.SS** [[THIS_ADDR]], align 8 1665 // CHECK17-NEXT: [[THIS1:%.*]] = load %struct.SS*, %struct.SS** [[THIS_ADDR]], align 8 1666 // CHECK17-NEXT: store i32 1, i32* [[COMP]], align 4 1667 // CHECK17-NEXT: [[TMP0:%.*]] = load i32, i32* [[COMP]], align 4 1668 // CHECK17-NEXT: [[CONV:%.*]] = bitcast i64* [[COMP_CASTED]] to i32* 1669 // CHECK17-NEXT: store i32 [[TMP0]], i32* [[CONV]], align 4 1670 // CHECK17-NEXT: [[TMP1:%.*]] = load i64, i64* [[COMP_CASTED]], align 8 1671 // CHECK17-NEXT: [[A:%.*]] = getelementptr inbounds [[STRUCT_SS:%.*]], %struct.SS* [[THIS1]], i32 0, i32 0 1672 // CHECK17-NEXT: [[TMP2:%.*]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 1673 // CHECK17-NEXT: [[TMP3:%.*]] = bitcast i8** [[TMP2]] to %struct.SS** 1674 // CHECK17-NEXT: store %struct.SS* [[THIS1]], %struct.SS** [[TMP3]], align 8 1675 // CHECK17-NEXT: [[TMP4:%.*]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 0 1676 // CHECK17-NEXT: [[TMP5:%.*]] = bitcast i8** [[TMP4]] to i32** 1677 // CHECK17-NEXT: store i32* [[A]], i32** [[TMP5]], align 8 1678 // CHECK17-NEXT: [[TMP6:%.*]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[DOTOFFLOAD_MAPPERS]], i64 0, i64 0 1679 // CHECK17-NEXT: store i8* null, i8** [[TMP6]], align 8 1680 // CHECK17-NEXT: [[TMP7:%.*]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 1 1681 // CHECK17-NEXT: [[TMP8:%.*]] = bitcast i8** [[TMP7]] to i64* 1682 // CHECK17-NEXT: store i64 [[TMP1]], i64* [[TMP8]], align 8 1683 // CHECK17-NEXT: [[TMP9:%.*]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 1 1684 // CHECK17-NEXT: [[TMP10:%.*]] = bitcast i8** [[TMP9]] to i64* 1685 // CHECK17-NEXT: store i64 [[TMP1]], i64* [[TMP10]], align 8 1686 // CHECK17-NEXT: [[TMP11:%.*]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[DOTOFFLOAD_MAPPERS]], i64 0, i64 1 1687 // CHECK17-NEXT: store i8* null, i8** [[TMP11]], align 8 1688 // CHECK17-NEXT: [[TMP12:%.*]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 1689 // CHECK17-NEXT: [[TMP13:%.*]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 0 1690 // CHECK17-NEXT: [[A2:%.*]] = getelementptr inbounds [[STRUCT_SS]], %struct.SS* [[THIS1]], i32 0, i32 0 1691 // CHECK17-NEXT: [[TMP14:%.*]] = load i32, i32* [[A2]], align 4 1692 // CHECK17-NEXT: [[TMP15:%.*]] = call i32 @__tgt_target_teams_mapper(%struct.ident_t* @[[GLOB1:[0-9]+]], i64 -1, i8* @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l161.region_id, i32 2, i8** [[TMP12]], i8** [[TMP13]], i64* getelementptr inbounds ([2 x i64], [2 x i64]* @.offload_sizes, i32 0, i32 0), i64* getelementptr inbounds ([2 x i64], [2 x i64]* @.offload_maptypes, i32 0, i32 0), i8** null, i8** null, i32 [[TMP14]], i32 123) 1693 // CHECK17-NEXT: [[TMP16:%.*]] = icmp ne i32 [[TMP15]], 0 1694 // CHECK17-NEXT: br i1 [[TMP16]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]] 1695 // CHECK17: omp_offload.failed: 1696 // CHECK17-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l161(%struct.SS* [[THIS1]], i64 [[TMP1]]) #[[ATTR2:[0-9]+]] 1697 // CHECK17-NEXT: br label [[OMP_OFFLOAD_CONT]] 1698 // CHECK17: omp_offload.cont: 1699 // CHECK17-NEXT: [[TMP17:%.*]] = load i32, i32* [[COMP]], align 4 1700 // CHECK17-NEXT: [[CONV4:%.*]] = bitcast i64* [[COMP_CASTED3]] to i32* 1701 // CHECK17-NEXT: store i32 [[TMP17]], i32* [[CONV4]], align 4 1702 // CHECK17-NEXT: [[TMP18:%.*]] = load i64, i64* [[COMP_CASTED3]], align 8 1703 // CHECK17-NEXT: [[B:%.*]] = getelementptr inbounds [[STRUCT_SS]], %struct.SS* [[THIS1]], i32 0, i32 1 1704 // CHECK17-NEXT: [[TMP19:%.*]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[DOTOFFLOAD_BASEPTRS5]], i32 0, i32 0 1705 // CHECK17-NEXT: [[TMP20:%.*]] = bitcast i8** [[TMP19]] to %struct.SS** 1706 // CHECK17-NEXT: store %struct.SS* [[THIS1]], %struct.SS** [[TMP20]], align 8 1707 // CHECK17-NEXT: [[TMP21:%.*]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[DOTOFFLOAD_PTRS6]], i32 0, i32 0 1708 // CHECK17-NEXT: [[TMP22:%.*]] = bitcast i8** [[TMP21]] to float** 1709 // CHECK17-NEXT: store float* [[B]], float** [[TMP22]], align 8 1710 // CHECK17-NEXT: [[TMP23:%.*]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[DOTOFFLOAD_MAPPERS7]], i64 0, i64 0 1711 // CHECK17-NEXT: store i8* null, i8** [[TMP23]], align 8 1712 // CHECK17-NEXT: [[TMP24:%.*]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[DOTOFFLOAD_BASEPTRS5]], i32 0, i32 1 1713 // CHECK17-NEXT: [[TMP25:%.*]] = bitcast i8** [[TMP24]] to i64* 1714 // CHECK17-NEXT: store i64 [[TMP18]], i64* [[TMP25]], align 8 1715 // CHECK17-NEXT: [[TMP26:%.*]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[DOTOFFLOAD_PTRS6]], i32 0, i32 1 1716 // CHECK17-NEXT: [[TMP27:%.*]] = bitcast i8** [[TMP26]] to i64* 1717 // CHECK17-NEXT: store i64 [[TMP18]], i64* [[TMP27]], align 8 1718 // CHECK17-NEXT: [[TMP28:%.*]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[DOTOFFLOAD_MAPPERS7]], i64 0, i64 1 1719 // CHECK17-NEXT: store i8* null, i8** [[TMP28]], align 8 1720 // CHECK17-NEXT: [[TMP29:%.*]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[DOTOFFLOAD_BASEPTRS5]], i32 0, i32 0 1721 // CHECK17-NEXT: [[TMP30:%.*]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[DOTOFFLOAD_PTRS6]], i32 0, i32 0 1722 // CHECK17-NEXT: [[B8:%.*]] = getelementptr inbounds [[STRUCT_SS]], %struct.SS* [[THIS1]], i32 0, i32 1 1723 // CHECK17-NEXT: [[TMP31:%.*]] = load float, float* [[B8]], align 4 1724 // CHECK17-NEXT: [[CONV9:%.*]] = fptosi float [[TMP31]] to i32 1725 // CHECK17-NEXT: [[ADD:%.*]] = add nsw i32 [[CONV9]], 123 1726 // CHECK17-NEXT: [[TMP32:%.*]] = call i32 @__tgt_target_teams_mapper(%struct.ident_t* @[[GLOB1]], i64 -1, i8* @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l169.region_id, i32 2, i8** [[TMP29]], i8** [[TMP30]], i64* getelementptr inbounds ([2 x i64], [2 x i64]* @.offload_sizes.2, i32 0, i32 0), i64* getelementptr inbounds ([2 x i64], [2 x i64]* @.offload_maptypes.3, i32 0, i32 0), i8** null, i8** null, i32 456, i32 [[ADD]]) 1727 // CHECK17-NEXT: [[TMP33:%.*]] = icmp ne i32 [[TMP32]], 0 1728 // CHECK17-NEXT: br i1 [[TMP33]], label [[OMP_OFFLOAD_FAILED10:%.*]], label [[OMP_OFFLOAD_CONT11:%.*]] 1729 // CHECK17: omp_offload.failed10: 1730 // CHECK17-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l169(%struct.SS* [[THIS1]], i64 [[TMP18]]) #[[ATTR2]] 1731 // CHECK17-NEXT: br label [[OMP_OFFLOAD_CONT11]] 1732 // CHECK17: omp_offload.cont11: 1733 // CHECK17-NEXT: [[TMP34:%.*]] = load i32, i32* [[COMP]], align 4 1734 // CHECK17-NEXT: ret i32 [[TMP34]] 1735 // 1736 // 1737 // CHECK17-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l161 1738 // CHECK17-SAME: (%struct.SS* noundef [[THIS:%.*]], i64 noundef [[COMP:%.*]]) #[[ATTR1:[0-9]+]] { 1739 // CHECK17-NEXT: entry: 1740 // CHECK17-NEXT: [[THIS_ADDR:%.*]] = alloca %struct.SS*, align 8 1741 // CHECK17-NEXT: [[COMP_ADDR:%.*]] = alloca i64, align 8 1742 // CHECK17-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1]]) 1743 // CHECK17-NEXT: store %struct.SS* [[THIS]], %struct.SS** [[THIS_ADDR]], align 8 1744 // CHECK17-NEXT: store i64 [[COMP]], i64* [[COMP_ADDR]], align 8 1745 // CHECK17-NEXT: [[TMP1:%.*]] = load %struct.SS*, %struct.SS** [[THIS_ADDR]], align 8 1746 // CHECK17-NEXT: [[CONV:%.*]] = bitcast i64* [[COMP_ADDR]] to i32* 1747 // CHECK17-NEXT: [[A:%.*]] = getelementptr inbounds [[STRUCT_SS:%.*]], %struct.SS* [[TMP1]], i32 0, i32 0 1748 // CHECK17-NEXT: [[TMP2:%.*]] = load i32, i32* [[A]], align 4 1749 // CHECK17-NEXT: call void @__kmpc_push_num_teams(%struct.ident_t* @[[GLOB1]], i32 [[TMP0]], i32 [[TMP2]], i32 123) 1750 // CHECK17-NEXT: call void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(%struct.ident_t* @[[GLOB1]], i32 1, void (i32*, i32*, ...)* bitcast (void (i32*, i32*, i32*)* @.omp_outlined. to void (i32*, i32*, ...)*), i32* [[CONV]]) 1751 // CHECK17-NEXT: ret void 1752 // 1753 // 1754 // CHECK17-LABEL: define {{[^@]+}}@.omp_outlined. 1755 // CHECK17-SAME: (i32* noalias noundef [[DOTGLOBAL_TID_:%.*]], i32* noalias noundef [[DOTBOUND_TID_:%.*]], i32* noundef nonnull align 4 dereferenceable(4) [[COMP:%.*]]) #[[ATTR1]] { 1756 // CHECK17-NEXT: entry: 1757 // CHECK17-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8 1758 // CHECK17-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8 1759 // CHECK17-NEXT: [[COMP_ADDR:%.*]] = alloca i32*, align 8 1760 // CHECK17-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 8 1761 // CHECK17-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 8 1762 // CHECK17-NEXT: store i32* [[COMP]], i32** [[COMP_ADDR]], align 8 1763 // CHECK17-NEXT: [[TMP0:%.*]] = load i32*, i32** [[COMP_ADDR]], align 8 1764 // CHECK17-NEXT: [[TMP1:%.*]] = load i32, i32* [[TMP0]], align 4 1765 // CHECK17-NEXT: [[INC:%.*]] = add nsw i32 [[TMP1]], 1 1766 // CHECK17-NEXT: store i32 [[INC]], i32* [[TMP0]], align 4 1767 // CHECK17-NEXT: ret void 1768 // 1769 // 1770 // CHECK17-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l169 1771 // CHECK17-SAME: (%struct.SS* noundef [[THIS:%.*]], i64 noundef [[COMP:%.*]]) #[[ATTR3:[0-9]+]] { 1772 // CHECK17-NEXT: entry: 1773 // CHECK17-NEXT: [[THIS_ADDR:%.*]] = alloca %struct.SS*, align 8 1774 // CHECK17-NEXT: [[COMP_ADDR:%.*]] = alloca i64, align 8 1775 // CHECK17-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1]]) 1776 // CHECK17-NEXT: store %struct.SS* [[THIS]], %struct.SS** [[THIS_ADDR]], align 8 1777 // CHECK17-NEXT: store i64 [[COMP]], i64* [[COMP_ADDR]], align 8 1778 // CHECK17-NEXT: [[TMP1:%.*]] = load %struct.SS*, %struct.SS** [[THIS_ADDR]], align 8 1779 // CHECK17-NEXT: [[CONV:%.*]] = bitcast i64* [[COMP_ADDR]] to i32* 1780 // CHECK17-NEXT: [[B:%.*]] = getelementptr inbounds [[STRUCT_SS:%.*]], %struct.SS* [[TMP1]], i32 0, i32 1 1781 // CHECK17-NEXT: [[TMP2:%.*]] = load float, float* [[B]], align 4 1782 // CHECK17-NEXT: [[CONV1:%.*]] = fptosi float [[TMP2]] to i32 1783 // CHECK17-NEXT: [[ADD:%.*]] = add nsw i32 [[CONV1]], 123 1784 // CHECK17-NEXT: call void @__kmpc_push_num_teams(%struct.ident_t* @[[GLOB1]], i32 [[TMP0]], i32 456, i32 [[ADD]]) 1785 // CHECK17-NEXT: call void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(%struct.ident_t* @[[GLOB1]], i32 1, void (i32*, i32*, ...)* bitcast (void (i32*, i32*, i32*)* @.omp_outlined..1 to void (i32*, i32*, ...)*), i32* [[CONV]]) 1786 // CHECK17-NEXT: ret void 1787 // 1788 // 1789 // CHECK17-LABEL: define {{[^@]+}}@.omp_outlined..1 1790 // CHECK17-SAME: (i32* noalias noundef [[DOTGLOBAL_TID_:%.*]], i32* noalias noundef [[DOTBOUND_TID_:%.*]], i32* noundef nonnull align 4 dereferenceable(4) [[COMP:%.*]]) #[[ATTR1]] { 1791 // CHECK17-NEXT: entry: 1792 // CHECK17-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8 1793 // CHECK17-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8 1794 // CHECK17-NEXT: [[COMP_ADDR:%.*]] = alloca i32*, align 8 1795 // CHECK17-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 8 1796 // CHECK17-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 8 1797 // CHECK17-NEXT: store i32* [[COMP]], i32** [[COMP_ADDR]], align 8 1798 // CHECK17-NEXT: [[TMP0:%.*]] = load i32*, i32** [[COMP_ADDR]], align 8 1799 // CHECK17-NEXT: [[TMP1:%.*]] = load i32, i32* [[TMP0]], align 4 1800 // CHECK17-NEXT: [[INC:%.*]] = add nsw i32 [[TMP1]], 1 1801 // CHECK17-NEXT: store i32 [[INC]], i32* [[TMP0]], align 4 1802 // CHECK17-NEXT: ret void 1803 // 1804 // 1805 // CHECK17-LABEL: define {{[^@]+}}@.omp_offloading.requires_reg 1806 // CHECK17-SAME: () #[[ATTR4:[0-9]+]] { 1807 // CHECK17-NEXT: entry: 1808 // CHECK17-NEXT: call void @__tgt_register_requires(i64 1) 1809 // CHECK17-NEXT: ret void 1810 // 1811 // 1812 // CHECK19-LABEL: define {{[^@]+}}@_Z21teams_template_structv 1813 // CHECK19-SAME: () #[[ATTR0:[0-9]+]] { 1814 // CHECK19-NEXT: entry: 1815 // CHECK19-NEXT: [[V:%.*]] = alloca [[STRUCT_SS:%.*]], align 4 1816 // CHECK19-NEXT: [[CALL:%.*]] = call noundef i32 @_ZN2SSIiLi123ELx456EE3fooEv(%struct.SS* noundef nonnull align 4 dereferenceable(8) [[V]]) 1817 // CHECK19-NEXT: ret i32 [[CALL]] 1818 // 1819 // 1820 // CHECK19-LABEL: define {{[^@]+}}@_ZN2SSIiLi123ELx456EE3fooEv 1821 // CHECK19-SAME: (%struct.SS* noundef nonnull align 4 dereferenceable(8) [[THIS:%.*]]) #[[ATTR0]] comdat align 2 { 1822 // CHECK19-NEXT: entry: 1823 // CHECK19-NEXT: [[THIS_ADDR:%.*]] = alloca %struct.SS*, align 4 1824 // CHECK19-NEXT: [[COMP:%.*]] = alloca i32, align 4 1825 // CHECK19-NEXT: [[COMP_CASTED:%.*]] = alloca i32, align 4 1826 // CHECK19-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [2 x i8*], align 4 1827 // CHECK19-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [2 x i8*], align 4 1828 // CHECK19-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [2 x i8*], align 4 1829 // CHECK19-NEXT: [[COMP_CASTED3:%.*]] = alloca i32, align 4 1830 // CHECK19-NEXT: [[DOTOFFLOAD_BASEPTRS4:%.*]] = alloca [2 x i8*], align 4 1831 // CHECK19-NEXT: [[DOTOFFLOAD_PTRS5:%.*]] = alloca [2 x i8*], align 4 1832 // CHECK19-NEXT: [[DOTOFFLOAD_MAPPERS6:%.*]] = alloca [2 x i8*], align 4 1833 // CHECK19-NEXT: store %struct.SS* [[THIS]], %struct.SS** [[THIS_ADDR]], align 4 1834 // CHECK19-NEXT: [[THIS1:%.*]] = load %struct.SS*, %struct.SS** [[THIS_ADDR]], align 4 1835 // CHECK19-NEXT: store i32 1, i32* [[COMP]], align 4 1836 // CHECK19-NEXT: [[TMP0:%.*]] = load i32, i32* [[COMP]], align 4 1837 // CHECK19-NEXT: store i32 [[TMP0]], i32* [[COMP_CASTED]], align 4 1838 // CHECK19-NEXT: [[TMP1:%.*]] = load i32, i32* [[COMP_CASTED]], align 4 1839 // CHECK19-NEXT: [[A:%.*]] = getelementptr inbounds [[STRUCT_SS:%.*]], %struct.SS* [[THIS1]], i32 0, i32 0 1840 // CHECK19-NEXT: [[TMP2:%.*]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 1841 // CHECK19-NEXT: [[TMP3:%.*]] = bitcast i8** [[TMP2]] to %struct.SS** 1842 // CHECK19-NEXT: store %struct.SS* [[THIS1]], %struct.SS** [[TMP3]], align 4 1843 // CHECK19-NEXT: [[TMP4:%.*]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 0 1844 // CHECK19-NEXT: [[TMP5:%.*]] = bitcast i8** [[TMP4]] to i32** 1845 // CHECK19-NEXT: store i32* [[A]], i32** [[TMP5]], align 4 1846 // CHECK19-NEXT: [[TMP6:%.*]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[DOTOFFLOAD_MAPPERS]], i32 0, i32 0 1847 // CHECK19-NEXT: store i8* null, i8** [[TMP6]], align 4 1848 // CHECK19-NEXT: [[TMP7:%.*]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 1 1849 // CHECK19-NEXT: [[TMP8:%.*]] = bitcast i8** [[TMP7]] to i32* 1850 // CHECK19-NEXT: store i32 [[TMP1]], i32* [[TMP8]], align 4 1851 // CHECK19-NEXT: [[TMP9:%.*]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 1 1852 // CHECK19-NEXT: [[TMP10:%.*]] = bitcast i8** [[TMP9]] to i32* 1853 // CHECK19-NEXT: store i32 [[TMP1]], i32* [[TMP10]], align 4 1854 // CHECK19-NEXT: [[TMP11:%.*]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[DOTOFFLOAD_MAPPERS]], i32 0, i32 1 1855 // CHECK19-NEXT: store i8* null, i8** [[TMP11]], align 4 1856 // CHECK19-NEXT: [[TMP12:%.*]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 1857 // CHECK19-NEXT: [[TMP13:%.*]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 0 1858 // CHECK19-NEXT: [[A2:%.*]] = getelementptr inbounds [[STRUCT_SS]], %struct.SS* [[THIS1]], i32 0, i32 0 1859 // CHECK19-NEXT: [[TMP14:%.*]] = load i32, i32* [[A2]], align 4 1860 // CHECK19-NEXT: [[TMP15:%.*]] = call i32 @__tgt_target_teams_mapper(%struct.ident_t* @[[GLOB1:[0-9]+]], i64 -1, i8* @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l161.region_id, i32 2, i8** [[TMP12]], i8** [[TMP13]], i64* getelementptr inbounds ([2 x i64], [2 x i64]* @.offload_sizes, i32 0, i32 0), i64* getelementptr inbounds ([2 x i64], [2 x i64]* @.offload_maptypes, i32 0, i32 0), i8** null, i8** null, i32 [[TMP14]], i32 123) 1861 // CHECK19-NEXT: [[TMP16:%.*]] = icmp ne i32 [[TMP15]], 0 1862 // CHECK19-NEXT: br i1 [[TMP16]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]] 1863 // CHECK19: omp_offload.failed: 1864 // CHECK19-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l161(%struct.SS* [[THIS1]], i32 [[TMP1]]) #[[ATTR2:[0-9]+]] 1865 // CHECK19-NEXT: br label [[OMP_OFFLOAD_CONT]] 1866 // CHECK19: omp_offload.cont: 1867 // CHECK19-NEXT: [[TMP17:%.*]] = load i32, i32* [[COMP]], align 4 1868 // CHECK19-NEXT: store i32 [[TMP17]], i32* [[COMP_CASTED3]], align 4 1869 // CHECK19-NEXT: [[TMP18:%.*]] = load i32, i32* [[COMP_CASTED3]], align 4 1870 // CHECK19-NEXT: [[B:%.*]] = getelementptr inbounds [[STRUCT_SS]], %struct.SS* [[THIS1]], i32 0, i32 1 1871 // CHECK19-NEXT: [[TMP19:%.*]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[DOTOFFLOAD_BASEPTRS4]], i32 0, i32 0 1872 // CHECK19-NEXT: [[TMP20:%.*]] = bitcast i8** [[TMP19]] to %struct.SS** 1873 // CHECK19-NEXT: store %struct.SS* [[THIS1]], %struct.SS** [[TMP20]], align 4 1874 // CHECK19-NEXT: [[TMP21:%.*]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[DOTOFFLOAD_PTRS5]], i32 0, i32 0 1875 // CHECK19-NEXT: [[TMP22:%.*]] = bitcast i8** [[TMP21]] to float** 1876 // CHECK19-NEXT: store float* [[B]], float** [[TMP22]], align 4 1877 // CHECK19-NEXT: [[TMP23:%.*]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[DOTOFFLOAD_MAPPERS6]], i32 0, i32 0 1878 // CHECK19-NEXT: store i8* null, i8** [[TMP23]], align 4 1879 // CHECK19-NEXT: [[TMP24:%.*]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[DOTOFFLOAD_BASEPTRS4]], i32 0, i32 1 1880 // CHECK19-NEXT: [[TMP25:%.*]] = bitcast i8** [[TMP24]] to i32* 1881 // CHECK19-NEXT: store i32 [[TMP18]], i32* [[TMP25]], align 4 1882 // CHECK19-NEXT: [[TMP26:%.*]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[DOTOFFLOAD_PTRS5]], i32 0, i32 1 1883 // CHECK19-NEXT: [[TMP27:%.*]] = bitcast i8** [[TMP26]] to i32* 1884 // CHECK19-NEXT: store i32 [[TMP18]], i32* [[TMP27]], align 4 1885 // CHECK19-NEXT: [[TMP28:%.*]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[DOTOFFLOAD_MAPPERS6]], i32 0, i32 1 1886 // CHECK19-NEXT: store i8* null, i8** [[TMP28]], align 4 1887 // CHECK19-NEXT: [[TMP29:%.*]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[DOTOFFLOAD_BASEPTRS4]], i32 0, i32 0 1888 // CHECK19-NEXT: [[TMP30:%.*]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[DOTOFFLOAD_PTRS5]], i32 0, i32 0 1889 // CHECK19-NEXT: [[B7:%.*]] = getelementptr inbounds [[STRUCT_SS]], %struct.SS* [[THIS1]], i32 0, i32 1 1890 // CHECK19-NEXT: [[TMP31:%.*]] = load float, float* [[B7]], align 4 1891 // CHECK19-NEXT: [[CONV:%.*]] = fptosi float [[TMP31]] to i32 1892 // CHECK19-NEXT: [[ADD:%.*]] = add nsw i32 [[CONV]], 123 1893 // CHECK19-NEXT: [[TMP32:%.*]] = call i32 @__tgt_target_teams_mapper(%struct.ident_t* @[[GLOB1]], i64 -1, i8* @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l169.region_id, i32 2, i8** [[TMP29]], i8** [[TMP30]], i64* getelementptr inbounds ([2 x i64], [2 x i64]* @.offload_sizes.2, i32 0, i32 0), i64* getelementptr inbounds ([2 x i64], [2 x i64]* @.offload_maptypes.3, i32 0, i32 0), i8** null, i8** null, i32 456, i32 [[ADD]]) 1894 // CHECK19-NEXT: [[TMP33:%.*]] = icmp ne i32 [[TMP32]], 0 1895 // CHECK19-NEXT: br i1 [[TMP33]], label [[OMP_OFFLOAD_FAILED8:%.*]], label [[OMP_OFFLOAD_CONT9:%.*]] 1896 // CHECK19: omp_offload.failed8: 1897 // CHECK19-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l169(%struct.SS* [[THIS1]], i32 [[TMP18]]) #[[ATTR2]] 1898 // CHECK19-NEXT: br label [[OMP_OFFLOAD_CONT9]] 1899 // CHECK19: omp_offload.cont9: 1900 // CHECK19-NEXT: [[TMP34:%.*]] = load i32, i32* [[COMP]], align 4 1901 // CHECK19-NEXT: ret i32 [[TMP34]] 1902 // 1903 // 1904 // CHECK19-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l161 1905 // CHECK19-SAME: (%struct.SS* noundef [[THIS:%.*]], i32 noundef [[COMP:%.*]]) #[[ATTR1:[0-9]+]] { 1906 // CHECK19-NEXT: entry: 1907 // CHECK19-NEXT: [[THIS_ADDR:%.*]] = alloca %struct.SS*, align 4 1908 // CHECK19-NEXT: [[COMP_ADDR:%.*]] = alloca i32, align 4 1909 // CHECK19-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1]]) 1910 // CHECK19-NEXT: store %struct.SS* [[THIS]], %struct.SS** [[THIS_ADDR]], align 4 1911 // CHECK19-NEXT: store i32 [[COMP]], i32* [[COMP_ADDR]], align 4 1912 // CHECK19-NEXT: [[TMP1:%.*]] = load %struct.SS*, %struct.SS** [[THIS_ADDR]], align 4 1913 // CHECK19-NEXT: [[A:%.*]] = getelementptr inbounds [[STRUCT_SS:%.*]], %struct.SS* [[TMP1]], i32 0, i32 0 1914 // CHECK19-NEXT: [[TMP2:%.*]] = load i32, i32* [[A]], align 4 1915 // CHECK19-NEXT: call void @__kmpc_push_num_teams(%struct.ident_t* @[[GLOB1]], i32 [[TMP0]], i32 [[TMP2]], i32 123) 1916 // CHECK19-NEXT: call void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(%struct.ident_t* @[[GLOB1]], i32 1, void (i32*, i32*, ...)* bitcast (void (i32*, i32*, i32*)* @.omp_outlined. to void (i32*, i32*, ...)*), i32* [[COMP_ADDR]]) 1917 // CHECK19-NEXT: ret void 1918 // 1919 // 1920 // CHECK19-LABEL: define {{[^@]+}}@.omp_outlined. 1921 // CHECK19-SAME: (i32* noalias noundef [[DOTGLOBAL_TID_:%.*]], i32* noalias noundef [[DOTBOUND_TID_:%.*]], i32* noundef nonnull align 4 dereferenceable(4) [[COMP:%.*]]) #[[ATTR1]] { 1922 // CHECK19-NEXT: entry: 1923 // CHECK19-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 4 1924 // CHECK19-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 4 1925 // CHECK19-NEXT: [[COMP_ADDR:%.*]] = alloca i32*, align 4 1926 // CHECK19-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 4 1927 // CHECK19-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 4 1928 // CHECK19-NEXT: store i32* [[COMP]], i32** [[COMP_ADDR]], align 4 1929 // CHECK19-NEXT: [[TMP0:%.*]] = load i32*, i32** [[COMP_ADDR]], align 4 1930 // CHECK19-NEXT: [[TMP1:%.*]] = load i32, i32* [[TMP0]], align 4 1931 // CHECK19-NEXT: [[INC:%.*]] = add nsw i32 [[TMP1]], 1 1932 // CHECK19-NEXT: store i32 [[INC]], i32* [[TMP0]], align 4 1933 // CHECK19-NEXT: ret void 1934 // 1935 // 1936 // CHECK19-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l169 1937 // CHECK19-SAME: (%struct.SS* noundef [[THIS:%.*]], i32 noundef [[COMP:%.*]]) #[[ATTR3:[0-9]+]] { 1938 // CHECK19-NEXT: entry: 1939 // CHECK19-NEXT: [[THIS_ADDR:%.*]] = alloca %struct.SS*, align 4 1940 // CHECK19-NEXT: [[COMP_ADDR:%.*]] = alloca i32, align 4 1941 // CHECK19-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1]]) 1942 // CHECK19-NEXT: store %struct.SS* [[THIS]], %struct.SS** [[THIS_ADDR]], align 4 1943 // CHECK19-NEXT: store i32 [[COMP]], i32* [[COMP_ADDR]], align 4 1944 // CHECK19-NEXT: [[TMP1:%.*]] = load %struct.SS*, %struct.SS** [[THIS_ADDR]], align 4 1945 // CHECK19-NEXT: [[B:%.*]] = getelementptr inbounds [[STRUCT_SS:%.*]], %struct.SS* [[TMP1]], i32 0, i32 1 1946 // CHECK19-NEXT: [[TMP2:%.*]] = load float, float* [[B]], align 4 1947 // CHECK19-NEXT: [[CONV:%.*]] = fptosi float [[TMP2]] to i32 1948 // CHECK19-NEXT: [[ADD:%.*]] = add nsw i32 [[CONV]], 123 1949 // CHECK19-NEXT: call void @__kmpc_push_num_teams(%struct.ident_t* @[[GLOB1]], i32 [[TMP0]], i32 456, i32 [[ADD]]) 1950 // CHECK19-NEXT: call void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(%struct.ident_t* @[[GLOB1]], i32 1, void (i32*, i32*, ...)* bitcast (void (i32*, i32*, i32*)* @.omp_outlined..1 to void (i32*, i32*, ...)*), i32* [[COMP_ADDR]]) 1951 // CHECK19-NEXT: ret void 1952 // 1953 // 1954 // CHECK19-LABEL: define {{[^@]+}}@.omp_outlined..1 1955 // CHECK19-SAME: (i32* noalias noundef [[DOTGLOBAL_TID_:%.*]], i32* noalias noundef [[DOTBOUND_TID_:%.*]], i32* noundef nonnull align 4 dereferenceable(4) [[COMP:%.*]]) #[[ATTR1]] { 1956 // CHECK19-NEXT: entry: 1957 // CHECK19-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 4 1958 // CHECK19-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 4 1959 // CHECK19-NEXT: [[COMP_ADDR:%.*]] = alloca i32*, align 4 1960 // CHECK19-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 4 1961 // CHECK19-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 4 1962 // CHECK19-NEXT: store i32* [[COMP]], i32** [[COMP_ADDR]], align 4 1963 // CHECK19-NEXT: [[TMP0:%.*]] = load i32*, i32** [[COMP_ADDR]], align 4 1964 // CHECK19-NEXT: [[TMP1:%.*]] = load i32, i32* [[TMP0]], align 4 1965 // CHECK19-NEXT: [[INC:%.*]] = add nsw i32 [[TMP1]], 1 1966 // CHECK19-NEXT: store i32 [[INC]], i32* [[TMP0]], align 4 1967 // CHECK19-NEXT: ret void 1968 // 1969 // 1970 // CHECK19-LABEL: define {{[^@]+}}@.omp_offloading.requires_reg 1971 // CHECK19-SAME: () #[[ATTR4:[0-9]+]] { 1972 // CHECK19-NEXT: entry: 1973 // CHECK19-NEXT: call void @__tgt_register_requires(i64 1) 1974 // CHECK19-NEXT: ret void 1975 // 1976 // 1977 // CHECK25-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l216 1978 // CHECK25-SAME: (i64 noundef [[ARGC:%.*]]) #[[ATTR0:[0-9]+]] { 1979 // CHECK25-NEXT: entry: 1980 // CHECK25-NEXT: [[ARGC_ADDR:%.*]] = alloca i64, align 8 1981 // CHECK25-NEXT: store i64 [[ARGC]], i64* [[ARGC_ADDR]], align 8 1982 // CHECK25-NEXT: [[CONV:%.*]] = bitcast i64* [[ARGC_ADDR]] to i32* 1983 // CHECK25-NEXT: call void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(%struct.ident_t* @[[GLOB1:[0-9]+]], i32 1, void (i32*, i32*, ...)* bitcast (void (i32*, i32*, i32*)* @.omp_outlined. to void (i32*, i32*, ...)*), i32* [[CONV]]) 1984 // CHECK25-NEXT: ret void 1985 // 1986 // 1987 // CHECK25-LABEL: define {{[^@]+}}@.omp_outlined. 1988 // CHECK25-SAME: (i32* noalias noundef [[DOTGLOBAL_TID_:%.*]], i32* noalias noundef [[DOTBOUND_TID_:%.*]], i32* noundef nonnull align 4 dereferenceable(4) [[ARGC:%.*]]) #[[ATTR0]] { 1989 // CHECK25-NEXT: entry: 1990 // CHECK25-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8 1991 // CHECK25-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8 1992 // CHECK25-NEXT: [[ARGC_ADDR:%.*]] = alloca i32*, align 8 1993 // CHECK25-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 8 1994 // CHECK25-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 8 1995 // CHECK25-NEXT: store i32* [[ARGC]], i32** [[ARGC_ADDR]], align 8 1996 // CHECK25-NEXT: [[TMP0:%.*]] = load i32*, i32** [[ARGC_ADDR]], align 8 1997 // CHECK25-NEXT: store i32 0, i32* [[TMP0]], align 4 1998 // CHECK25-NEXT: ret void 1999 // 2000 // 2001 // CHECK25-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIPPcEiT__l209 2002 // CHECK25-SAME: (i8** noundef [[ARGC:%.*]]) #[[ATTR0]] { 2003 // CHECK25-NEXT: entry: 2004 // CHECK25-NEXT: [[ARGC_ADDR:%.*]] = alloca i8**, align 8 2005 // CHECK25-NEXT: store i8** [[ARGC]], i8*** [[ARGC_ADDR]], align 8 2006 // CHECK25-NEXT: call void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(%struct.ident_t* @[[GLOB1]], i32 1, void (i32*, i32*, ...)* bitcast (void (i32*, i32*, i8***)* @.omp_outlined..1 to void (i32*, i32*, ...)*), i8*** [[ARGC_ADDR]]) 2007 // CHECK25-NEXT: ret void 2008 // 2009 // 2010 // CHECK25-LABEL: define {{[^@]+}}@.omp_outlined..1 2011 // CHECK25-SAME: (i32* noalias noundef [[DOTGLOBAL_TID_:%.*]], i32* noalias noundef [[DOTBOUND_TID_:%.*]], i8*** noundef nonnull align 8 dereferenceable(8) [[ARGC:%.*]]) #[[ATTR0]] { 2012 // CHECK25-NEXT: entry: 2013 // CHECK25-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8 2014 // CHECK25-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8 2015 // CHECK25-NEXT: [[ARGC_ADDR:%.*]] = alloca i8***, align 8 2016 // CHECK25-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 8 2017 // CHECK25-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 8 2018 // CHECK25-NEXT: store i8*** [[ARGC]], i8**** [[ARGC_ADDR]], align 8 2019 // CHECK25-NEXT: [[TMP0:%.*]] = load i8***, i8**** [[ARGC_ADDR]], align 8 2020 // CHECK25-NEXT: store i8** null, i8*** [[TMP0]], align 8 2021 // CHECK25-NEXT: ret void 2022 // 2023 // 2024 // CHECK27-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l216 2025 // CHECK27-SAME: (i32 noundef [[ARGC:%.*]]) #[[ATTR0:[0-9]+]] { 2026 // CHECK27-NEXT: entry: 2027 // CHECK27-NEXT: [[ARGC_ADDR:%.*]] = alloca i32, align 4 2028 // CHECK27-NEXT: store i32 [[ARGC]], i32* [[ARGC_ADDR]], align 4 2029 // CHECK27-NEXT: call void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(%struct.ident_t* @[[GLOB1:[0-9]+]], i32 1, void (i32*, i32*, ...)* bitcast (void (i32*, i32*, i32*)* @.omp_outlined. to void (i32*, i32*, ...)*), i32* [[ARGC_ADDR]]) 2030 // CHECK27-NEXT: ret void 2031 // 2032 // 2033 // CHECK27-LABEL: define {{[^@]+}}@.omp_outlined. 2034 // CHECK27-SAME: (i32* noalias noundef [[DOTGLOBAL_TID_:%.*]], i32* noalias noundef [[DOTBOUND_TID_:%.*]], i32* noundef nonnull align 4 dereferenceable(4) [[ARGC:%.*]]) #[[ATTR0]] { 2035 // CHECK27-NEXT: entry: 2036 // CHECK27-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 4 2037 // CHECK27-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 4 2038 // CHECK27-NEXT: [[ARGC_ADDR:%.*]] = alloca i32*, align 4 2039 // CHECK27-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 4 2040 // CHECK27-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 4 2041 // CHECK27-NEXT: store i32* [[ARGC]], i32** [[ARGC_ADDR]], align 4 2042 // CHECK27-NEXT: [[TMP0:%.*]] = load i32*, i32** [[ARGC_ADDR]], align 4 2043 // CHECK27-NEXT: store i32 0, i32* [[TMP0]], align 4 2044 // CHECK27-NEXT: ret void 2045 // 2046 // 2047 // CHECK27-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIPPcEiT__l209 2048 // CHECK27-SAME: (i8** noundef [[ARGC:%.*]]) #[[ATTR0]] { 2049 // CHECK27-NEXT: entry: 2050 // CHECK27-NEXT: [[ARGC_ADDR:%.*]] = alloca i8**, align 4 2051 // CHECK27-NEXT: store i8** [[ARGC]], i8*** [[ARGC_ADDR]], align 4 2052 // CHECK27-NEXT: call void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(%struct.ident_t* @[[GLOB1]], i32 1, void (i32*, i32*, ...)* bitcast (void (i32*, i32*, i8***)* @.omp_outlined..1 to void (i32*, i32*, ...)*), i8*** [[ARGC_ADDR]]) 2053 // CHECK27-NEXT: ret void 2054 // 2055 // 2056 // CHECK27-LABEL: define {{[^@]+}}@.omp_outlined..1 2057 // CHECK27-SAME: (i32* noalias noundef [[DOTGLOBAL_TID_:%.*]], i32* noalias noundef [[DOTBOUND_TID_:%.*]], i8*** noundef nonnull align 4 dereferenceable(4) [[ARGC:%.*]]) #[[ATTR0]] { 2058 // CHECK27-NEXT: entry: 2059 // CHECK27-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 4 2060 // CHECK27-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 4 2061 // CHECK27-NEXT: [[ARGC_ADDR:%.*]] = alloca i8***, align 4 2062 // CHECK27-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 4 2063 // CHECK27-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 4 2064 // CHECK27-NEXT: store i8*** [[ARGC]], i8**** [[ARGC_ADDR]], align 4 2065 // CHECK27-NEXT: [[TMP0:%.*]] = load i8***, i8**** [[ARGC_ADDR]], align 4 2066 // CHECK27-NEXT: store i8** null, i8*** [[TMP0]], align 4 2067 // CHECK27-NEXT: ret void 2068 // 2069 // 2070 // CHECK33-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l265 2071 // CHECK33-SAME: (i64 noundef [[A:%.*]], i64 noundef [[B:%.*]], i64 noundef [[ARGC:%.*]]) #[[ATTR0:[0-9]+]] { 2072 // CHECK33-NEXT: entry: 2073 // CHECK33-NEXT: [[A_ADDR:%.*]] = alloca i64, align 8 2074 // CHECK33-NEXT: [[B_ADDR:%.*]] = alloca i64, align 8 2075 // CHECK33-NEXT: [[ARGC_ADDR:%.*]] = alloca i64, align 8 2076 // CHECK33-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1:[0-9]+]]) 2077 // CHECK33-NEXT: store i64 [[A]], i64* [[A_ADDR]], align 8 2078 // CHECK33-NEXT: store i64 [[B]], i64* [[B_ADDR]], align 8 2079 // CHECK33-NEXT: store i64 [[ARGC]], i64* [[ARGC_ADDR]], align 8 2080 // CHECK33-NEXT: [[CONV:%.*]] = bitcast i64* [[A_ADDR]] to i32* 2081 // CHECK33-NEXT: [[CONV1:%.*]] = bitcast i64* [[B_ADDR]] to i32* 2082 // CHECK33-NEXT: [[CONV2:%.*]] = bitcast i64* [[ARGC_ADDR]] to i32* 2083 // CHECK33-NEXT: [[TMP1:%.*]] = load i32, i32* [[CONV]], align 4 2084 // CHECK33-NEXT: [[TMP2:%.*]] = load i32, i32* [[CONV1]], align 4 2085 // CHECK33-NEXT: call void @__kmpc_push_num_teams(%struct.ident_t* @[[GLOB1]], i32 [[TMP0]], i32 [[TMP1]], i32 [[TMP2]]) 2086 // CHECK33-NEXT: call void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(%struct.ident_t* @[[GLOB1]], i32 1, void (i32*, i32*, ...)* bitcast (void (i32*, i32*, i32*)* @.omp_outlined. to void (i32*, i32*, ...)*), i32* [[CONV2]]) 2087 // CHECK33-NEXT: ret void 2088 // 2089 // 2090 // CHECK33-LABEL: define {{[^@]+}}@.omp_outlined. 2091 // CHECK33-SAME: (i32* noalias noundef [[DOTGLOBAL_TID_:%.*]], i32* noalias noundef [[DOTBOUND_TID_:%.*]], i32* noundef nonnull align 4 dereferenceable(4) [[ARGC:%.*]]) #[[ATTR0]] { 2092 // CHECK33-NEXT: entry: 2093 // CHECK33-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8 2094 // CHECK33-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8 2095 // CHECK33-NEXT: [[ARGC_ADDR:%.*]] = alloca i32*, align 8 2096 // CHECK33-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 8 2097 // CHECK33-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 8 2098 // CHECK33-NEXT: store i32* [[ARGC]], i32** [[ARGC_ADDR]], align 8 2099 // CHECK33-NEXT: [[TMP0:%.*]] = load i32*, i32** [[ARGC_ADDR]], align 8 2100 // CHECK33-NEXT: store i32 0, i32* [[TMP0]], align 4 2101 // CHECK33-NEXT: ret void 2102 // 2103 // 2104 // CHECK33-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIPPcEiT__l254 2105 // CHECK33-SAME: (i64 noundef [[A:%.*]], i64 noundef [[B:%.*]], i8** noundef [[ARGC:%.*]]) #[[ATTR0]] { 2106 // CHECK33-NEXT: entry: 2107 // CHECK33-NEXT: [[A_ADDR:%.*]] = alloca i64, align 8 2108 // CHECK33-NEXT: [[B_ADDR:%.*]] = alloca i64, align 8 2109 // CHECK33-NEXT: [[ARGC_ADDR:%.*]] = alloca i8**, align 8 2110 // CHECK33-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1]]) 2111 // CHECK33-NEXT: store i64 [[A]], i64* [[A_ADDR]], align 8 2112 // CHECK33-NEXT: store i64 [[B]], i64* [[B_ADDR]], align 8 2113 // CHECK33-NEXT: store i8** [[ARGC]], i8*** [[ARGC_ADDR]], align 8 2114 // CHECK33-NEXT: [[CONV:%.*]] = bitcast i64* [[A_ADDR]] to i32* 2115 // CHECK33-NEXT: [[CONV1:%.*]] = bitcast i64* [[B_ADDR]] to i32* 2116 // CHECK33-NEXT: [[TMP1:%.*]] = load i32, i32* [[CONV]], align 4 2117 // CHECK33-NEXT: [[TMP2:%.*]] = load i32, i32* [[CONV1]], align 4 2118 // CHECK33-NEXT: call void @__kmpc_push_num_teams(%struct.ident_t* @[[GLOB1]], i32 [[TMP0]], i32 [[TMP1]], i32 [[TMP2]]) 2119 // CHECK33-NEXT: call void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(%struct.ident_t* @[[GLOB1]], i32 1, void (i32*, i32*, ...)* bitcast (void (i32*, i32*, i8***)* @.omp_outlined..1 to void (i32*, i32*, ...)*), i8*** [[ARGC_ADDR]]) 2120 // CHECK33-NEXT: ret void 2121 // 2122 // 2123 // CHECK33-LABEL: define {{[^@]+}}@.omp_outlined..1 2124 // CHECK33-SAME: (i32* noalias noundef [[DOTGLOBAL_TID_:%.*]], i32* noalias noundef [[DOTBOUND_TID_:%.*]], i8*** noundef nonnull align 8 dereferenceable(8) [[ARGC:%.*]]) #[[ATTR0]] { 2125 // CHECK33-NEXT: entry: 2126 // CHECK33-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8 2127 // CHECK33-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8 2128 // CHECK33-NEXT: [[ARGC_ADDR:%.*]] = alloca i8***, align 8 2129 // CHECK33-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 8 2130 // CHECK33-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 8 2131 // CHECK33-NEXT: store i8*** [[ARGC]], i8**** [[ARGC_ADDR]], align 8 2132 // CHECK33-NEXT: [[TMP0:%.*]] = load i8***, i8**** [[ARGC_ADDR]], align 8 2133 // CHECK33-NEXT: store i8** null, i8*** [[TMP0]], align 8 2134 // CHECK33-NEXT: ret void 2135 // 2136 // 2137 // CHECK35-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l265 2138 // CHECK35-SAME: (i32 noundef [[A:%.*]], i32 noundef [[B:%.*]], i32 noundef [[ARGC:%.*]]) #[[ATTR0:[0-9]+]] { 2139 // CHECK35-NEXT: entry: 2140 // CHECK35-NEXT: [[A_ADDR:%.*]] = alloca i32, align 4 2141 // CHECK35-NEXT: [[B_ADDR:%.*]] = alloca i32, align 4 2142 // CHECK35-NEXT: [[ARGC_ADDR:%.*]] = alloca i32, align 4 2143 // CHECK35-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1:[0-9]+]]) 2144 // CHECK35-NEXT: store i32 [[A]], i32* [[A_ADDR]], align 4 2145 // CHECK35-NEXT: store i32 [[B]], i32* [[B_ADDR]], align 4 2146 // CHECK35-NEXT: store i32 [[ARGC]], i32* [[ARGC_ADDR]], align 4 2147 // CHECK35-NEXT: [[TMP1:%.*]] = load i32, i32* [[A_ADDR]], align 4 2148 // CHECK35-NEXT: [[TMP2:%.*]] = load i32, i32* [[B_ADDR]], align 4 2149 // CHECK35-NEXT: call void @__kmpc_push_num_teams(%struct.ident_t* @[[GLOB1]], i32 [[TMP0]], i32 [[TMP1]], i32 [[TMP2]]) 2150 // CHECK35-NEXT: call void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(%struct.ident_t* @[[GLOB1]], i32 1, void (i32*, i32*, ...)* bitcast (void (i32*, i32*, i32*)* @.omp_outlined. to void (i32*, i32*, ...)*), i32* [[ARGC_ADDR]]) 2151 // CHECK35-NEXT: ret void 2152 // 2153 // 2154 // CHECK35-LABEL: define {{[^@]+}}@.omp_outlined. 2155 // CHECK35-SAME: (i32* noalias noundef [[DOTGLOBAL_TID_:%.*]], i32* noalias noundef [[DOTBOUND_TID_:%.*]], i32* noundef nonnull align 4 dereferenceable(4) [[ARGC:%.*]]) #[[ATTR0]] { 2156 // CHECK35-NEXT: entry: 2157 // CHECK35-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 4 2158 // CHECK35-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 4 2159 // CHECK35-NEXT: [[ARGC_ADDR:%.*]] = alloca i32*, align 4 2160 // CHECK35-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 4 2161 // CHECK35-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 4 2162 // CHECK35-NEXT: store i32* [[ARGC]], i32** [[ARGC_ADDR]], align 4 2163 // CHECK35-NEXT: [[TMP0:%.*]] = load i32*, i32** [[ARGC_ADDR]], align 4 2164 // CHECK35-NEXT: store i32 0, i32* [[TMP0]], align 4 2165 // CHECK35-NEXT: ret void 2166 // 2167 // 2168 // CHECK35-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIPPcEiT__l254 2169 // CHECK35-SAME: (i32 noundef [[A:%.*]], i32 noundef [[B:%.*]], i8** noundef [[ARGC:%.*]]) #[[ATTR0]] { 2170 // CHECK35-NEXT: entry: 2171 // CHECK35-NEXT: [[A_ADDR:%.*]] = alloca i32, align 4 2172 // CHECK35-NEXT: [[B_ADDR:%.*]] = alloca i32, align 4 2173 // CHECK35-NEXT: [[ARGC_ADDR:%.*]] = alloca i8**, align 4 2174 // CHECK35-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1]]) 2175 // CHECK35-NEXT: store i32 [[A]], i32* [[A_ADDR]], align 4 2176 // CHECK35-NEXT: store i32 [[B]], i32* [[B_ADDR]], align 4 2177 // CHECK35-NEXT: store i8** [[ARGC]], i8*** [[ARGC_ADDR]], align 4 2178 // CHECK35-NEXT: [[TMP1:%.*]] = load i32, i32* [[A_ADDR]], align 4 2179 // CHECK35-NEXT: [[TMP2:%.*]] = load i32, i32* [[B_ADDR]], align 4 2180 // CHECK35-NEXT: call void @__kmpc_push_num_teams(%struct.ident_t* @[[GLOB1]], i32 [[TMP0]], i32 [[TMP1]], i32 [[TMP2]]) 2181 // CHECK35-NEXT: call void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(%struct.ident_t* @[[GLOB1]], i32 1, void (i32*, i32*, ...)* bitcast (void (i32*, i32*, i8***)* @.omp_outlined..1 to void (i32*, i32*, ...)*), i8*** [[ARGC_ADDR]]) 2182 // CHECK35-NEXT: ret void 2183 // 2184 // 2185 // CHECK35-LABEL: define {{[^@]+}}@.omp_outlined..1 2186 // CHECK35-SAME: (i32* noalias noundef [[DOTGLOBAL_TID_:%.*]], i32* noalias noundef [[DOTBOUND_TID_:%.*]], i8*** noundef nonnull align 4 dereferenceable(4) [[ARGC:%.*]]) #[[ATTR0]] { 2187 // CHECK35-NEXT: entry: 2188 // CHECK35-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 4 2189 // CHECK35-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 4 2190 // CHECK35-NEXT: [[ARGC_ADDR:%.*]] = alloca i8***, align 4 2191 // CHECK35-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 4 2192 // CHECK35-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 4 2193 // CHECK35-NEXT: store i8*** [[ARGC]], i8**** [[ARGC_ADDR]], align 4 2194 // CHECK35-NEXT: [[TMP0:%.*]] = load i8***, i8**** [[ARGC_ADDR]], align 4 2195 // CHECK35-NEXT: store i8** null, i8*** [[TMP0]], align 4 2196 // CHECK35-NEXT: ret void 2197 // 2198 // 2199 // CHECK41-LABEL: define {{[^@]+}}@_Z3foov 2200 // CHECK41-SAME: () #[[ATTR0:[0-9]+]] { 2201 // CHECK41-NEXT: entry: 2202 // CHECK41-NEXT: call void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(%struct.ident_t* @[[GLOB1:[0-9]+]], i32 0, void (i32*, i32*, ...)* bitcast (void (i32*, i32*)* @.omp_outlined. to void (i32*, i32*, ...)*)) 2203 // CHECK41-NEXT: ret void 2204 // 2205 // 2206 // CHECK41-LABEL: define {{[^@]+}}@.omp_outlined. 2207 // CHECK41-SAME: (i32* noalias noundef [[DOTGLOBAL_TID_:%.*]], i32* noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR1:[0-9]+]] { 2208 // CHECK41-NEXT: entry: 2209 // CHECK41-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8 2210 // CHECK41-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8 2211 // CHECK41-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 8 2212 // CHECK41-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 8 2213 // CHECK41-NEXT: ret void 2214 // 2215 // 2216 // CHECK43-LABEL: define {{[^@]+}}@_Z3foov 2217 // CHECK43-SAME: () #[[ATTR0:[0-9]+]] { 2218 // CHECK43-NEXT: entry: 2219 // CHECK43-NEXT: call void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(%struct.ident_t* @[[GLOB1:[0-9]+]], i32 0, void (i32*, i32*, ...)* bitcast (void (i32*, i32*)* @.omp_outlined. to void (i32*, i32*, ...)*)) 2220 // CHECK43-NEXT: ret void 2221 // 2222 // 2223 // CHECK43-LABEL: define {{[^@]+}}@.omp_outlined. 2224 // CHECK43-SAME: (i32* noalias noundef [[DOTGLOBAL_TID_:%.*]], i32* noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR1:[0-9]+]] { 2225 // CHECK43-NEXT: entry: 2226 // CHECK43-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 4 2227 // CHECK43-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 4 2228 // CHECK43-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 4 2229 // CHECK43-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 4 2230 // CHECK43-NEXT: ret void 2231 // 2232