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 -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 -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 -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 -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 -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 -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 -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 -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 -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 -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 -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 -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 -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 -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 -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 -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 -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 -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 -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 -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 -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 -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 -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 -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 -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 -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 -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 -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 -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 -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 -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 -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 -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 -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 -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 -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 -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 -DCK4 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --check-prefix=CHECK25 188 // RUN: %clang_cc1 -DCK4 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o %t %s 189 // RUN: %clang_cc1 -DCK4 -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix=CHECK25 190 // RUN: %clang_cc1 -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 -DCK4 -verify -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s --check-prefix=CHECK27 192 // RUN: %clang_cc1 -DCK4 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-x86-host.bc -o %t %s 193 // RUN: %clang_cc1 -DCK4 -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-x86-host.bc -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix=CHECK27 194 195 // RUN: %clang_cc1 -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 -DCK4 -verify -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}" 197 // RUN: %clang_cc1 -DCK4 -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o %t %s 198 // RUN: %clang_cc1 -DCK4 -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}" 199 // RUN: %clang_cc1 -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 -DCK4 -verify -fopenmp-simd -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}" 201 // RUN: %clang_cc1 -DCK4 -fopenmp-simd -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-x86-host.bc -o %t %s 202 // RUN: %clang_cc1 -DCK4 -fopenmp-simd -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-x86-host.bc -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}" 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 -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 -DCK5 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --check-prefix=CHECK33 230 // RUN: %clang_cc1 -DCK5 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o %t %s 231 // RUN: %clang_cc1 -DCK5 -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix=CHECK33 232 // RUN: %clang_cc1 -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 -DCK5 -verify -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s --check-prefix=CHECK35 234 // RUN: %clang_cc1 -DCK5 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-x86-host.bc -o %t %s 235 // RUN: %clang_cc1 -DCK5 -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-x86-host.bc -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix=CHECK35 236 237 // RUN: %clang_cc1 -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 -DCK5 -verify -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}" 239 // RUN: %clang_cc1 -DCK5 -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o %t %s 240 // RUN: %clang_cc1 -DCK5 -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}" 241 // RUN: %clang_cc1 -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 -DCK5 -verify -fopenmp-simd -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}" 243 // RUN: %clang_cc1 -DCK5 -fopenmp-simd -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-x86-host.bc -o %t %s 244 // RUN: %clang_cc1 -DCK5 -fopenmp-simd -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-x86-host.bc -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}" 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 -DCK6 -verify -fopenmp -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 -DCK6 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s 280 // RUN: %clang_cc1 -DCK6 -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 -DCK6 -verify -fopenmp -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 -DCK6 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s 283 // RUN: %clang_cc1 -DCK6 -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=CHECK43 284 285 // RUN: %clang_cc1 -DCK6 -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}}" 286 // RUN: %clang_cc1 -DCK6 -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s 287 // RUN: %clang_cc1 -DCK6 -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}}" 288 // RUN: %clang_cc1 -DCK6 -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}}" 289 // RUN: %clang_cc1 -DCK6 -fopenmp-simd -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s 290 // RUN: %clang_cc1 -DCK6 -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}}" 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 ptr], align 8 310 // CHECK1-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x ptr], align 8 311 // CHECK1-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [1 x ptr], align 8 312 // CHECK1-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8 313 // CHECK1-NEXT: [[COMP_CASTED1:%.*]] = alloca i64, align 8 314 // CHECK1-NEXT: [[DOTOFFLOAD_BASEPTRS2:%.*]] = alloca [1 x ptr], align 8 315 // CHECK1-NEXT: [[DOTOFFLOAD_PTRS3:%.*]] = alloca [1 x ptr], align 8 316 // CHECK1-NEXT: [[DOTOFFLOAD_MAPPERS4:%.*]] = alloca [1 x ptr], align 8 317 // CHECK1-NEXT: [[KERNEL_ARGS5:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 8 318 // CHECK1-NEXT: [[LA_CASTED:%.*]] = alloca i64, align 8 319 // CHECK1-NEXT: [[COMP_CASTED8:%.*]] = alloca i64, align 8 320 // CHECK1-NEXT: [[DOTOFFLOAD_BASEPTRS9:%.*]] = alloca [2 x ptr], align 8 321 // CHECK1-NEXT: [[DOTOFFLOAD_PTRS10:%.*]] = alloca [2 x ptr], align 8 322 // CHECK1-NEXT: [[DOTOFFLOAD_MAPPERS11:%.*]] = alloca [2 x ptr], align 8 323 // CHECK1-NEXT: [[KERNEL_ARGS12:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 8 324 // CHECK1-NEXT: [[LA_CASTED15:%.*]] = alloca i64, align 8 325 // CHECK1-NEXT: [[COMP_CASTED16:%.*]] = alloca i64, align 8 326 // CHECK1-NEXT: [[DOTOFFLOAD_BASEPTRS17:%.*]] = alloca [2 x ptr], align 8 327 // CHECK1-NEXT: [[DOTOFFLOAD_PTRS18:%.*]] = alloca [2 x ptr], align 8 328 // CHECK1-NEXT: [[DOTOFFLOAD_MAPPERS19:%.*]] = alloca [2 x ptr], align 8 329 // CHECK1-NEXT: [[KERNEL_ARGS20:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 8 330 // CHECK1-NEXT: [[GBLA_CASTED:%.*]] = alloca i64, align 8 331 // CHECK1-NEXT: [[A_CASTED:%.*]] = alloca i64, align 8 332 // CHECK1-NEXT: [[GBLB_CASTED:%.*]] = alloca i64, align 8 333 // CHECK1-NEXT: [[LC_CASTED:%.*]] = alloca i64, align 8 334 // CHECK1-NEXT: [[COMP_CASTED23:%.*]] = alloca i64, align 8 335 // CHECK1-NEXT: [[DOTOFFLOAD_BASEPTRS24:%.*]] = alloca [5 x ptr], align 8 336 // CHECK1-NEXT: [[DOTOFFLOAD_PTRS25:%.*]] = alloca [5 x ptr], align 8 337 // CHECK1-NEXT: [[DOTOFFLOAD_MAPPERS26:%.*]] = alloca [5 x ptr], align 8 338 // CHECK1-NEXT: [[KERNEL_ARGS28:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 8 339 // CHECK1-NEXT: [[TMP:%.*]] = alloca ptr, align 8 340 // CHECK1-NEXT: [[GBLC_CASTED:%.*]] = alloca i64, align 8 341 // CHECK1-NEXT: [[COMP_CASTED31:%.*]] = alloca i64, align 8 342 // CHECK1-NEXT: [[DOTOFFLOAD_BASEPTRS32:%.*]] = alloca [2 x ptr], align 8 343 // CHECK1-NEXT: [[DOTOFFLOAD_PTRS33:%.*]] = alloca [2 x ptr], align 8 344 // CHECK1-NEXT: [[DOTOFFLOAD_MAPPERS34:%.*]] = alloca [2 x ptr], align 8 345 // CHECK1-NEXT: [[_TMP35:%.*]] = alloca ptr, align 8 346 // CHECK1-NEXT: [[_TMP37:%.*]] = alloca ptr, align 8 347 // CHECK1-NEXT: [[KERNEL_ARGS39:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 8 348 // CHECK1-NEXT: store i32 [[A]], ptr [[A_ADDR]], align 4 349 // CHECK1-NEXT: store i32 1, ptr [[COMP]], align 4 350 // CHECK1-NEXT: store i32 23, ptr [[LA]], align 4 351 // CHECK1-NEXT: store float 2.500000e+01, ptr [[LC]], align 4 352 // CHECK1-NEXT: [[TMP0:%.*]] = load i32, ptr [[COMP]], align 4 353 // CHECK1-NEXT: store i32 [[TMP0]], ptr [[COMP_CASTED]], align 4 354 // CHECK1-NEXT: [[TMP1:%.*]] = load i64, ptr [[COMP_CASTED]], align 8 355 // CHECK1-NEXT: [[TMP2:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 356 // CHECK1-NEXT: store i64 [[TMP1]], ptr [[TMP2]], align 8 357 // CHECK1-NEXT: [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0 358 // CHECK1-NEXT: store i64 [[TMP1]], ptr [[TMP3]], align 8 359 // CHECK1-NEXT: [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 0 360 // CHECK1-NEXT: store ptr null, ptr [[TMP4]], align 8 361 // CHECK1-NEXT: [[TMP5:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 362 // CHECK1-NEXT: [[TMP6:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0 363 // CHECK1-NEXT: [[TMP7:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0 364 // CHECK1-NEXT: store i32 3, ptr [[TMP7]], align 4 365 // CHECK1-NEXT: [[TMP8:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1 366 // CHECK1-NEXT: store i32 1, ptr [[TMP8]], align 4 367 // CHECK1-NEXT: [[TMP9:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2 368 // CHECK1-NEXT: store ptr [[TMP5]], ptr [[TMP9]], align 8 369 // CHECK1-NEXT: [[TMP10:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3 370 // CHECK1-NEXT: store ptr [[TMP6]], ptr [[TMP10]], align 8 371 // CHECK1-NEXT: [[TMP11:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4 372 // CHECK1-NEXT: store ptr @.offload_sizes, ptr [[TMP11]], align 8 373 // CHECK1-NEXT: [[TMP12:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5 374 // CHECK1-NEXT: store ptr @.offload_maptypes, ptr [[TMP12]], align 8 375 // CHECK1-NEXT: [[TMP13:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6 376 // CHECK1-NEXT: store ptr null, ptr [[TMP13]], align 8 377 // CHECK1-NEXT: [[TMP14:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7 378 // CHECK1-NEXT: store ptr null, ptr [[TMP14]], align 8 379 // CHECK1-NEXT: [[TMP15:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8 380 // CHECK1-NEXT: store i64 0, ptr [[TMP15]], align 8 381 // CHECK1-NEXT: [[TMP16:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9 382 // CHECK1-NEXT: store i64 0, ptr [[TMP16]], align 8 383 // CHECK1-NEXT: [[TMP17:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10 384 // CHECK1-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP17]], align 4 385 // CHECK1-NEXT: [[TMP18:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11 386 // CHECK1-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP18]], align 4 387 // CHECK1-NEXT: [[TMP19:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12 388 // CHECK1-NEXT: store i32 0, ptr [[TMP19]], align 4 389 // CHECK1-NEXT: [[TMP20:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1:[0-9]+]], i64 -1, i32 0, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l31.region_id, ptr [[KERNEL_ARGS]]) 390 // CHECK1-NEXT: [[TMP21:%.*]] = icmp ne i32 [[TMP20]], 0 391 // CHECK1-NEXT: br i1 [[TMP21]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]] 392 // CHECK1: omp_offload.failed: 393 // CHECK1-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l31(i64 [[TMP1]]) #[[ATTR2:[0-9]+]] 394 // CHECK1-NEXT: br label [[OMP_OFFLOAD_CONT]] 395 // CHECK1: omp_offload.cont: 396 // CHECK1-NEXT: [[TMP22:%.*]] = load i32, ptr [[COMP]], align 4 397 // CHECK1-NEXT: store i32 [[TMP22]], ptr [[COMP_CASTED1]], align 4 398 // CHECK1-NEXT: [[TMP23:%.*]] = load i64, ptr [[COMP_CASTED1]], align 8 399 // CHECK1-NEXT: [[TMP24:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS2]], i32 0, i32 0 400 // CHECK1-NEXT: store i64 [[TMP23]], ptr [[TMP24]], align 8 401 // CHECK1-NEXT: [[TMP25:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS3]], i32 0, i32 0 402 // CHECK1-NEXT: store i64 [[TMP23]], ptr [[TMP25]], align 8 403 // CHECK1-NEXT: [[TMP26:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS4]], i64 0, i64 0 404 // CHECK1-NEXT: store ptr null, ptr [[TMP26]], align 8 405 // CHECK1-NEXT: [[TMP27:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS2]], i32 0, i32 0 406 // CHECK1-NEXT: [[TMP28:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS3]], i32 0, i32 0 407 // CHECK1-NEXT: [[TMP29:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 0 408 // CHECK1-NEXT: store i32 3, ptr [[TMP29]], align 4 409 // CHECK1-NEXT: [[TMP30:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 1 410 // CHECK1-NEXT: store i32 1, ptr [[TMP30]], align 4 411 // CHECK1-NEXT: [[TMP31:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 2 412 // CHECK1-NEXT: store ptr [[TMP27]], ptr [[TMP31]], align 8 413 // CHECK1-NEXT: [[TMP32:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 3 414 // CHECK1-NEXT: store ptr [[TMP28]], ptr [[TMP32]], align 8 415 // CHECK1-NEXT: [[TMP33:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 4 416 // CHECK1-NEXT: store ptr @.offload_sizes.1, ptr [[TMP33]], align 8 417 // CHECK1-NEXT: [[TMP34:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 5 418 // CHECK1-NEXT: store ptr @.offload_maptypes.2, ptr [[TMP34]], align 8 419 // CHECK1-NEXT: [[TMP35:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 6 420 // CHECK1-NEXT: store ptr null, ptr [[TMP35]], align 8 421 // CHECK1-NEXT: [[TMP36:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 7 422 // CHECK1-NEXT: store ptr null, ptr [[TMP36]], align 8 423 // CHECK1-NEXT: [[TMP37:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 8 424 // CHECK1-NEXT: store i64 0, ptr [[TMP37]], align 8 425 // CHECK1-NEXT: [[TMP38:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 9 426 // CHECK1-NEXT: store i64 0, ptr [[TMP38]], align 8 427 // CHECK1-NEXT: [[TMP39:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 10 428 // CHECK1-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP39]], align 4 429 // CHECK1-NEXT: [[TMP40:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 11 430 // CHECK1-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP40]], align 4 431 // CHECK1-NEXT: [[TMP41:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 12 432 // CHECK1-NEXT: store i32 0, ptr [[TMP41]], align 4 433 // CHECK1-NEXT: [[TMP42:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 -1, i32 0, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l37.region_id, ptr [[KERNEL_ARGS5]]) 434 // CHECK1-NEXT: [[TMP43:%.*]] = icmp ne i32 [[TMP42]], 0 435 // CHECK1-NEXT: br i1 [[TMP43]], label [[OMP_OFFLOAD_FAILED6:%.*]], label [[OMP_OFFLOAD_CONT7:%.*]] 436 // CHECK1: omp_offload.failed6: 437 // CHECK1-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l37(i64 [[TMP23]]) #[[ATTR2]] 438 // CHECK1-NEXT: br label [[OMP_OFFLOAD_CONT7]] 439 // CHECK1: omp_offload.cont7: 440 // CHECK1-NEXT: [[TMP44:%.*]] = load i32, ptr [[LA]], align 4 441 // CHECK1-NEXT: store i32 [[TMP44]], ptr [[LA_CASTED]], align 4 442 // CHECK1-NEXT: [[TMP45:%.*]] = load i64, ptr [[LA_CASTED]], align 8 443 // CHECK1-NEXT: [[TMP46:%.*]] = load i32, ptr [[COMP]], align 4 444 // CHECK1-NEXT: store i32 [[TMP46]], ptr [[COMP_CASTED8]], align 4 445 // CHECK1-NEXT: [[TMP47:%.*]] = load i64, ptr [[COMP_CASTED8]], align 8 446 // CHECK1-NEXT: [[TMP48:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS9]], i32 0, i32 0 447 // CHECK1-NEXT: store i64 [[TMP45]], ptr [[TMP48]], align 8 448 // CHECK1-NEXT: [[TMP49:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS10]], i32 0, i32 0 449 // CHECK1-NEXT: store i64 [[TMP45]], ptr [[TMP49]], align 8 450 // CHECK1-NEXT: [[TMP50:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS11]], i64 0, i64 0 451 // CHECK1-NEXT: store ptr null, ptr [[TMP50]], align 8 452 // CHECK1-NEXT: [[TMP51:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS9]], i32 0, i32 1 453 // CHECK1-NEXT: store i64 [[TMP47]], ptr [[TMP51]], align 8 454 // CHECK1-NEXT: [[TMP52:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS10]], i32 0, i32 1 455 // CHECK1-NEXT: store i64 [[TMP47]], ptr [[TMP52]], align 8 456 // CHECK1-NEXT: [[TMP53:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS11]], i64 0, i64 1 457 // CHECK1-NEXT: store ptr null, ptr [[TMP53]], align 8 458 // CHECK1-NEXT: [[TMP54:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS9]], i32 0, i32 0 459 // CHECK1-NEXT: [[TMP55:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS10]], i32 0, i32 0 460 // CHECK1-NEXT: [[TMP56:%.*]] = load i32, ptr [[LA]], align 4 461 // CHECK1-NEXT: [[TMP57:%.*]] = insertvalue [3 x i32] zeroinitializer, i32 [[TMP56]], 0 462 // CHECK1-NEXT: [[TMP58:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS12]], i32 0, i32 0 463 // CHECK1-NEXT: store i32 3, ptr [[TMP58]], align 4 464 // CHECK1-NEXT: [[TMP59:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS12]], i32 0, i32 1 465 // CHECK1-NEXT: store i32 2, ptr [[TMP59]], align 4 466 // CHECK1-NEXT: [[TMP60:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS12]], i32 0, i32 2 467 // CHECK1-NEXT: store ptr [[TMP54]], ptr [[TMP60]], align 8 468 // CHECK1-NEXT: [[TMP61:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS12]], i32 0, i32 3 469 // CHECK1-NEXT: store ptr [[TMP55]], ptr [[TMP61]], align 8 470 // CHECK1-NEXT: [[TMP62:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS12]], i32 0, i32 4 471 // CHECK1-NEXT: store ptr @.offload_sizes.3, ptr [[TMP62]], align 8 472 // CHECK1-NEXT: [[TMP63:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS12]], i32 0, i32 5 473 // CHECK1-NEXT: store ptr @.offload_maptypes.4, ptr [[TMP63]], align 8 474 // CHECK1-NEXT: [[TMP64:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS12]], i32 0, i32 6 475 // CHECK1-NEXT: store ptr null, ptr [[TMP64]], align 8 476 // CHECK1-NEXT: [[TMP65:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS12]], i32 0, i32 7 477 // CHECK1-NEXT: store ptr null, ptr [[TMP65]], align 8 478 // CHECK1-NEXT: [[TMP66:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS12]], i32 0, i32 8 479 // CHECK1-NEXT: store i64 0, ptr [[TMP66]], align 8 480 // CHECK1-NEXT: [[TMP67:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS12]], i32 0, i32 9 481 // CHECK1-NEXT: store i64 0, ptr [[TMP67]], align 8 482 // CHECK1-NEXT: [[TMP68:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS12]], i32 0, i32 10 483 // CHECK1-NEXT: store [3 x i32] [[TMP57]], ptr [[TMP68]], align 4 484 // CHECK1-NEXT: [[TMP69:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS12]], i32 0, i32 11 485 // CHECK1-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP69]], align 4 486 // CHECK1-NEXT: [[TMP70:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS12]], i32 0, i32 12 487 // CHECK1-NEXT: store i32 0, ptr [[TMP70]], align 4 488 // CHECK1-NEXT: [[TMP71:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 -1, i32 [[TMP56]], i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l46.region_id, ptr [[KERNEL_ARGS12]]) 489 // CHECK1-NEXT: [[TMP72:%.*]] = icmp ne i32 [[TMP71]], 0 490 // CHECK1-NEXT: br i1 [[TMP72]], label [[OMP_OFFLOAD_FAILED13:%.*]], label [[OMP_OFFLOAD_CONT14:%.*]] 491 // CHECK1: omp_offload.failed13: 492 // CHECK1-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l46(i64 [[TMP45]], i64 [[TMP47]]) #[[ATTR2]] 493 // CHECK1-NEXT: br label [[OMP_OFFLOAD_CONT14]] 494 // CHECK1: omp_offload.cont14: 495 // CHECK1-NEXT: [[TMP73:%.*]] = load i32, ptr [[LA]], align 4 496 // CHECK1-NEXT: store i32 [[TMP73]], ptr [[LA_CASTED15]], align 4 497 // CHECK1-NEXT: [[TMP74:%.*]] = load i64, ptr [[LA_CASTED15]], align 8 498 // CHECK1-NEXT: [[TMP75:%.*]] = load i32, ptr [[COMP]], align 4 499 // CHECK1-NEXT: store i32 [[TMP75]], ptr [[COMP_CASTED16]], align 4 500 // CHECK1-NEXT: [[TMP76:%.*]] = load i64, ptr [[COMP_CASTED16]], align 8 501 // CHECK1-NEXT: [[TMP77:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS17]], i32 0, i32 0 502 // CHECK1-NEXT: store i64 [[TMP74]], ptr [[TMP77]], align 8 503 // CHECK1-NEXT: [[TMP78:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS18]], i32 0, i32 0 504 // CHECK1-NEXT: store i64 [[TMP74]], ptr [[TMP78]], align 8 505 // CHECK1-NEXT: [[TMP79:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS19]], i64 0, i64 0 506 // CHECK1-NEXT: store ptr null, ptr [[TMP79]], align 8 507 // CHECK1-NEXT: [[TMP80:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS17]], i32 0, i32 1 508 // CHECK1-NEXT: store i64 [[TMP76]], ptr [[TMP80]], align 8 509 // CHECK1-NEXT: [[TMP81:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS18]], i32 0, i32 1 510 // CHECK1-NEXT: store i64 [[TMP76]], ptr [[TMP81]], align 8 511 // CHECK1-NEXT: [[TMP82:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS19]], i64 0, i64 1 512 // CHECK1-NEXT: store ptr null, ptr [[TMP82]], align 8 513 // CHECK1-NEXT: [[TMP83:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS17]], i32 0, i32 0 514 // CHECK1-NEXT: [[TMP84:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS18]], i32 0, i32 0 515 // CHECK1-NEXT: [[TMP85:%.*]] = load i32, ptr [[LA]], align 4 516 // CHECK1-NEXT: [[TMP86:%.*]] = insertvalue [3 x i32] zeroinitializer, i32 [[TMP85]], 0 517 // CHECK1-NEXT: [[TMP87:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS20]], i32 0, i32 0 518 // CHECK1-NEXT: store i32 3, ptr [[TMP87]], align 4 519 // CHECK1-NEXT: [[TMP88:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS20]], i32 0, i32 1 520 // CHECK1-NEXT: store i32 2, ptr [[TMP88]], align 4 521 // CHECK1-NEXT: [[TMP89:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS20]], i32 0, i32 2 522 // CHECK1-NEXT: store ptr [[TMP83]], ptr [[TMP89]], align 8 523 // CHECK1-NEXT: [[TMP90:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS20]], i32 0, i32 3 524 // CHECK1-NEXT: store ptr [[TMP84]], ptr [[TMP90]], align 8 525 // CHECK1-NEXT: [[TMP91:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS20]], i32 0, i32 4 526 // CHECK1-NEXT: store ptr @.offload_sizes.5, ptr [[TMP91]], align 8 527 // CHECK1-NEXT: [[TMP92:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS20]], i32 0, i32 5 528 // CHECK1-NEXT: store ptr @.offload_maptypes.6, ptr [[TMP92]], align 8 529 // CHECK1-NEXT: [[TMP93:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS20]], i32 0, i32 6 530 // CHECK1-NEXT: store ptr null, ptr [[TMP93]], align 8 531 // CHECK1-NEXT: [[TMP94:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS20]], i32 0, i32 7 532 // CHECK1-NEXT: store ptr null, ptr [[TMP94]], align 8 533 // CHECK1-NEXT: [[TMP95:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS20]], i32 0, i32 8 534 // CHECK1-NEXT: store i64 0, ptr [[TMP95]], align 8 535 // CHECK1-NEXT: [[TMP96:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS20]], i32 0, i32 9 536 // CHECK1-NEXT: store i64 0, ptr [[TMP96]], align 8 537 // CHECK1-NEXT: [[TMP97:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS20]], i32 0, i32 10 538 // CHECK1-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP97]], align 4 539 // CHECK1-NEXT: [[TMP98:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS20]], i32 0, i32 11 540 // CHECK1-NEXT: store [3 x i32] [[TMP86]], ptr [[TMP98]], align 4 541 // CHECK1-NEXT: [[TMP99:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS20]], i32 0, i32 12 542 // CHECK1-NEXT: store i32 0, ptr [[TMP99]], align 4 543 // CHECK1-NEXT: [[TMP100:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 -1, i32 0, i32 [[TMP85]], ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l53.region_id, ptr [[KERNEL_ARGS20]]) 544 // CHECK1-NEXT: [[TMP101:%.*]] = icmp ne i32 [[TMP100]], 0 545 // CHECK1-NEXT: br i1 [[TMP101]], label [[OMP_OFFLOAD_FAILED21:%.*]], label [[OMP_OFFLOAD_CONT22:%.*]] 546 // CHECK1: omp_offload.failed21: 547 // CHECK1-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l53(i64 [[TMP74]], i64 [[TMP76]]) #[[ATTR2]] 548 // CHECK1-NEXT: br label [[OMP_OFFLOAD_CONT22]] 549 // CHECK1: omp_offload.cont22: 550 // CHECK1-NEXT: [[TMP102:%.*]] = load i32, ptr @Gbla, align 4 551 // CHECK1-NEXT: store i32 [[TMP102]], ptr [[GBLA_CASTED]], align 4 552 // CHECK1-NEXT: [[TMP103:%.*]] = load i64, ptr [[GBLA_CASTED]], align 8 553 // CHECK1-NEXT: [[TMP104:%.*]] = load i32, ptr [[A_ADDR]], align 4 554 // CHECK1-NEXT: store i32 [[TMP104]], ptr [[A_CASTED]], align 4 555 // CHECK1-NEXT: [[TMP105:%.*]] = load i64, ptr [[A_CASTED]], align 8 556 // CHECK1-NEXT: [[TMP106:%.*]] = load i64, ptr @Gblb, align 8 557 // CHECK1-NEXT: store i64 [[TMP106]], ptr [[GBLB_CASTED]], align 8 558 // CHECK1-NEXT: [[TMP107:%.*]] = load i64, ptr [[GBLB_CASTED]], align 8 559 // CHECK1-NEXT: [[TMP108:%.*]] = load float, ptr [[LC]], align 4 560 // CHECK1-NEXT: store float [[TMP108]], ptr [[LC_CASTED]], align 4 561 // CHECK1-NEXT: [[TMP109:%.*]] = load i64, ptr [[LC_CASTED]], align 8 562 // CHECK1-NEXT: [[TMP110:%.*]] = load i32, ptr [[COMP]], align 4 563 // CHECK1-NEXT: store i32 [[TMP110]], ptr [[COMP_CASTED23]], align 4 564 // CHECK1-NEXT: [[TMP111:%.*]] = load i64, ptr [[COMP_CASTED23]], align 8 565 // CHECK1-NEXT: [[TMP112:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_BASEPTRS24]], i32 0, i32 0 566 // CHECK1-NEXT: store i64 [[TMP103]], ptr [[TMP112]], align 8 567 // CHECK1-NEXT: [[TMP113:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_PTRS25]], i32 0, i32 0 568 // CHECK1-NEXT: store i64 [[TMP103]], ptr [[TMP113]], align 8 569 // CHECK1-NEXT: [[TMP114:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_MAPPERS26]], i64 0, i64 0 570 // CHECK1-NEXT: store ptr null, ptr [[TMP114]], align 8 571 // CHECK1-NEXT: [[TMP115:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_BASEPTRS24]], i32 0, i32 1 572 // CHECK1-NEXT: store i64 [[TMP105]], ptr [[TMP115]], align 8 573 // CHECK1-NEXT: [[TMP116:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_PTRS25]], i32 0, i32 1 574 // CHECK1-NEXT: store i64 [[TMP105]], ptr [[TMP116]], align 8 575 // CHECK1-NEXT: [[TMP117:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_MAPPERS26]], i64 0, i64 1 576 // CHECK1-NEXT: store ptr null, ptr [[TMP117]], align 8 577 // CHECK1-NEXT: [[TMP118:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_BASEPTRS24]], i32 0, i32 2 578 // CHECK1-NEXT: store i64 [[TMP107]], ptr [[TMP118]], align 8 579 // CHECK1-NEXT: [[TMP119:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_PTRS25]], i32 0, i32 2 580 // CHECK1-NEXT: store i64 [[TMP107]], ptr [[TMP119]], align 8 581 // CHECK1-NEXT: [[TMP120:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_MAPPERS26]], i64 0, i64 2 582 // CHECK1-NEXT: store ptr null, ptr [[TMP120]], align 8 583 // CHECK1-NEXT: [[TMP121:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_BASEPTRS24]], i32 0, i32 3 584 // CHECK1-NEXT: store i64 [[TMP109]], ptr [[TMP121]], align 8 585 // CHECK1-NEXT: [[TMP122:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_PTRS25]], i32 0, i32 3 586 // CHECK1-NEXT: store i64 [[TMP109]], ptr [[TMP122]], align 8 587 // CHECK1-NEXT: [[TMP123:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_MAPPERS26]], i64 0, i64 3 588 // CHECK1-NEXT: store ptr null, ptr [[TMP123]], align 8 589 // CHECK1-NEXT: [[TMP124:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_BASEPTRS24]], i32 0, i32 4 590 // CHECK1-NEXT: store i64 [[TMP111]], ptr [[TMP124]], align 8 591 // CHECK1-NEXT: [[TMP125:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_PTRS25]], i32 0, i32 4 592 // CHECK1-NEXT: store i64 [[TMP111]], ptr [[TMP125]], align 8 593 // CHECK1-NEXT: [[TMP126:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_MAPPERS26]], i64 0, i64 4 594 // CHECK1-NEXT: store ptr null, ptr [[TMP126]], align 8 595 // CHECK1-NEXT: [[TMP127:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_BASEPTRS24]], i32 0, i32 0 596 // CHECK1-NEXT: [[TMP128:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_PTRS25]], i32 0, i32 0 597 // CHECK1-NEXT: [[TMP129:%.*]] = load i32, ptr @Gbla, align 4 598 // CHECK1-NEXT: [[TMP130:%.*]] = load i32, ptr [[A_ADDR]], align 4 599 // CHECK1-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP129]], [[TMP130]] 600 // CHECK1-NEXT: [[TMP131:%.*]] = load i64, ptr @Gblb, align 8 601 // CHECK1-NEXT: [[TMP132:%.*]] = load float, ptr [[LC]], align 4 602 // CHECK1-NEXT: [[CONV:%.*]] = fptosi float [[TMP132]] to i64 603 // CHECK1-NEXT: [[ADD27:%.*]] = add nsw i64 [[TMP131]], [[CONV]] 604 // CHECK1-NEXT: [[TMP133:%.*]] = trunc i64 [[ADD27]] to i32 605 // CHECK1-NEXT: [[TMP134:%.*]] = insertvalue [3 x i32] zeroinitializer, i32 [[ADD]], 0 606 // CHECK1-NEXT: [[TMP135:%.*]] = insertvalue [3 x i32] zeroinitializer, i32 [[TMP133]], 0 607 // CHECK1-NEXT: [[TMP136:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS28]], i32 0, i32 0 608 // CHECK1-NEXT: store i32 3, ptr [[TMP136]], align 4 609 // CHECK1-NEXT: [[TMP137:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS28]], i32 0, i32 1 610 // CHECK1-NEXT: store i32 5, ptr [[TMP137]], align 4 611 // CHECK1-NEXT: [[TMP138:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS28]], i32 0, i32 2 612 // CHECK1-NEXT: store ptr [[TMP127]], ptr [[TMP138]], align 8 613 // CHECK1-NEXT: [[TMP139:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS28]], i32 0, i32 3 614 // CHECK1-NEXT: store ptr [[TMP128]], ptr [[TMP139]], align 8 615 // CHECK1-NEXT: [[TMP140:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS28]], i32 0, i32 4 616 // CHECK1-NEXT: store ptr @.offload_sizes.7, ptr [[TMP140]], align 8 617 // CHECK1-NEXT: [[TMP141:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS28]], i32 0, i32 5 618 // CHECK1-NEXT: store ptr @.offload_maptypes.8, ptr [[TMP141]], align 8 619 // CHECK1-NEXT: [[TMP142:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS28]], i32 0, i32 6 620 // CHECK1-NEXT: store ptr null, ptr [[TMP142]], align 8 621 // CHECK1-NEXT: [[TMP143:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS28]], i32 0, i32 7 622 // CHECK1-NEXT: store ptr null, ptr [[TMP143]], align 8 623 // CHECK1-NEXT: [[TMP144:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS28]], i32 0, i32 8 624 // CHECK1-NEXT: store i64 0, ptr [[TMP144]], align 8 625 // CHECK1-NEXT: [[TMP145:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS28]], i32 0, i32 9 626 // CHECK1-NEXT: store i64 0, ptr [[TMP145]], align 8 627 // CHECK1-NEXT: [[TMP146:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS28]], i32 0, i32 10 628 // CHECK1-NEXT: store [3 x i32] [[TMP134]], ptr [[TMP146]], align 4 629 // CHECK1-NEXT: [[TMP147:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS28]], i32 0, i32 11 630 // CHECK1-NEXT: store [3 x i32] [[TMP135]], ptr [[TMP147]], align 4 631 // CHECK1-NEXT: [[TMP148:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS28]], i32 0, i32 12 632 // CHECK1-NEXT: store i32 0, ptr [[TMP148]], align 4 633 // CHECK1-NEXT: [[TMP149:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 -1, i32 [[ADD]], i32 [[TMP133]], ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l62.region_id, ptr [[KERNEL_ARGS28]]) 634 // CHECK1-NEXT: [[TMP150:%.*]] = icmp ne i32 [[TMP149]], 0 635 // CHECK1-NEXT: br i1 [[TMP150]], label [[OMP_OFFLOAD_FAILED29:%.*]], label [[OMP_OFFLOAD_CONT30:%.*]] 636 // CHECK1: omp_offload.failed29: 637 // CHECK1-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l62(i64 [[TMP103]], i64 [[TMP105]], i64 [[TMP107]], i64 [[TMP109]], i64 [[TMP111]]) #[[ATTR2]] 638 // CHECK1-NEXT: br label [[OMP_OFFLOAD_CONT30]] 639 // CHECK1: omp_offload.cont30: 640 // CHECK1-NEXT: [[TMP151:%.*]] = load ptr, ptr @Gblc, align 8 641 // CHECK1-NEXT: store ptr [[TMP151]], ptr [[TMP]], align 8 642 // CHECK1-NEXT: [[TMP152:%.*]] = load i32, ptr @Gbla, align 4 643 // CHECK1-NEXT: store i32 [[TMP152]], ptr [[GBLC_CASTED]], align 4 644 // CHECK1-NEXT: [[TMP153:%.*]] = load i64, ptr [[GBLC_CASTED]], align 8 645 // CHECK1-NEXT: [[TMP154:%.*]] = load i32, ptr [[COMP]], align 4 646 // CHECK1-NEXT: store i32 [[TMP154]], ptr [[COMP_CASTED31]], align 4 647 // CHECK1-NEXT: [[TMP155:%.*]] = load i64, ptr [[COMP_CASTED31]], align 8 648 // CHECK1-NEXT: [[TMP156:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS32]], i32 0, i32 0 649 // CHECK1-NEXT: store i64 [[TMP153]], ptr [[TMP156]], align 8 650 // CHECK1-NEXT: [[TMP157:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS33]], i32 0, i32 0 651 // CHECK1-NEXT: store i64 [[TMP153]], ptr [[TMP157]], align 8 652 // CHECK1-NEXT: [[TMP158:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS34]], i64 0, i64 0 653 // CHECK1-NEXT: store ptr null, ptr [[TMP158]], align 8 654 // CHECK1-NEXT: [[TMP159:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS32]], i32 0, i32 1 655 // CHECK1-NEXT: store i64 [[TMP155]], ptr [[TMP159]], align 8 656 // CHECK1-NEXT: [[TMP160:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS33]], i32 0, i32 1 657 // CHECK1-NEXT: store i64 [[TMP155]], ptr [[TMP160]], align 8 658 // CHECK1-NEXT: [[TMP161:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS34]], i64 0, i64 1 659 // CHECK1-NEXT: store ptr null, ptr [[TMP161]], align 8 660 // CHECK1-NEXT: [[TMP162:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS32]], i32 0, i32 0 661 // CHECK1-NEXT: [[TMP163:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS33]], i32 0, i32 0 662 // CHECK1-NEXT: [[TMP164:%.*]] = load ptr, ptr @Gblc, align 8 663 // CHECK1-NEXT: store ptr [[TMP164]], ptr [[_TMP35]], align 8 664 // CHECK1-NEXT: [[TMP165:%.*]] = load i32, ptr @Gbla, align 4 665 // CHECK1-NEXT: [[ADD36:%.*]] = add nsw i32 [[TMP165]], 1 666 // CHECK1-NEXT: [[TMP166:%.*]] = load ptr, ptr @Gblc, align 8 667 // CHECK1-NEXT: store ptr [[TMP166]], ptr [[_TMP37]], align 8 668 // CHECK1-NEXT: [[TMP167:%.*]] = load i32, ptr @Gbla, align 4 669 // CHECK1-NEXT: [[ADD38:%.*]] = add nsw i32 [[TMP167]], 2 670 // CHECK1-NEXT: [[TMP168:%.*]] = insertvalue [3 x i32] zeroinitializer, i32 [[ADD36]], 0 671 // CHECK1-NEXT: [[TMP169:%.*]] = insertvalue [3 x i32] zeroinitializer, i32 [[ADD38]], 0 672 // CHECK1-NEXT: [[TMP170:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS39]], i32 0, i32 0 673 // CHECK1-NEXT: store i32 3, ptr [[TMP170]], align 4 674 // CHECK1-NEXT: [[TMP171:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS39]], i32 0, i32 1 675 // CHECK1-NEXT: store i32 2, ptr [[TMP171]], align 4 676 // CHECK1-NEXT: [[TMP172:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS39]], i32 0, i32 2 677 // CHECK1-NEXT: store ptr [[TMP162]], ptr [[TMP172]], align 8 678 // CHECK1-NEXT: [[TMP173:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS39]], i32 0, i32 3 679 // CHECK1-NEXT: store ptr [[TMP163]], ptr [[TMP173]], align 8 680 // CHECK1-NEXT: [[TMP174:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS39]], i32 0, i32 4 681 // CHECK1-NEXT: store ptr @.offload_sizes.9, ptr [[TMP174]], align 8 682 // CHECK1-NEXT: [[TMP175:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS39]], i32 0, i32 5 683 // CHECK1-NEXT: store ptr @.offload_maptypes.10, ptr [[TMP175]], align 8 684 // CHECK1-NEXT: [[TMP176:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS39]], i32 0, i32 6 685 // CHECK1-NEXT: store ptr null, ptr [[TMP176]], align 8 686 // CHECK1-NEXT: [[TMP177:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS39]], i32 0, i32 7 687 // CHECK1-NEXT: store ptr null, ptr [[TMP177]], align 8 688 // CHECK1-NEXT: [[TMP178:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS39]], i32 0, i32 8 689 // CHECK1-NEXT: store i64 0, ptr [[TMP178]], align 8 690 // CHECK1-NEXT: [[TMP179:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS39]], i32 0, i32 9 691 // CHECK1-NEXT: store i64 0, ptr [[TMP179]], align 8 692 // CHECK1-NEXT: [[TMP180:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS39]], i32 0, i32 10 693 // CHECK1-NEXT: store [3 x i32] [[TMP168]], ptr [[TMP180]], align 4 694 // CHECK1-NEXT: [[TMP181:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS39]], i32 0, i32 11 695 // CHECK1-NEXT: store [3 x i32] [[TMP169]], ptr [[TMP181]], align 4 696 // CHECK1-NEXT: [[TMP182:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS39]], i32 0, i32 12 697 // CHECK1-NEXT: store i32 0, ptr [[TMP182]], align 4 698 // CHECK1-NEXT: [[TMP183:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 -1, i32 [[ADD36]], i32 [[ADD38]], ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l71.region_id, ptr [[KERNEL_ARGS39]]) 699 // CHECK1-NEXT: [[TMP184:%.*]] = icmp ne i32 [[TMP183]], 0 700 // CHECK1-NEXT: br i1 [[TMP184]], label [[OMP_OFFLOAD_FAILED40:%.*]], label [[OMP_OFFLOAD_CONT41:%.*]] 701 // CHECK1: omp_offload.failed40: 702 // CHECK1-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l71(i64 [[TMP153]], i64 [[TMP155]]) #[[ATTR2]] 703 // CHECK1-NEXT: br label [[OMP_OFFLOAD_CONT41]] 704 // CHECK1: omp_offload.cont41: 705 // CHECK1-NEXT: [[TMP185:%.*]] = load i32, ptr [[COMP]], align 4 706 // CHECK1-NEXT: ret i32 [[TMP185]] 707 // 708 // 709 // CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l31 710 // CHECK1-SAME: (i64 noundef [[COMP:%.*]]) #[[ATTR1:[0-9]+]] { 711 // CHECK1-NEXT: entry: 712 // CHECK1-NEXT: [[COMP_ADDR:%.*]] = alloca i64, align 8 713 // CHECK1-NEXT: store i64 [[COMP]], ptr [[COMP_ADDR]], align 8 714 // CHECK1-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB1]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l31.omp_outlined, ptr [[COMP_ADDR]]) 715 // CHECK1-NEXT: ret void 716 // 717 // 718 // CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l31.omp_outlined 719 // CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[COMP:%.*]]) #[[ATTR1]] { 720 // CHECK1-NEXT: entry: 721 // CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8 722 // CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8 723 // CHECK1-NEXT: [[COMP_ADDR:%.*]] = alloca ptr, align 8 724 // CHECK1-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8 725 // CHECK1-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8 726 // CHECK1-NEXT: store ptr [[COMP]], ptr [[COMP_ADDR]], align 8 727 // CHECK1-NEXT: [[TMP0:%.*]] = load ptr, ptr [[COMP_ADDR]], align 8 728 // CHECK1-NEXT: [[TMP1:%.*]] = load i32, ptr [[TMP0]], align 4 729 // CHECK1-NEXT: [[INC:%.*]] = add nsw i32 [[TMP1]], 1 730 // CHECK1-NEXT: store i32 [[INC]], ptr [[TMP0]], align 4 731 // CHECK1-NEXT: ret void 732 // 733 // 734 // CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l37 735 // CHECK1-SAME: (i64 noundef [[COMP:%.*]]) #[[ATTR1]] { 736 // CHECK1-NEXT: entry: 737 // CHECK1-NEXT: [[COMP_ADDR:%.*]] = alloca i64, align 8 738 // CHECK1-NEXT: store i64 [[COMP]], ptr [[COMP_ADDR]], align 8 739 // CHECK1-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB1]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l37.omp_outlined, ptr [[COMP_ADDR]]) 740 // CHECK1-NEXT: ret void 741 // 742 // 743 // CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l37.omp_outlined 744 // CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[COMP:%.*]]) #[[ATTR1]] { 745 // CHECK1-NEXT: entry: 746 // CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8 747 // CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8 748 // CHECK1-NEXT: [[COMP_ADDR:%.*]] = alloca ptr, align 8 749 // CHECK1-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8 750 // CHECK1-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8 751 // CHECK1-NEXT: store ptr [[COMP]], ptr [[COMP_ADDR]], align 8 752 // CHECK1-NEXT: [[TMP0:%.*]] = load ptr, ptr [[COMP_ADDR]], align 8 753 // CHECK1-NEXT: [[TMP1:%.*]] = load i32, ptr [[TMP0]], align 4 754 // CHECK1-NEXT: [[INC:%.*]] = add nsw i32 [[TMP1]], 1 755 // CHECK1-NEXT: store i32 [[INC]], ptr [[TMP0]], align 4 756 // CHECK1-NEXT: ret void 757 // 758 // 759 // CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l46 760 // CHECK1-SAME: (i64 noundef [[LA:%.*]], i64 noundef [[COMP:%.*]]) #[[ATTR1]] { 761 // CHECK1-NEXT: entry: 762 // CHECK1-NEXT: [[LA_ADDR:%.*]] = alloca i64, align 8 763 // CHECK1-NEXT: [[COMP_ADDR:%.*]] = alloca i64, align 8 764 // CHECK1-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]]) 765 // CHECK1-NEXT: store i64 [[LA]], ptr [[LA_ADDR]], align 8 766 // CHECK1-NEXT: store i64 [[COMP]], ptr [[COMP_ADDR]], align 8 767 // CHECK1-NEXT: [[TMP1:%.*]] = load i32, ptr [[LA_ADDR]], align 4 768 // CHECK1-NEXT: call void @__kmpc_push_num_teams(ptr @[[GLOB1]], i32 [[TMP0]], i32 [[TMP1]], i32 0) 769 // CHECK1-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB1]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l46.omp_outlined, ptr [[COMP_ADDR]]) 770 // CHECK1-NEXT: ret void 771 // 772 // 773 // CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l46.omp_outlined 774 // CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[COMP:%.*]]) #[[ATTR1]] { 775 // CHECK1-NEXT: entry: 776 // CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8 777 // CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8 778 // CHECK1-NEXT: [[COMP_ADDR:%.*]] = alloca ptr, align 8 779 // CHECK1-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8 780 // CHECK1-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8 781 // CHECK1-NEXT: store ptr [[COMP]], ptr [[COMP_ADDR]], align 8 782 // CHECK1-NEXT: [[TMP0:%.*]] = load ptr, ptr [[COMP_ADDR]], align 8 783 // CHECK1-NEXT: [[TMP1:%.*]] = load i32, ptr [[TMP0]], align 4 784 // CHECK1-NEXT: [[INC:%.*]] = add nsw i32 [[TMP1]], 1 785 // CHECK1-NEXT: store i32 [[INC]], ptr [[TMP0]], align 4 786 // CHECK1-NEXT: ret void 787 // 788 // 789 // CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l53 790 // CHECK1-SAME: (i64 noundef [[LA:%.*]], i64 noundef [[COMP:%.*]]) #[[ATTR1]] { 791 // CHECK1-NEXT: entry: 792 // CHECK1-NEXT: [[LA_ADDR:%.*]] = alloca i64, align 8 793 // CHECK1-NEXT: [[COMP_ADDR:%.*]] = alloca i64, align 8 794 // CHECK1-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]]) 795 // CHECK1-NEXT: store i64 [[LA]], ptr [[LA_ADDR]], align 8 796 // CHECK1-NEXT: store i64 [[COMP]], ptr [[COMP_ADDR]], align 8 797 // CHECK1-NEXT: [[TMP1:%.*]] = load i32, ptr [[LA_ADDR]], align 4 798 // CHECK1-NEXT: call void @__kmpc_push_num_teams(ptr @[[GLOB1]], i32 [[TMP0]], i32 0, i32 [[TMP1]]) 799 // CHECK1-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB1]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l53.omp_outlined, ptr [[COMP_ADDR]]) 800 // CHECK1-NEXT: ret void 801 // 802 // 803 // CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l53.omp_outlined 804 // CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[COMP:%.*]]) #[[ATTR1]] { 805 // CHECK1-NEXT: entry: 806 // CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8 807 // CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8 808 // CHECK1-NEXT: [[COMP_ADDR:%.*]] = alloca ptr, align 8 809 // CHECK1-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8 810 // CHECK1-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8 811 // CHECK1-NEXT: store ptr [[COMP]], ptr [[COMP_ADDR]], align 8 812 // CHECK1-NEXT: [[TMP0:%.*]] = load ptr, ptr [[COMP_ADDR]], align 8 813 // CHECK1-NEXT: [[TMP1:%.*]] = load i32, ptr [[TMP0]], align 4 814 // CHECK1-NEXT: [[INC:%.*]] = add nsw i32 [[TMP1]], 1 815 // CHECK1-NEXT: store i32 [[INC]], ptr [[TMP0]], align 4 816 // CHECK1-NEXT: ret void 817 // 818 // 819 // CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l62 820 // CHECK1-SAME: (i64 noundef [[GBLA:%.*]], i64 noundef [[A:%.*]], i64 noundef [[GBLB:%.*]], i64 noundef [[LC:%.*]], i64 noundef [[COMP:%.*]]) #[[ATTR1]] { 821 // CHECK1-NEXT: entry: 822 // CHECK1-NEXT: [[GBLA_ADDR:%.*]] = alloca i64, align 8 823 // CHECK1-NEXT: [[A_ADDR:%.*]] = alloca i64, align 8 824 // CHECK1-NEXT: [[GBLB_ADDR:%.*]] = alloca i64, align 8 825 // CHECK1-NEXT: [[LC_ADDR:%.*]] = alloca i64, align 8 826 // CHECK1-NEXT: [[COMP_ADDR:%.*]] = alloca i64, align 8 827 // CHECK1-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]]) 828 // CHECK1-NEXT: store i64 [[GBLA]], ptr [[GBLA_ADDR]], align 8 829 // CHECK1-NEXT: store i64 [[A]], ptr [[A_ADDR]], align 8 830 // CHECK1-NEXT: store i64 [[GBLB]], ptr [[GBLB_ADDR]], align 8 831 // CHECK1-NEXT: store i64 [[LC]], ptr [[LC_ADDR]], align 8 832 // CHECK1-NEXT: store i64 [[COMP]], ptr [[COMP_ADDR]], align 8 833 // CHECK1-NEXT: [[TMP1:%.*]] = load i32, ptr [[GBLA_ADDR]], align 4 834 // CHECK1-NEXT: [[TMP2:%.*]] = load i32, ptr [[A_ADDR]], align 4 835 // CHECK1-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP1]], [[TMP2]] 836 // CHECK1-NEXT: [[TMP3:%.*]] = load i64, ptr [[GBLB_ADDR]], align 8 837 // CHECK1-NEXT: [[TMP4:%.*]] = load float, ptr [[LC_ADDR]], align 4 838 // CHECK1-NEXT: [[CONV:%.*]] = fptosi float [[TMP4]] to i64 839 // CHECK1-NEXT: [[ADD1:%.*]] = add nsw i64 [[TMP3]], [[CONV]] 840 // CHECK1-NEXT: [[TMP5:%.*]] = trunc i64 [[ADD1]] to i32 841 // CHECK1-NEXT: call void @__kmpc_push_num_teams(ptr @[[GLOB1]], i32 [[TMP0]], i32 [[ADD]], i32 [[TMP5]]) 842 // CHECK1-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB1]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l62.omp_outlined, ptr [[COMP_ADDR]]) 843 // CHECK1-NEXT: ret void 844 // 845 // 846 // CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l62.omp_outlined 847 // CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[COMP:%.*]]) #[[ATTR1]] { 848 // CHECK1-NEXT: entry: 849 // CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8 850 // CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8 851 // CHECK1-NEXT: [[COMP_ADDR:%.*]] = alloca ptr, align 8 852 // CHECK1-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8 853 // CHECK1-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8 854 // CHECK1-NEXT: store ptr [[COMP]], ptr [[COMP_ADDR]], align 8 855 // CHECK1-NEXT: [[TMP0:%.*]] = load ptr, ptr [[COMP_ADDR]], align 8 856 // CHECK1-NEXT: [[TMP1:%.*]] = load i32, ptr [[TMP0]], align 4 857 // CHECK1-NEXT: [[INC:%.*]] = add nsw i32 [[TMP1]], 1 858 // CHECK1-NEXT: store i32 [[INC]], ptr [[TMP0]], align 4 859 // CHECK1-NEXT: ret void 860 // 861 // 862 // CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l71 863 // CHECK1-SAME: (i64 noundef [[GBLC:%.*]], i64 noundef [[COMP:%.*]]) #[[ATTR1]] { 864 // CHECK1-NEXT: entry: 865 // CHECK1-NEXT: [[GBLC_ADDR:%.*]] = alloca i64, align 8 866 // CHECK1-NEXT: [[COMP_ADDR:%.*]] = alloca i64, align 8 867 // CHECK1-NEXT: [[TMP:%.*]] = alloca ptr, align 8 868 // CHECK1-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]]) 869 // CHECK1-NEXT: store i64 [[GBLC]], ptr [[GBLC_ADDR]], align 8 870 // CHECK1-NEXT: store i64 [[COMP]], ptr [[COMP_ADDR]], align 8 871 // CHECK1-NEXT: store ptr [[GBLC_ADDR]], ptr [[TMP]], align 8 872 // CHECK1-NEXT: [[TMP1:%.*]] = load i32, ptr @Gbla, align 4 873 // CHECK1-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP1]], 1 874 // CHECK1-NEXT: [[TMP2:%.*]] = load i32, ptr @Gbla, align 4 875 // CHECK1-NEXT: [[ADD1:%.*]] = add nsw i32 [[TMP2]], 2 876 // CHECK1-NEXT: call void @__kmpc_push_num_teams(ptr @[[GLOB1]], i32 [[TMP0]], i32 [[ADD]], i32 [[ADD1]]) 877 // CHECK1-NEXT: [[TMP3:%.*]] = load ptr, ptr [[TMP]], align 8 878 // CHECK1-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB1]], i32 2, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l71.omp_outlined, ptr [[COMP_ADDR]], ptr [[TMP3]]) 879 // CHECK1-NEXT: ret void 880 // 881 // 882 // CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l71.omp_outlined 883 // CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[COMP:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[GBLC:%.*]]) #[[ATTR1]] { 884 // CHECK1-NEXT: entry: 885 // CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8 886 // CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8 887 // CHECK1-NEXT: [[COMP_ADDR:%.*]] = alloca ptr, align 8 888 // CHECK1-NEXT: [[GBLC_ADDR:%.*]] = alloca ptr, align 8 889 // CHECK1-NEXT: [[TMP:%.*]] = alloca ptr, align 8 890 // CHECK1-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8 891 // CHECK1-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8 892 // CHECK1-NEXT: store ptr [[COMP]], ptr [[COMP_ADDR]], align 8 893 // CHECK1-NEXT: store ptr [[GBLC]], ptr [[GBLC_ADDR]], align 8 894 // CHECK1-NEXT: [[TMP0:%.*]] = load ptr, ptr [[COMP_ADDR]], align 8 895 // CHECK1-NEXT: [[TMP1:%.*]] = load ptr, ptr [[GBLC_ADDR]], align 8 896 // CHECK1-NEXT: store ptr [[TMP1]], ptr [[TMP]], align 8 897 // CHECK1-NEXT: [[TMP2:%.*]] = load i32, ptr @Gbla, align 4 898 // CHECK1-NEXT: [[TMP3:%.*]] = load i32, ptr [[TMP0]], align 4 899 // CHECK1-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP3]], [[TMP2]] 900 // CHECK1-NEXT: store i32 [[ADD]], ptr [[TMP0]], align 4 901 // CHECK1-NEXT: ret void 902 // 903 // 904 // CHECK3-LABEL: define {{[^@]+}}@_Z27teams_argument_global_locali 905 // CHECK3-SAME: (i32 noundef [[A:%.*]]) #[[ATTR0:[0-9]+]] { 906 // CHECK3-NEXT: entry: 907 // CHECK3-NEXT: [[A_ADDR:%.*]] = alloca i32, align 4 908 // CHECK3-NEXT: [[COMP:%.*]] = alloca i32, align 4 909 // CHECK3-NEXT: [[LA:%.*]] = alloca i32, align 4 910 // CHECK3-NEXT: [[LC:%.*]] = alloca float, align 4 911 // CHECK3-NEXT: [[COMP_CASTED:%.*]] = alloca i32, align 4 912 // CHECK3-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x ptr], align 4 913 // CHECK3-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x ptr], align 4 914 // CHECK3-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [1 x ptr], align 4 915 // CHECK3-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8 916 // CHECK3-NEXT: [[COMP_CASTED1:%.*]] = alloca i32, align 4 917 // CHECK3-NEXT: [[DOTOFFLOAD_BASEPTRS2:%.*]] = alloca [1 x ptr], align 4 918 // CHECK3-NEXT: [[DOTOFFLOAD_PTRS3:%.*]] = alloca [1 x ptr], align 4 919 // CHECK3-NEXT: [[DOTOFFLOAD_MAPPERS4:%.*]] = alloca [1 x ptr], align 4 920 // CHECK3-NEXT: [[KERNEL_ARGS5:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 8 921 // CHECK3-NEXT: [[LA_CASTED:%.*]] = alloca i32, align 4 922 // CHECK3-NEXT: [[COMP_CASTED8:%.*]] = alloca i32, align 4 923 // CHECK3-NEXT: [[DOTOFFLOAD_BASEPTRS9:%.*]] = alloca [2 x ptr], align 4 924 // CHECK3-NEXT: [[DOTOFFLOAD_PTRS10:%.*]] = alloca [2 x ptr], align 4 925 // CHECK3-NEXT: [[DOTOFFLOAD_MAPPERS11:%.*]] = alloca [2 x ptr], align 4 926 // CHECK3-NEXT: [[KERNEL_ARGS12:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 8 927 // CHECK3-NEXT: [[LA_CASTED15:%.*]] = alloca i32, align 4 928 // CHECK3-NEXT: [[COMP_CASTED16:%.*]] = alloca i32, align 4 929 // CHECK3-NEXT: [[DOTOFFLOAD_BASEPTRS17:%.*]] = alloca [2 x ptr], align 4 930 // CHECK3-NEXT: [[DOTOFFLOAD_PTRS18:%.*]] = alloca [2 x ptr], align 4 931 // CHECK3-NEXT: [[DOTOFFLOAD_MAPPERS19:%.*]] = alloca [2 x ptr], align 4 932 // CHECK3-NEXT: [[KERNEL_ARGS20:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 8 933 // CHECK3-NEXT: [[GBLA_CASTED:%.*]] = alloca i32, align 4 934 // CHECK3-NEXT: [[A_CASTED:%.*]] = alloca i32, align 4 935 // CHECK3-NEXT: [[LC_CASTED:%.*]] = alloca i32, align 4 936 // CHECK3-NEXT: [[COMP_CASTED23:%.*]] = alloca i32, align 4 937 // CHECK3-NEXT: [[DOTOFFLOAD_BASEPTRS24:%.*]] = alloca [5 x ptr], align 4 938 // CHECK3-NEXT: [[DOTOFFLOAD_PTRS25:%.*]] = alloca [5 x ptr], align 4 939 // CHECK3-NEXT: [[DOTOFFLOAD_MAPPERS26:%.*]] = alloca [5 x ptr], align 4 940 // CHECK3-NEXT: [[KERNEL_ARGS28:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 8 941 // CHECK3-NEXT: [[TMP:%.*]] = alloca ptr, align 4 942 // CHECK3-NEXT: [[GBLC_CASTED:%.*]] = alloca i32, align 4 943 // CHECK3-NEXT: [[COMP_CASTED31:%.*]] = alloca i32, align 4 944 // CHECK3-NEXT: [[DOTOFFLOAD_BASEPTRS32:%.*]] = alloca [2 x ptr], align 4 945 // CHECK3-NEXT: [[DOTOFFLOAD_PTRS33:%.*]] = alloca [2 x ptr], align 4 946 // CHECK3-NEXT: [[DOTOFFLOAD_MAPPERS34:%.*]] = alloca [2 x ptr], align 4 947 // CHECK3-NEXT: [[_TMP35:%.*]] = alloca ptr, align 4 948 // CHECK3-NEXT: [[_TMP37:%.*]] = alloca ptr, align 4 949 // CHECK3-NEXT: [[KERNEL_ARGS39:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 8 950 // CHECK3-NEXT: store i32 [[A]], ptr [[A_ADDR]], align 4 951 // CHECK3-NEXT: store i32 1, ptr [[COMP]], align 4 952 // CHECK3-NEXT: store i32 23, ptr [[LA]], align 4 953 // CHECK3-NEXT: store float 2.500000e+01, ptr [[LC]], align 4 954 // CHECK3-NEXT: [[TMP0:%.*]] = load i32, ptr [[COMP]], align 4 955 // CHECK3-NEXT: store i32 [[TMP0]], ptr [[COMP_CASTED]], align 4 956 // CHECK3-NEXT: [[TMP1:%.*]] = load i32, ptr [[COMP_CASTED]], align 4 957 // CHECK3-NEXT: [[TMP2:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 958 // CHECK3-NEXT: store i32 [[TMP1]], ptr [[TMP2]], align 4 959 // CHECK3-NEXT: [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0 960 // CHECK3-NEXT: store i32 [[TMP1]], ptr [[TMP3]], align 4 961 // CHECK3-NEXT: [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 0 962 // CHECK3-NEXT: store ptr null, ptr [[TMP4]], align 4 963 // CHECK3-NEXT: [[TMP5:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 964 // CHECK3-NEXT: [[TMP6:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0 965 // CHECK3-NEXT: [[TMP7:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0 966 // CHECK3-NEXT: store i32 3, ptr [[TMP7]], align 4 967 // CHECK3-NEXT: [[TMP8:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1 968 // CHECK3-NEXT: store i32 1, ptr [[TMP8]], align 4 969 // CHECK3-NEXT: [[TMP9:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2 970 // CHECK3-NEXT: store ptr [[TMP5]], ptr [[TMP9]], align 4 971 // CHECK3-NEXT: [[TMP10:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3 972 // CHECK3-NEXT: store ptr [[TMP6]], ptr [[TMP10]], align 4 973 // CHECK3-NEXT: [[TMP11:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4 974 // CHECK3-NEXT: store ptr @.offload_sizes, ptr [[TMP11]], align 4 975 // CHECK3-NEXT: [[TMP12:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5 976 // CHECK3-NEXT: store ptr @.offload_maptypes, ptr [[TMP12]], align 4 977 // CHECK3-NEXT: [[TMP13:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6 978 // CHECK3-NEXT: store ptr null, ptr [[TMP13]], align 4 979 // CHECK3-NEXT: [[TMP14:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7 980 // CHECK3-NEXT: store ptr null, ptr [[TMP14]], align 4 981 // CHECK3-NEXT: [[TMP15:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8 982 // CHECK3-NEXT: store i64 0, ptr [[TMP15]], align 8 983 // CHECK3-NEXT: [[TMP16:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9 984 // CHECK3-NEXT: store i64 0, ptr [[TMP16]], align 8 985 // CHECK3-NEXT: [[TMP17:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10 986 // CHECK3-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP17]], align 4 987 // CHECK3-NEXT: [[TMP18:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11 988 // CHECK3-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP18]], align 4 989 // CHECK3-NEXT: [[TMP19:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12 990 // CHECK3-NEXT: store i32 0, ptr [[TMP19]], align 4 991 // CHECK3-NEXT: [[TMP20:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1:[0-9]+]], i64 -1, i32 0, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l31.region_id, ptr [[KERNEL_ARGS]]) 992 // CHECK3-NEXT: [[TMP21:%.*]] = icmp ne i32 [[TMP20]], 0 993 // CHECK3-NEXT: br i1 [[TMP21]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]] 994 // CHECK3: omp_offload.failed: 995 // CHECK3-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l31(i32 [[TMP1]]) #[[ATTR2:[0-9]+]] 996 // CHECK3-NEXT: br label [[OMP_OFFLOAD_CONT]] 997 // CHECK3: omp_offload.cont: 998 // CHECK3-NEXT: [[TMP22:%.*]] = load i32, ptr [[COMP]], align 4 999 // CHECK3-NEXT: store i32 [[TMP22]], ptr [[COMP_CASTED1]], align 4 1000 // CHECK3-NEXT: [[TMP23:%.*]] = load i32, ptr [[COMP_CASTED1]], align 4 1001 // CHECK3-NEXT: [[TMP24:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS2]], i32 0, i32 0 1002 // CHECK3-NEXT: store i32 [[TMP23]], ptr [[TMP24]], align 4 1003 // CHECK3-NEXT: [[TMP25:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS3]], i32 0, i32 0 1004 // CHECK3-NEXT: store i32 [[TMP23]], ptr [[TMP25]], align 4 1005 // CHECK3-NEXT: [[TMP26:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS4]], i32 0, i32 0 1006 // CHECK3-NEXT: store ptr null, ptr [[TMP26]], align 4 1007 // CHECK3-NEXT: [[TMP27:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS2]], i32 0, i32 0 1008 // CHECK3-NEXT: [[TMP28:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS3]], i32 0, i32 0 1009 // CHECK3-NEXT: [[TMP29:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 0 1010 // CHECK3-NEXT: store i32 3, ptr [[TMP29]], align 4 1011 // CHECK3-NEXT: [[TMP30:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 1 1012 // CHECK3-NEXT: store i32 1, ptr [[TMP30]], align 4 1013 // CHECK3-NEXT: [[TMP31:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 2 1014 // CHECK3-NEXT: store ptr [[TMP27]], ptr [[TMP31]], align 4 1015 // CHECK3-NEXT: [[TMP32:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 3 1016 // CHECK3-NEXT: store ptr [[TMP28]], ptr [[TMP32]], align 4 1017 // CHECK3-NEXT: [[TMP33:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 4 1018 // CHECK3-NEXT: store ptr @.offload_sizes.1, ptr [[TMP33]], align 4 1019 // CHECK3-NEXT: [[TMP34:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 5 1020 // CHECK3-NEXT: store ptr @.offload_maptypes.2, ptr [[TMP34]], align 4 1021 // CHECK3-NEXT: [[TMP35:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 6 1022 // CHECK3-NEXT: store ptr null, ptr [[TMP35]], align 4 1023 // CHECK3-NEXT: [[TMP36:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 7 1024 // CHECK3-NEXT: store ptr null, ptr [[TMP36]], align 4 1025 // CHECK3-NEXT: [[TMP37:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 8 1026 // CHECK3-NEXT: store i64 0, ptr [[TMP37]], align 8 1027 // CHECK3-NEXT: [[TMP38:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 9 1028 // CHECK3-NEXT: store i64 0, ptr [[TMP38]], align 8 1029 // CHECK3-NEXT: [[TMP39:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 10 1030 // CHECK3-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP39]], align 4 1031 // CHECK3-NEXT: [[TMP40:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 11 1032 // CHECK3-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP40]], align 4 1033 // CHECK3-NEXT: [[TMP41:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 12 1034 // CHECK3-NEXT: store i32 0, ptr [[TMP41]], align 4 1035 // CHECK3-NEXT: [[TMP42:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 -1, i32 0, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l37.region_id, ptr [[KERNEL_ARGS5]]) 1036 // CHECK3-NEXT: [[TMP43:%.*]] = icmp ne i32 [[TMP42]], 0 1037 // CHECK3-NEXT: br i1 [[TMP43]], label [[OMP_OFFLOAD_FAILED6:%.*]], label [[OMP_OFFLOAD_CONT7:%.*]] 1038 // CHECK3: omp_offload.failed6: 1039 // CHECK3-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l37(i32 [[TMP23]]) #[[ATTR2]] 1040 // CHECK3-NEXT: br label [[OMP_OFFLOAD_CONT7]] 1041 // CHECK3: omp_offload.cont7: 1042 // CHECK3-NEXT: [[TMP44:%.*]] = load i32, ptr [[LA]], align 4 1043 // CHECK3-NEXT: store i32 [[TMP44]], ptr [[LA_CASTED]], align 4 1044 // CHECK3-NEXT: [[TMP45:%.*]] = load i32, ptr [[LA_CASTED]], align 4 1045 // CHECK3-NEXT: [[TMP46:%.*]] = load i32, ptr [[COMP]], align 4 1046 // CHECK3-NEXT: store i32 [[TMP46]], ptr [[COMP_CASTED8]], align 4 1047 // CHECK3-NEXT: [[TMP47:%.*]] = load i32, ptr [[COMP_CASTED8]], align 4 1048 // CHECK3-NEXT: [[TMP48:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS9]], i32 0, i32 0 1049 // CHECK3-NEXT: store i32 [[TMP45]], ptr [[TMP48]], align 4 1050 // CHECK3-NEXT: [[TMP49:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS10]], i32 0, i32 0 1051 // CHECK3-NEXT: store i32 [[TMP45]], ptr [[TMP49]], align 4 1052 // CHECK3-NEXT: [[TMP50:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS11]], i32 0, i32 0 1053 // CHECK3-NEXT: store ptr null, ptr [[TMP50]], align 4 1054 // CHECK3-NEXT: [[TMP51:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS9]], i32 0, i32 1 1055 // CHECK3-NEXT: store i32 [[TMP47]], ptr [[TMP51]], align 4 1056 // CHECK3-NEXT: [[TMP52:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS10]], i32 0, i32 1 1057 // CHECK3-NEXT: store i32 [[TMP47]], ptr [[TMP52]], align 4 1058 // CHECK3-NEXT: [[TMP53:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS11]], i32 0, i32 1 1059 // CHECK3-NEXT: store ptr null, ptr [[TMP53]], align 4 1060 // CHECK3-NEXT: [[TMP54:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS9]], i32 0, i32 0 1061 // CHECK3-NEXT: [[TMP55:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS10]], i32 0, i32 0 1062 // CHECK3-NEXT: [[TMP56:%.*]] = load i32, ptr [[LA]], align 4 1063 // CHECK3-NEXT: [[TMP57:%.*]] = insertvalue [3 x i32] zeroinitializer, i32 [[TMP56]], 0 1064 // CHECK3-NEXT: [[TMP58:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS12]], i32 0, i32 0 1065 // CHECK3-NEXT: store i32 3, ptr [[TMP58]], align 4 1066 // CHECK3-NEXT: [[TMP59:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS12]], i32 0, i32 1 1067 // CHECK3-NEXT: store i32 2, ptr [[TMP59]], align 4 1068 // CHECK3-NEXT: [[TMP60:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS12]], i32 0, i32 2 1069 // CHECK3-NEXT: store ptr [[TMP54]], ptr [[TMP60]], align 4 1070 // CHECK3-NEXT: [[TMP61:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS12]], i32 0, i32 3 1071 // CHECK3-NEXT: store ptr [[TMP55]], ptr [[TMP61]], align 4 1072 // CHECK3-NEXT: [[TMP62:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS12]], i32 0, i32 4 1073 // CHECK3-NEXT: store ptr @.offload_sizes.3, ptr [[TMP62]], align 4 1074 // CHECK3-NEXT: [[TMP63:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS12]], i32 0, i32 5 1075 // CHECK3-NEXT: store ptr @.offload_maptypes.4, ptr [[TMP63]], align 4 1076 // CHECK3-NEXT: [[TMP64:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS12]], i32 0, i32 6 1077 // CHECK3-NEXT: store ptr null, ptr [[TMP64]], align 4 1078 // CHECK3-NEXT: [[TMP65:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS12]], i32 0, i32 7 1079 // CHECK3-NEXT: store ptr null, ptr [[TMP65]], align 4 1080 // CHECK3-NEXT: [[TMP66:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS12]], i32 0, i32 8 1081 // CHECK3-NEXT: store i64 0, ptr [[TMP66]], align 8 1082 // CHECK3-NEXT: [[TMP67:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS12]], i32 0, i32 9 1083 // CHECK3-NEXT: store i64 0, ptr [[TMP67]], align 8 1084 // CHECK3-NEXT: [[TMP68:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS12]], i32 0, i32 10 1085 // CHECK3-NEXT: store [3 x i32] [[TMP57]], ptr [[TMP68]], align 4 1086 // CHECK3-NEXT: [[TMP69:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS12]], i32 0, i32 11 1087 // CHECK3-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP69]], align 4 1088 // CHECK3-NEXT: [[TMP70:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS12]], i32 0, i32 12 1089 // CHECK3-NEXT: store i32 0, ptr [[TMP70]], align 4 1090 // CHECK3-NEXT: [[TMP71:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 -1, i32 [[TMP56]], i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l46.region_id, ptr [[KERNEL_ARGS12]]) 1091 // CHECK3-NEXT: [[TMP72:%.*]] = icmp ne i32 [[TMP71]], 0 1092 // CHECK3-NEXT: br i1 [[TMP72]], label [[OMP_OFFLOAD_FAILED13:%.*]], label [[OMP_OFFLOAD_CONT14:%.*]] 1093 // CHECK3: omp_offload.failed13: 1094 // CHECK3-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l46(i32 [[TMP45]], i32 [[TMP47]]) #[[ATTR2]] 1095 // CHECK3-NEXT: br label [[OMP_OFFLOAD_CONT14]] 1096 // CHECK3: omp_offload.cont14: 1097 // CHECK3-NEXT: [[TMP73:%.*]] = load i32, ptr [[LA]], align 4 1098 // CHECK3-NEXT: store i32 [[TMP73]], ptr [[LA_CASTED15]], align 4 1099 // CHECK3-NEXT: [[TMP74:%.*]] = load i32, ptr [[LA_CASTED15]], align 4 1100 // CHECK3-NEXT: [[TMP75:%.*]] = load i32, ptr [[COMP]], align 4 1101 // CHECK3-NEXT: store i32 [[TMP75]], ptr [[COMP_CASTED16]], align 4 1102 // CHECK3-NEXT: [[TMP76:%.*]] = load i32, ptr [[COMP_CASTED16]], align 4 1103 // CHECK3-NEXT: [[TMP77:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS17]], i32 0, i32 0 1104 // CHECK3-NEXT: store i32 [[TMP74]], ptr [[TMP77]], align 4 1105 // CHECK3-NEXT: [[TMP78:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS18]], i32 0, i32 0 1106 // CHECK3-NEXT: store i32 [[TMP74]], ptr [[TMP78]], align 4 1107 // CHECK3-NEXT: [[TMP79:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS19]], i32 0, i32 0 1108 // CHECK3-NEXT: store ptr null, ptr [[TMP79]], align 4 1109 // CHECK3-NEXT: [[TMP80:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS17]], i32 0, i32 1 1110 // CHECK3-NEXT: store i32 [[TMP76]], ptr [[TMP80]], align 4 1111 // CHECK3-NEXT: [[TMP81:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS18]], i32 0, i32 1 1112 // CHECK3-NEXT: store i32 [[TMP76]], ptr [[TMP81]], align 4 1113 // CHECK3-NEXT: [[TMP82:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS19]], i32 0, i32 1 1114 // CHECK3-NEXT: store ptr null, ptr [[TMP82]], align 4 1115 // CHECK3-NEXT: [[TMP83:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS17]], i32 0, i32 0 1116 // CHECK3-NEXT: [[TMP84:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS18]], i32 0, i32 0 1117 // CHECK3-NEXT: [[TMP85:%.*]] = load i32, ptr [[LA]], align 4 1118 // CHECK3-NEXT: [[TMP86:%.*]] = insertvalue [3 x i32] zeroinitializer, i32 [[TMP85]], 0 1119 // CHECK3-NEXT: [[TMP87:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS20]], i32 0, i32 0 1120 // CHECK3-NEXT: store i32 3, ptr [[TMP87]], align 4 1121 // CHECK3-NEXT: [[TMP88:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS20]], i32 0, i32 1 1122 // CHECK3-NEXT: store i32 2, ptr [[TMP88]], align 4 1123 // CHECK3-NEXT: [[TMP89:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS20]], i32 0, i32 2 1124 // CHECK3-NEXT: store ptr [[TMP83]], ptr [[TMP89]], align 4 1125 // CHECK3-NEXT: [[TMP90:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS20]], i32 0, i32 3 1126 // CHECK3-NEXT: store ptr [[TMP84]], ptr [[TMP90]], align 4 1127 // CHECK3-NEXT: [[TMP91:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS20]], i32 0, i32 4 1128 // CHECK3-NEXT: store ptr @.offload_sizes.5, ptr [[TMP91]], align 4 1129 // CHECK3-NEXT: [[TMP92:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS20]], i32 0, i32 5 1130 // CHECK3-NEXT: store ptr @.offload_maptypes.6, ptr [[TMP92]], align 4 1131 // CHECK3-NEXT: [[TMP93:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS20]], i32 0, i32 6 1132 // CHECK3-NEXT: store ptr null, ptr [[TMP93]], align 4 1133 // CHECK3-NEXT: [[TMP94:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS20]], i32 0, i32 7 1134 // CHECK3-NEXT: store ptr null, ptr [[TMP94]], align 4 1135 // CHECK3-NEXT: [[TMP95:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS20]], i32 0, i32 8 1136 // CHECK3-NEXT: store i64 0, ptr [[TMP95]], align 8 1137 // CHECK3-NEXT: [[TMP96:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS20]], i32 0, i32 9 1138 // CHECK3-NEXT: store i64 0, ptr [[TMP96]], align 8 1139 // CHECK3-NEXT: [[TMP97:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS20]], i32 0, i32 10 1140 // CHECK3-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP97]], align 4 1141 // CHECK3-NEXT: [[TMP98:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS20]], i32 0, i32 11 1142 // CHECK3-NEXT: store [3 x i32] [[TMP86]], ptr [[TMP98]], align 4 1143 // CHECK3-NEXT: [[TMP99:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS20]], i32 0, i32 12 1144 // CHECK3-NEXT: store i32 0, ptr [[TMP99]], align 4 1145 // CHECK3-NEXT: [[TMP100:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 -1, i32 0, i32 [[TMP85]], ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l53.region_id, ptr [[KERNEL_ARGS20]]) 1146 // CHECK3-NEXT: [[TMP101:%.*]] = icmp ne i32 [[TMP100]], 0 1147 // CHECK3-NEXT: br i1 [[TMP101]], label [[OMP_OFFLOAD_FAILED21:%.*]], label [[OMP_OFFLOAD_CONT22:%.*]] 1148 // CHECK3: omp_offload.failed21: 1149 // CHECK3-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l53(i32 [[TMP74]], i32 [[TMP76]]) #[[ATTR2]] 1150 // CHECK3-NEXT: br label [[OMP_OFFLOAD_CONT22]] 1151 // CHECK3: omp_offload.cont22: 1152 // CHECK3-NEXT: [[TMP102:%.*]] = load i32, ptr @Gbla, align 4 1153 // CHECK3-NEXT: store i32 [[TMP102]], ptr [[GBLA_CASTED]], align 4 1154 // CHECK3-NEXT: [[TMP103:%.*]] = load i32, ptr [[GBLA_CASTED]], align 4 1155 // CHECK3-NEXT: [[TMP104:%.*]] = load i32, ptr [[A_ADDR]], align 4 1156 // CHECK3-NEXT: store i32 [[TMP104]], ptr [[A_CASTED]], align 4 1157 // CHECK3-NEXT: [[TMP105:%.*]] = load i32, ptr [[A_CASTED]], align 4 1158 // CHECK3-NEXT: [[TMP106:%.*]] = load float, ptr [[LC]], align 4 1159 // CHECK3-NEXT: store float [[TMP106]], ptr [[LC_CASTED]], align 4 1160 // CHECK3-NEXT: [[TMP107:%.*]] = load i32, ptr [[LC_CASTED]], align 4 1161 // CHECK3-NEXT: [[TMP108:%.*]] = load i32, ptr [[COMP]], align 4 1162 // CHECK3-NEXT: store i32 [[TMP108]], ptr [[COMP_CASTED23]], align 4 1163 // CHECK3-NEXT: [[TMP109:%.*]] = load i32, ptr [[COMP_CASTED23]], align 4 1164 // CHECK3-NEXT: [[TMP110:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_BASEPTRS24]], i32 0, i32 0 1165 // CHECK3-NEXT: store i32 [[TMP103]], ptr [[TMP110]], align 4 1166 // CHECK3-NEXT: [[TMP111:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_PTRS25]], i32 0, i32 0 1167 // CHECK3-NEXT: store i32 [[TMP103]], ptr [[TMP111]], align 4 1168 // CHECK3-NEXT: [[TMP112:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_MAPPERS26]], i32 0, i32 0 1169 // CHECK3-NEXT: store ptr null, ptr [[TMP112]], align 4 1170 // CHECK3-NEXT: [[TMP113:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_BASEPTRS24]], i32 0, i32 1 1171 // CHECK3-NEXT: store i32 [[TMP105]], ptr [[TMP113]], align 4 1172 // CHECK3-NEXT: [[TMP114:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_PTRS25]], i32 0, i32 1 1173 // CHECK3-NEXT: store i32 [[TMP105]], ptr [[TMP114]], align 4 1174 // CHECK3-NEXT: [[TMP115:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_MAPPERS26]], i32 0, i32 1 1175 // CHECK3-NEXT: store ptr null, ptr [[TMP115]], align 4 1176 // CHECK3-NEXT: [[TMP116:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_BASEPTRS24]], i32 0, i32 2 1177 // CHECK3-NEXT: store ptr @Gblb, ptr [[TMP116]], align 4 1178 // CHECK3-NEXT: [[TMP117:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_PTRS25]], i32 0, i32 2 1179 // CHECK3-NEXT: store ptr @Gblb, ptr [[TMP117]], align 4 1180 // CHECK3-NEXT: [[TMP118:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_MAPPERS26]], i32 0, i32 2 1181 // CHECK3-NEXT: store ptr null, ptr [[TMP118]], align 4 1182 // CHECK3-NEXT: [[TMP119:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_BASEPTRS24]], i32 0, i32 3 1183 // CHECK3-NEXT: store i32 [[TMP107]], ptr [[TMP119]], align 4 1184 // CHECK3-NEXT: [[TMP120:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_PTRS25]], i32 0, i32 3 1185 // CHECK3-NEXT: store i32 [[TMP107]], ptr [[TMP120]], align 4 1186 // CHECK3-NEXT: [[TMP121:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_MAPPERS26]], i32 0, i32 3 1187 // CHECK3-NEXT: store ptr null, ptr [[TMP121]], align 4 1188 // CHECK3-NEXT: [[TMP122:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_BASEPTRS24]], i32 0, i32 4 1189 // CHECK3-NEXT: store i32 [[TMP109]], ptr [[TMP122]], align 4 1190 // CHECK3-NEXT: [[TMP123:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_PTRS25]], i32 0, i32 4 1191 // CHECK3-NEXT: store i32 [[TMP109]], ptr [[TMP123]], align 4 1192 // CHECK3-NEXT: [[TMP124:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_MAPPERS26]], i32 0, i32 4 1193 // CHECK3-NEXT: store ptr null, ptr [[TMP124]], align 4 1194 // CHECK3-NEXT: [[TMP125:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_BASEPTRS24]], i32 0, i32 0 1195 // CHECK3-NEXT: [[TMP126:%.*]] = getelementptr inbounds [5 x ptr], ptr [[DOTOFFLOAD_PTRS25]], i32 0, i32 0 1196 // CHECK3-NEXT: [[TMP127:%.*]] = load i32, ptr @Gbla, align 4 1197 // CHECK3-NEXT: [[TMP128:%.*]] = load i32, ptr [[A_ADDR]], align 4 1198 // CHECK3-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP127]], [[TMP128]] 1199 // CHECK3-NEXT: [[TMP129:%.*]] = load i64, ptr @Gblb, align 8 1200 // CHECK3-NEXT: [[TMP130:%.*]] = load float, ptr [[LC]], align 4 1201 // CHECK3-NEXT: [[CONV:%.*]] = fptosi float [[TMP130]] to i64 1202 // CHECK3-NEXT: [[ADD27:%.*]] = add nsw i64 [[TMP129]], [[CONV]] 1203 // CHECK3-NEXT: [[TMP131:%.*]] = trunc i64 [[ADD27]] to i32 1204 // CHECK3-NEXT: [[TMP132:%.*]] = insertvalue [3 x i32] zeroinitializer, i32 [[ADD]], 0 1205 // CHECK3-NEXT: [[TMP133:%.*]] = insertvalue [3 x i32] zeroinitializer, i32 [[TMP131]], 0 1206 // CHECK3-NEXT: [[TMP134:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS28]], i32 0, i32 0 1207 // CHECK3-NEXT: store i32 3, ptr [[TMP134]], align 4 1208 // CHECK3-NEXT: [[TMP135:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS28]], i32 0, i32 1 1209 // CHECK3-NEXT: store i32 5, ptr [[TMP135]], align 4 1210 // CHECK3-NEXT: [[TMP136:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS28]], i32 0, i32 2 1211 // CHECK3-NEXT: store ptr [[TMP125]], ptr [[TMP136]], align 4 1212 // CHECK3-NEXT: [[TMP137:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS28]], i32 0, i32 3 1213 // CHECK3-NEXT: store ptr [[TMP126]], ptr [[TMP137]], align 4 1214 // CHECK3-NEXT: [[TMP138:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS28]], i32 0, i32 4 1215 // CHECK3-NEXT: store ptr @.offload_sizes.7, ptr [[TMP138]], align 4 1216 // CHECK3-NEXT: [[TMP139:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS28]], i32 0, i32 5 1217 // CHECK3-NEXT: store ptr @.offload_maptypes.8, ptr [[TMP139]], align 4 1218 // CHECK3-NEXT: [[TMP140:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS28]], i32 0, i32 6 1219 // CHECK3-NEXT: store ptr null, ptr [[TMP140]], align 4 1220 // CHECK3-NEXT: [[TMP141:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS28]], i32 0, i32 7 1221 // CHECK3-NEXT: store ptr null, ptr [[TMP141]], align 4 1222 // CHECK3-NEXT: [[TMP142:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS28]], i32 0, i32 8 1223 // CHECK3-NEXT: store i64 0, ptr [[TMP142]], align 8 1224 // CHECK3-NEXT: [[TMP143:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS28]], i32 0, i32 9 1225 // CHECK3-NEXT: store i64 0, ptr [[TMP143]], align 8 1226 // CHECK3-NEXT: [[TMP144:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS28]], i32 0, i32 10 1227 // CHECK3-NEXT: store [3 x i32] [[TMP132]], ptr [[TMP144]], align 4 1228 // CHECK3-NEXT: [[TMP145:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS28]], i32 0, i32 11 1229 // CHECK3-NEXT: store [3 x i32] [[TMP133]], ptr [[TMP145]], align 4 1230 // CHECK3-NEXT: [[TMP146:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS28]], i32 0, i32 12 1231 // CHECK3-NEXT: store i32 0, ptr [[TMP146]], align 4 1232 // CHECK3-NEXT: [[TMP147:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 -1, i32 [[ADD]], i32 [[TMP131]], ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l62.region_id, ptr [[KERNEL_ARGS28]]) 1233 // CHECK3-NEXT: [[TMP148:%.*]] = icmp ne i32 [[TMP147]], 0 1234 // CHECK3-NEXT: br i1 [[TMP148]], label [[OMP_OFFLOAD_FAILED29:%.*]], label [[OMP_OFFLOAD_CONT30:%.*]] 1235 // CHECK3: omp_offload.failed29: 1236 // CHECK3-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l62(i32 [[TMP103]], i32 [[TMP105]], ptr @Gblb, i32 [[TMP107]], i32 [[TMP109]]) #[[ATTR2]] 1237 // CHECK3-NEXT: br label [[OMP_OFFLOAD_CONT30]] 1238 // CHECK3: omp_offload.cont30: 1239 // CHECK3-NEXT: [[TMP149:%.*]] = load ptr, ptr @Gblc, align 4 1240 // CHECK3-NEXT: store ptr [[TMP149]], ptr [[TMP]], align 4 1241 // CHECK3-NEXT: [[TMP150:%.*]] = load i32, ptr @Gbla, align 4 1242 // CHECK3-NEXT: store i32 [[TMP150]], ptr [[GBLC_CASTED]], align 4 1243 // CHECK3-NEXT: [[TMP151:%.*]] = load i32, ptr [[GBLC_CASTED]], align 4 1244 // CHECK3-NEXT: [[TMP152:%.*]] = load i32, ptr [[COMP]], align 4 1245 // CHECK3-NEXT: store i32 [[TMP152]], ptr [[COMP_CASTED31]], align 4 1246 // CHECK3-NEXT: [[TMP153:%.*]] = load i32, ptr [[COMP_CASTED31]], align 4 1247 // CHECK3-NEXT: [[TMP154:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS32]], i32 0, i32 0 1248 // CHECK3-NEXT: store i32 [[TMP151]], ptr [[TMP154]], align 4 1249 // CHECK3-NEXT: [[TMP155:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS33]], i32 0, i32 0 1250 // CHECK3-NEXT: store i32 [[TMP151]], ptr [[TMP155]], align 4 1251 // CHECK3-NEXT: [[TMP156:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS34]], i32 0, i32 0 1252 // CHECK3-NEXT: store ptr null, ptr [[TMP156]], align 4 1253 // CHECK3-NEXT: [[TMP157:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS32]], i32 0, i32 1 1254 // CHECK3-NEXT: store i32 [[TMP153]], ptr [[TMP157]], align 4 1255 // CHECK3-NEXT: [[TMP158:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS33]], i32 0, i32 1 1256 // CHECK3-NEXT: store i32 [[TMP153]], ptr [[TMP158]], align 4 1257 // CHECK3-NEXT: [[TMP159:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS34]], i32 0, i32 1 1258 // CHECK3-NEXT: store ptr null, ptr [[TMP159]], align 4 1259 // CHECK3-NEXT: [[TMP160:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS32]], i32 0, i32 0 1260 // CHECK3-NEXT: [[TMP161:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS33]], i32 0, i32 0 1261 // CHECK3-NEXT: [[TMP162:%.*]] = load ptr, ptr @Gblc, align 4 1262 // CHECK3-NEXT: store ptr [[TMP162]], ptr [[_TMP35]], align 4 1263 // CHECK3-NEXT: [[TMP163:%.*]] = load i32, ptr @Gbla, align 4 1264 // CHECK3-NEXT: [[ADD36:%.*]] = add nsw i32 [[TMP163]], 1 1265 // CHECK3-NEXT: [[TMP164:%.*]] = load ptr, ptr @Gblc, align 4 1266 // CHECK3-NEXT: store ptr [[TMP164]], ptr [[_TMP37]], align 4 1267 // CHECK3-NEXT: [[TMP165:%.*]] = load i32, ptr @Gbla, align 4 1268 // CHECK3-NEXT: [[ADD38:%.*]] = add nsw i32 [[TMP165]], 2 1269 // CHECK3-NEXT: [[TMP166:%.*]] = insertvalue [3 x i32] zeroinitializer, i32 [[ADD36]], 0 1270 // CHECK3-NEXT: [[TMP167:%.*]] = insertvalue [3 x i32] zeroinitializer, i32 [[ADD38]], 0 1271 // CHECK3-NEXT: [[TMP168:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS39]], i32 0, i32 0 1272 // CHECK3-NEXT: store i32 3, ptr [[TMP168]], align 4 1273 // CHECK3-NEXT: [[TMP169:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS39]], i32 0, i32 1 1274 // CHECK3-NEXT: store i32 2, ptr [[TMP169]], align 4 1275 // CHECK3-NEXT: [[TMP170:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS39]], i32 0, i32 2 1276 // CHECK3-NEXT: store ptr [[TMP160]], ptr [[TMP170]], align 4 1277 // CHECK3-NEXT: [[TMP171:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS39]], i32 0, i32 3 1278 // CHECK3-NEXT: store ptr [[TMP161]], ptr [[TMP171]], align 4 1279 // CHECK3-NEXT: [[TMP172:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS39]], i32 0, i32 4 1280 // CHECK3-NEXT: store ptr @.offload_sizes.9, ptr [[TMP172]], align 4 1281 // CHECK3-NEXT: [[TMP173:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS39]], i32 0, i32 5 1282 // CHECK3-NEXT: store ptr @.offload_maptypes.10, ptr [[TMP173]], align 4 1283 // CHECK3-NEXT: [[TMP174:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS39]], i32 0, i32 6 1284 // CHECK3-NEXT: store ptr null, ptr [[TMP174]], align 4 1285 // CHECK3-NEXT: [[TMP175:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS39]], i32 0, i32 7 1286 // CHECK3-NEXT: store ptr null, ptr [[TMP175]], align 4 1287 // CHECK3-NEXT: [[TMP176:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS39]], i32 0, i32 8 1288 // CHECK3-NEXT: store i64 0, ptr [[TMP176]], align 8 1289 // CHECK3-NEXT: [[TMP177:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS39]], i32 0, i32 9 1290 // CHECK3-NEXT: store i64 0, ptr [[TMP177]], align 8 1291 // CHECK3-NEXT: [[TMP178:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS39]], i32 0, i32 10 1292 // CHECK3-NEXT: store [3 x i32] [[TMP166]], ptr [[TMP178]], align 4 1293 // CHECK3-NEXT: [[TMP179:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS39]], i32 0, i32 11 1294 // CHECK3-NEXT: store [3 x i32] [[TMP167]], ptr [[TMP179]], align 4 1295 // CHECK3-NEXT: [[TMP180:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS39]], i32 0, i32 12 1296 // CHECK3-NEXT: store i32 0, ptr [[TMP180]], align 4 1297 // CHECK3-NEXT: [[TMP181:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 -1, i32 [[ADD36]], i32 [[ADD38]], ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l71.region_id, ptr [[KERNEL_ARGS39]]) 1298 // CHECK3-NEXT: [[TMP182:%.*]] = icmp ne i32 [[TMP181]], 0 1299 // CHECK3-NEXT: br i1 [[TMP182]], label [[OMP_OFFLOAD_FAILED40:%.*]], label [[OMP_OFFLOAD_CONT41:%.*]] 1300 // CHECK3: omp_offload.failed40: 1301 // CHECK3-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l71(i32 [[TMP151]], i32 [[TMP153]]) #[[ATTR2]] 1302 // CHECK3-NEXT: br label [[OMP_OFFLOAD_CONT41]] 1303 // CHECK3: omp_offload.cont41: 1304 // CHECK3-NEXT: [[TMP183:%.*]] = load i32, ptr [[COMP]], align 4 1305 // CHECK3-NEXT: ret i32 [[TMP183]] 1306 // 1307 // 1308 // CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l31 1309 // CHECK3-SAME: (i32 noundef [[COMP:%.*]]) #[[ATTR1:[0-9]+]] { 1310 // CHECK3-NEXT: entry: 1311 // CHECK3-NEXT: [[COMP_ADDR:%.*]] = alloca i32, align 4 1312 // CHECK3-NEXT: store i32 [[COMP]], ptr [[COMP_ADDR]], align 4 1313 // CHECK3-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB1]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l31.omp_outlined, ptr [[COMP_ADDR]]) 1314 // CHECK3-NEXT: ret void 1315 // 1316 // 1317 // CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l31.omp_outlined 1318 // CHECK3-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[COMP:%.*]]) #[[ATTR1]] { 1319 // CHECK3-NEXT: entry: 1320 // CHECK3-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 4 1321 // CHECK3-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 4 1322 // CHECK3-NEXT: [[COMP_ADDR:%.*]] = alloca ptr, align 4 1323 // CHECK3-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 4 1324 // CHECK3-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 4 1325 // CHECK3-NEXT: store ptr [[COMP]], ptr [[COMP_ADDR]], align 4 1326 // CHECK3-NEXT: [[TMP0:%.*]] = load ptr, ptr [[COMP_ADDR]], align 4 1327 // CHECK3-NEXT: [[TMP1:%.*]] = load i32, ptr [[TMP0]], align 4 1328 // CHECK3-NEXT: [[INC:%.*]] = add nsw i32 [[TMP1]], 1 1329 // CHECK3-NEXT: store i32 [[INC]], ptr [[TMP0]], align 4 1330 // CHECK3-NEXT: ret void 1331 // 1332 // 1333 // CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l37 1334 // CHECK3-SAME: (i32 noundef [[COMP:%.*]]) #[[ATTR1]] { 1335 // CHECK3-NEXT: entry: 1336 // CHECK3-NEXT: [[COMP_ADDR:%.*]] = alloca i32, align 4 1337 // CHECK3-NEXT: store i32 [[COMP]], ptr [[COMP_ADDR]], align 4 1338 // CHECK3-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB1]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l37.omp_outlined, ptr [[COMP_ADDR]]) 1339 // CHECK3-NEXT: ret void 1340 // 1341 // 1342 // CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l37.omp_outlined 1343 // CHECK3-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[COMP:%.*]]) #[[ATTR1]] { 1344 // CHECK3-NEXT: entry: 1345 // CHECK3-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 4 1346 // CHECK3-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 4 1347 // CHECK3-NEXT: [[COMP_ADDR:%.*]] = alloca ptr, align 4 1348 // CHECK3-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 4 1349 // CHECK3-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 4 1350 // CHECK3-NEXT: store ptr [[COMP]], ptr [[COMP_ADDR]], align 4 1351 // CHECK3-NEXT: [[TMP0:%.*]] = load ptr, ptr [[COMP_ADDR]], align 4 1352 // CHECK3-NEXT: [[TMP1:%.*]] = load i32, ptr [[TMP0]], align 4 1353 // CHECK3-NEXT: [[INC:%.*]] = add nsw i32 [[TMP1]], 1 1354 // CHECK3-NEXT: store i32 [[INC]], ptr [[TMP0]], align 4 1355 // CHECK3-NEXT: ret void 1356 // 1357 // 1358 // CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l46 1359 // CHECK3-SAME: (i32 noundef [[LA:%.*]], i32 noundef [[COMP:%.*]]) #[[ATTR1]] { 1360 // CHECK3-NEXT: entry: 1361 // CHECK3-NEXT: [[LA_ADDR:%.*]] = alloca i32, align 4 1362 // CHECK3-NEXT: [[COMP_ADDR:%.*]] = alloca i32, align 4 1363 // CHECK3-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]]) 1364 // CHECK3-NEXT: store i32 [[LA]], ptr [[LA_ADDR]], align 4 1365 // CHECK3-NEXT: store i32 [[COMP]], ptr [[COMP_ADDR]], align 4 1366 // CHECK3-NEXT: [[TMP1:%.*]] = load i32, ptr [[LA_ADDR]], align 4 1367 // CHECK3-NEXT: call void @__kmpc_push_num_teams(ptr @[[GLOB1]], i32 [[TMP0]], i32 [[TMP1]], i32 0) 1368 // CHECK3-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB1]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l46.omp_outlined, ptr [[COMP_ADDR]]) 1369 // CHECK3-NEXT: ret void 1370 // 1371 // 1372 // CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l46.omp_outlined 1373 // CHECK3-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[COMP:%.*]]) #[[ATTR1]] { 1374 // CHECK3-NEXT: entry: 1375 // CHECK3-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 4 1376 // CHECK3-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 4 1377 // CHECK3-NEXT: [[COMP_ADDR:%.*]] = alloca ptr, align 4 1378 // CHECK3-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 4 1379 // CHECK3-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 4 1380 // CHECK3-NEXT: store ptr [[COMP]], ptr [[COMP_ADDR]], align 4 1381 // CHECK3-NEXT: [[TMP0:%.*]] = load ptr, ptr [[COMP_ADDR]], align 4 1382 // CHECK3-NEXT: [[TMP1:%.*]] = load i32, ptr [[TMP0]], align 4 1383 // CHECK3-NEXT: [[INC:%.*]] = add nsw i32 [[TMP1]], 1 1384 // CHECK3-NEXT: store i32 [[INC]], ptr [[TMP0]], align 4 1385 // CHECK3-NEXT: ret void 1386 // 1387 // 1388 // CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l53 1389 // CHECK3-SAME: (i32 noundef [[LA:%.*]], i32 noundef [[COMP:%.*]]) #[[ATTR1]] { 1390 // CHECK3-NEXT: entry: 1391 // CHECK3-NEXT: [[LA_ADDR:%.*]] = alloca i32, align 4 1392 // CHECK3-NEXT: [[COMP_ADDR:%.*]] = alloca i32, align 4 1393 // CHECK3-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]]) 1394 // CHECK3-NEXT: store i32 [[LA]], ptr [[LA_ADDR]], align 4 1395 // CHECK3-NEXT: store i32 [[COMP]], ptr [[COMP_ADDR]], align 4 1396 // CHECK3-NEXT: [[TMP1:%.*]] = load i32, ptr [[LA_ADDR]], align 4 1397 // CHECK3-NEXT: call void @__kmpc_push_num_teams(ptr @[[GLOB1]], i32 [[TMP0]], i32 0, i32 [[TMP1]]) 1398 // CHECK3-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB1]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l53.omp_outlined, ptr [[COMP_ADDR]]) 1399 // CHECK3-NEXT: ret void 1400 // 1401 // 1402 // CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l53.omp_outlined 1403 // CHECK3-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[COMP:%.*]]) #[[ATTR1]] { 1404 // CHECK3-NEXT: entry: 1405 // CHECK3-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 4 1406 // CHECK3-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 4 1407 // CHECK3-NEXT: [[COMP_ADDR:%.*]] = alloca ptr, align 4 1408 // CHECK3-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 4 1409 // CHECK3-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 4 1410 // CHECK3-NEXT: store ptr [[COMP]], ptr [[COMP_ADDR]], align 4 1411 // CHECK3-NEXT: [[TMP0:%.*]] = load ptr, ptr [[COMP_ADDR]], align 4 1412 // CHECK3-NEXT: [[TMP1:%.*]] = load i32, ptr [[TMP0]], align 4 1413 // CHECK3-NEXT: [[INC:%.*]] = add nsw i32 [[TMP1]], 1 1414 // CHECK3-NEXT: store i32 [[INC]], ptr [[TMP0]], align 4 1415 // CHECK3-NEXT: ret void 1416 // 1417 // 1418 // CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l62 1419 // CHECK3-SAME: (i32 noundef [[GBLA:%.*]], i32 noundef [[A:%.*]], ptr noundef nonnull align 4 dereferenceable(8) [[GBLB:%.*]], i32 noundef [[LC:%.*]], i32 noundef [[COMP:%.*]]) #[[ATTR1]] { 1420 // CHECK3-NEXT: entry: 1421 // CHECK3-NEXT: [[GBLA_ADDR:%.*]] = alloca i32, align 4 1422 // CHECK3-NEXT: [[A_ADDR:%.*]] = alloca i32, align 4 1423 // CHECK3-NEXT: [[GBLB_ADDR:%.*]] = alloca ptr, align 4 1424 // CHECK3-NEXT: [[LC_ADDR:%.*]] = alloca i32, align 4 1425 // CHECK3-NEXT: [[COMP_ADDR:%.*]] = alloca i32, align 4 1426 // CHECK3-NEXT: [[GBLB1:%.*]] = alloca i64, align 8 1427 // CHECK3-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]]) 1428 // CHECK3-NEXT: store i32 [[GBLA]], ptr [[GBLA_ADDR]], align 4 1429 // CHECK3-NEXT: store i32 [[A]], ptr [[A_ADDR]], align 4 1430 // CHECK3-NEXT: store ptr [[GBLB]], ptr [[GBLB_ADDR]], align 4 1431 // CHECK3-NEXT: store i32 [[LC]], ptr [[LC_ADDR]], align 4 1432 // CHECK3-NEXT: store i32 [[COMP]], ptr [[COMP_ADDR]], align 4 1433 // CHECK3-NEXT: [[TMP1:%.*]] = load ptr, ptr [[GBLB_ADDR]], align 4 1434 // CHECK3-NEXT: [[TMP2:%.*]] = load i64, ptr [[TMP1]], align 8 1435 // CHECK3-NEXT: store i64 [[TMP2]], ptr [[GBLB1]], align 8 1436 // CHECK3-NEXT: [[TMP3:%.*]] = load i32, ptr [[GBLA_ADDR]], align 4 1437 // CHECK3-NEXT: [[TMP4:%.*]] = load i32, ptr [[A_ADDR]], align 4 1438 // CHECK3-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP3]], [[TMP4]] 1439 // CHECK3-NEXT: [[TMP5:%.*]] = load i64, ptr [[GBLB1]], align 8 1440 // CHECK3-NEXT: [[TMP6:%.*]] = load float, ptr [[LC_ADDR]], align 4 1441 // CHECK3-NEXT: [[CONV:%.*]] = fptosi float [[TMP6]] to i64 1442 // CHECK3-NEXT: [[ADD2:%.*]] = add nsw i64 [[TMP5]], [[CONV]] 1443 // CHECK3-NEXT: [[TMP7:%.*]] = trunc i64 [[ADD2]] to i32 1444 // CHECK3-NEXT: call void @__kmpc_push_num_teams(ptr @[[GLOB1]], i32 [[TMP0]], i32 [[ADD]], i32 [[TMP7]]) 1445 // CHECK3-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB1]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l62.omp_outlined, ptr [[COMP_ADDR]]) 1446 // CHECK3-NEXT: ret void 1447 // 1448 // 1449 // CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l62.omp_outlined 1450 // CHECK3-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[COMP:%.*]]) #[[ATTR1]] { 1451 // CHECK3-NEXT: entry: 1452 // CHECK3-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 4 1453 // CHECK3-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 4 1454 // CHECK3-NEXT: [[COMP_ADDR:%.*]] = alloca ptr, align 4 1455 // CHECK3-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 4 1456 // CHECK3-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 4 1457 // CHECK3-NEXT: store ptr [[COMP]], ptr [[COMP_ADDR]], align 4 1458 // CHECK3-NEXT: [[TMP0:%.*]] = load ptr, ptr [[COMP_ADDR]], align 4 1459 // CHECK3-NEXT: [[TMP1:%.*]] = load i32, ptr [[TMP0]], align 4 1460 // CHECK3-NEXT: [[INC:%.*]] = add nsw i32 [[TMP1]], 1 1461 // CHECK3-NEXT: store i32 [[INC]], ptr [[TMP0]], align 4 1462 // CHECK3-NEXT: ret void 1463 // 1464 // 1465 // CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l71 1466 // CHECK3-SAME: (i32 noundef [[GBLC:%.*]], i32 noundef [[COMP:%.*]]) #[[ATTR1]] { 1467 // CHECK3-NEXT: entry: 1468 // CHECK3-NEXT: [[GBLC_ADDR:%.*]] = alloca i32, align 4 1469 // CHECK3-NEXT: [[COMP_ADDR:%.*]] = alloca i32, align 4 1470 // CHECK3-NEXT: [[TMP:%.*]] = alloca ptr, align 4 1471 // CHECK3-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]]) 1472 // CHECK3-NEXT: store i32 [[GBLC]], ptr [[GBLC_ADDR]], align 4 1473 // CHECK3-NEXT: store i32 [[COMP]], ptr [[COMP_ADDR]], align 4 1474 // CHECK3-NEXT: store ptr [[GBLC_ADDR]], ptr [[TMP]], align 4 1475 // CHECK3-NEXT: [[TMP1:%.*]] = load i32, ptr @Gbla, align 4 1476 // CHECK3-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP1]], 1 1477 // CHECK3-NEXT: [[TMP2:%.*]] = load i32, ptr @Gbla, align 4 1478 // CHECK3-NEXT: [[ADD1:%.*]] = add nsw i32 [[TMP2]], 2 1479 // CHECK3-NEXT: call void @__kmpc_push_num_teams(ptr @[[GLOB1]], i32 [[TMP0]], i32 [[ADD]], i32 [[ADD1]]) 1480 // CHECK3-NEXT: [[TMP3:%.*]] = load ptr, ptr [[TMP]], align 4 1481 // CHECK3-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB1]], i32 2, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l71.omp_outlined, ptr [[COMP_ADDR]], ptr [[TMP3]]) 1482 // CHECK3-NEXT: ret void 1483 // 1484 // 1485 // CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z27teams_argument_global_locali_l71.omp_outlined 1486 // CHECK3-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[COMP:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[GBLC:%.*]]) #[[ATTR1]] { 1487 // CHECK3-NEXT: entry: 1488 // CHECK3-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 4 1489 // CHECK3-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 4 1490 // CHECK3-NEXT: [[COMP_ADDR:%.*]] = alloca ptr, align 4 1491 // CHECK3-NEXT: [[GBLC_ADDR:%.*]] = alloca ptr, align 4 1492 // CHECK3-NEXT: [[TMP:%.*]] = alloca ptr, align 4 1493 // CHECK3-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 4 1494 // CHECK3-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 4 1495 // CHECK3-NEXT: store ptr [[COMP]], ptr [[COMP_ADDR]], align 4 1496 // CHECK3-NEXT: store ptr [[GBLC]], ptr [[GBLC_ADDR]], align 4 1497 // CHECK3-NEXT: [[TMP0:%.*]] = load ptr, ptr [[COMP_ADDR]], align 4 1498 // CHECK3-NEXT: [[TMP1:%.*]] = load ptr, ptr [[GBLC_ADDR]], align 4 1499 // CHECK3-NEXT: store ptr [[TMP1]], ptr [[TMP]], align 4 1500 // CHECK3-NEXT: [[TMP2:%.*]] = load i32, ptr @Gbla, align 4 1501 // CHECK3-NEXT: [[TMP3:%.*]] = load i32, ptr [[TMP0]], align 4 1502 // CHECK3-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP3]], [[TMP2]] 1503 // CHECK3-NEXT: store i32 [[ADD]], ptr [[TMP0]], align 4 1504 // CHECK3-NEXT: ret void 1505 // 1506 // 1507 // CHECK9-LABEL: define {{[^@]+}}@_Z18teams_template_argv 1508 // CHECK9-SAME: () #[[ATTR0:[0-9]+]] { 1509 // CHECK9-NEXT: entry: 1510 // CHECK9-NEXT: [[COMP:%.*]] = alloca i32, align 4 1511 // CHECK9-NEXT: [[LA:%.*]] = alloca [[STRUCT_SS:%.*]], align 4 1512 // CHECK9-NEXT: [[LB:%.*]] = alloca [[STRUCT_SS_0:%.*]], align 8 1513 // CHECK9-NEXT: [[COMP_CASTED:%.*]] = alloca i64, align 8 1514 // CHECK9-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [3 x ptr], align 8 1515 // CHECK9-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [3 x ptr], align 8 1516 // CHECK9-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [3 x ptr], align 8 1517 // CHECK9-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8 1518 // CHECK9-NEXT: [[COMP_CASTED1:%.*]] = alloca i64, align 8 1519 // CHECK9-NEXT: [[DOTOFFLOAD_BASEPTRS2:%.*]] = alloca [3 x ptr], align 8 1520 // CHECK9-NEXT: [[DOTOFFLOAD_PTRS3:%.*]] = alloca [3 x ptr], align 8 1521 // CHECK9-NEXT: [[DOTOFFLOAD_MAPPERS4:%.*]] = alloca [3 x ptr], align 8 1522 // CHECK9-NEXT: [[KERNEL_ARGS7:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 8 1523 // CHECK9-NEXT: store i32 1, ptr [[COMP]], align 4 1524 // CHECK9-NEXT: [[TMP0:%.*]] = load i32, ptr [[COMP]], align 4 1525 // CHECK9-NEXT: store i32 [[TMP0]], ptr [[COMP_CASTED]], align 4 1526 // CHECK9-NEXT: [[TMP1:%.*]] = load i64, ptr [[COMP_CASTED]], align 8 1527 // CHECK9-NEXT: [[TMP2:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 1528 // CHECK9-NEXT: store ptr @Gbla, ptr [[TMP2]], align 8 1529 // CHECK9-NEXT: [[TMP3:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0 1530 // CHECK9-NEXT: store ptr @Gbla, ptr [[TMP3]], align 8 1531 // CHECK9-NEXT: [[TMP4:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 0 1532 // CHECK9-NEXT: store ptr null, ptr [[TMP4]], align 8 1533 // CHECK9-NEXT: [[TMP5:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 1 1534 // CHECK9-NEXT: store ptr [[LA]], ptr [[TMP5]], align 8 1535 // CHECK9-NEXT: [[TMP6:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 1 1536 // CHECK9-NEXT: store ptr [[LA]], ptr [[TMP6]], align 8 1537 // CHECK9-NEXT: [[TMP7:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 1 1538 // CHECK9-NEXT: store ptr null, ptr [[TMP7]], align 8 1539 // CHECK9-NEXT: [[TMP8:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 2 1540 // CHECK9-NEXT: store i64 [[TMP1]], ptr [[TMP8]], align 8 1541 // CHECK9-NEXT: [[TMP9:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 2 1542 // CHECK9-NEXT: store i64 [[TMP1]], ptr [[TMP9]], align 8 1543 // CHECK9-NEXT: [[TMP10:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 2 1544 // CHECK9-NEXT: store ptr null, ptr [[TMP10]], align 8 1545 // CHECK9-NEXT: [[TMP11:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 1546 // CHECK9-NEXT: [[TMP12:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0 1547 // CHECK9-NEXT: [[TMP13:%.*]] = load i32, ptr @Gbla, align 4 1548 // CHECK9-NEXT: [[B:%.*]] = getelementptr inbounds nuw [[STRUCT_SS]], ptr [[LA]], i32 0, i32 1 1549 // CHECK9-NEXT: [[TMP14:%.*]] = load float, ptr [[B]], align 4 1550 // CHECK9-NEXT: [[CONV:%.*]] = fptosi float [[TMP14]] to i64 1551 // CHECK9-NEXT: [[TMP15:%.*]] = trunc i64 [[CONV]] to i32 1552 // CHECK9-NEXT: [[TMP16:%.*]] = insertvalue [3 x i32] zeroinitializer, i32 [[TMP13]], 0 1553 // CHECK9-NEXT: [[TMP17:%.*]] = insertvalue [3 x i32] zeroinitializer, i32 [[TMP15]], 0 1554 // CHECK9-NEXT: [[TMP18:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0 1555 // CHECK9-NEXT: store i32 3, ptr [[TMP18]], align 4 1556 // CHECK9-NEXT: [[TMP19:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1 1557 // CHECK9-NEXT: store i32 3, ptr [[TMP19]], align 4 1558 // CHECK9-NEXT: [[TMP20:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2 1559 // CHECK9-NEXT: store ptr [[TMP11]], ptr [[TMP20]], align 8 1560 // CHECK9-NEXT: [[TMP21:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3 1561 // CHECK9-NEXT: store ptr [[TMP12]], ptr [[TMP21]], align 8 1562 // CHECK9-NEXT: [[TMP22:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4 1563 // CHECK9-NEXT: store ptr @.offload_sizes, ptr [[TMP22]], align 8 1564 // CHECK9-NEXT: [[TMP23:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5 1565 // CHECK9-NEXT: store ptr @.offload_maptypes, ptr [[TMP23]], align 8 1566 // CHECK9-NEXT: [[TMP24:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6 1567 // CHECK9-NEXT: store ptr null, ptr [[TMP24]], align 8 1568 // CHECK9-NEXT: [[TMP25:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7 1569 // CHECK9-NEXT: store ptr null, ptr [[TMP25]], align 8 1570 // CHECK9-NEXT: [[TMP26:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8 1571 // CHECK9-NEXT: store i64 0, ptr [[TMP26]], align 8 1572 // CHECK9-NEXT: [[TMP27:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9 1573 // CHECK9-NEXT: store i64 0, ptr [[TMP27]], align 8 1574 // CHECK9-NEXT: [[TMP28:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10 1575 // CHECK9-NEXT: store [3 x i32] [[TMP16]], ptr [[TMP28]], align 4 1576 // CHECK9-NEXT: [[TMP29:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11 1577 // CHECK9-NEXT: store [3 x i32] [[TMP17]], ptr [[TMP29]], align 4 1578 // CHECK9-NEXT: [[TMP30:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12 1579 // CHECK9-NEXT: store i32 0, ptr [[TMP30]], align 4 1580 // CHECK9-NEXT: [[TMP31:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1:[0-9]+]], i64 -1, i32 [[TMP13]], i32 [[TMP15]], ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z18teams_template_argv_l116.region_id, ptr [[KERNEL_ARGS]]) 1581 // CHECK9-NEXT: [[TMP32:%.*]] = icmp ne i32 [[TMP31]], 0 1582 // CHECK9-NEXT: br i1 [[TMP32]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]] 1583 // CHECK9: omp_offload.failed: 1584 // CHECK9-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z18teams_template_argv_l116(ptr @Gbla, ptr [[LA]], i64 [[TMP1]]) #[[ATTR2:[0-9]+]] 1585 // CHECK9-NEXT: br label [[OMP_OFFLOAD_CONT]] 1586 // CHECK9: omp_offload.cont: 1587 // CHECK9-NEXT: [[TMP33:%.*]] = load i32, ptr [[COMP]], align 4 1588 // CHECK9-NEXT: store i32 [[TMP33]], ptr [[COMP_CASTED1]], align 4 1589 // CHECK9-NEXT: [[TMP34:%.*]] = load i64, ptr [[COMP_CASTED1]], align 8 1590 // CHECK9-NEXT: [[TMP35:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS2]], i32 0, i32 0 1591 // CHECK9-NEXT: store ptr [[LB]], ptr [[TMP35]], align 8 1592 // CHECK9-NEXT: [[TMP36:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS3]], i32 0, i32 0 1593 // CHECK9-NEXT: store ptr [[LB]], ptr [[TMP36]], align 8 1594 // CHECK9-NEXT: [[TMP37:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS4]], i64 0, i64 0 1595 // CHECK9-NEXT: store ptr null, ptr [[TMP37]], align 8 1596 // CHECK9-NEXT: [[TMP38:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS2]], i32 0, i32 1 1597 // CHECK9-NEXT: store ptr @Gblb, ptr [[TMP38]], align 8 1598 // CHECK9-NEXT: [[TMP39:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS3]], i32 0, i32 1 1599 // CHECK9-NEXT: store ptr @Gblb, ptr [[TMP39]], align 8 1600 // CHECK9-NEXT: [[TMP40:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS4]], i64 0, i64 1 1601 // CHECK9-NEXT: store ptr null, ptr [[TMP40]], align 8 1602 // CHECK9-NEXT: [[TMP41:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS2]], i32 0, i32 2 1603 // CHECK9-NEXT: store i64 [[TMP34]], ptr [[TMP41]], align 8 1604 // CHECK9-NEXT: [[TMP42:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS3]], i32 0, i32 2 1605 // CHECK9-NEXT: store i64 [[TMP34]], ptr [[TMP42]], align 8 1606 // CHECK9-NEXT: [[TMP43:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS4]], i64 0, i64 2 1607 // CHECK9-NEXT: store ptr null, ptr [[TMP43]], align 8 1608 // CHECK9-NEXT: [[TMP44:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS2]], i32 0, i32 0 1609 // CHECK9-NEXT: [[TMP45:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS3]], i32 0, i32 0 1610 // CHECK9-NEXT: [[B5:%.*]] = getelementptr inbounds nuw [[STRUCT_SS_0]], ptr [[LB]], i32 0, i32 1 1611 // CHECK9-NEXT: [[TMP46:%.*]] = load float, ptr [[B5]], align 8 1612 // CHECK9-NEXT: [[CONV6:%.*]] = fptosi float [[TMP46]] to i64 1613 // CHECK9-NEXT: [[TMP47:%.*]] = trunc i64 [[CONV6]] to i32 1614 // CHECK9-NEXT: [[TMP48:%.*]] = load i64, ptr @Gblb, align 8 1615 // CHECK9-NEXT: [[TMP49:%.*]] = trunc i64 [[TMP48]] to i32 1616 // CHECK9-NEXT: [[TMP50:%.*]] = insertvalue [3 x i32] zeroinitializer, i32 [[TMP47]], 0 1617 // CHECK9-NEXT: [[TMP51:%.*]] = insertvalue [3 x i32] zeroinitializer, i32 [[TMP49]], 0 1618 // CHECK9-NEXT: [[TMP52:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 0 1619 // CHECK9-NEXT: store i32 3, ptr [[TMP52]], align 4 1620 // CHECK9-NEXT: [[TMP53:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 1 1621 // CHECK9-NEXT: store i32 3, ptr [[TMP53]], align 4 1622 // CHECK9-NEXT: [[TMP54:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 2 1623 // CHECK9-NEXT: store ptr [[TMP44]], ptr [[TMP54]], align 8 1624 // CHECK9-NEXT: [[TMP55:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 3 1625 // CHECK9-NEXT: store ptr [[TMP45]], ptr [[TMP55]], align 8 1626 // CHECK9-NEXT: [[TMP56:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 4 1627 // CHECK9-NEXT: store ptr @.offload_sizes.1, ptr [[TMP56]], align 8 1628 // CHECK9-NEXT: [[TMP57:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 5 1629 // CHECK9-NEXT: store ptr @.offload_maptypes.2, ptr [[TMP57]], align 8 1630 // CHECK9-NEXT: [[TMP58:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 6 1631 // CHECK9-NEXT: store ptr null, ptr [[TMP58]], align 8 1632 // CHECK9-NEXT: [[TMP59:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 7 1633 // CHECK9-NEXT: store ptr null, ptr [[TMP59]], align 8 1634 // CHECK9-NEXT: [[TMP60:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 8 1635 // CHECK9-NEXT: store i64 0, ptr [[TMP60]], align 8 1636 // CHECK9-NEXT: [[TMP61:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 9 1637 // CHECK9-NEXT: store i64 0, ptr [[TMP61]], align 8 1638 // CHECK9-NEXT: [[TMP62:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 10 1639 // CHECK9-NEXT: store [3 x i32] [[TMP50]], ptr [[TMP62]], align 4 1640 // CHECK9-NEXT: [[TMP63:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 11 1641 // CHECK9-NEXT: store [3 x i32] [[TMP51]], ptr [[TMP63]], align 4 1642 // CHECK9-NEXT: [[TMP64:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 12 1643 // CHECK9-NEXT: store i32 0, ptr [[TMP64]], align 4 1644 // CHECK9-NEXT: [[TMP65:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 -1, i32 [[TMP47]], i32 [[TMP49]], ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z18teams_template_argv_l125.region_id, ptr [[KERNEL_ARGS7]]) 1645 // CHECK9-NEXT: [[TMP66:%.*]] = icmp ne i32 [[TMP65]], 0 1646 // CHECK9-NEXT: br i1 [[TMP66]], label [[OMP_OFFLOAD_FAILED8:%.*]], label [[OMP_OFFLOAD_CONT9:%.*]] 1647 // CHECK9: omp_offload.failed8: 1648 // CHECK9-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z18teams_template_argv_l125(ptr [[LB]], ptr @Gblb, i64 [[TMP34]]) #[[ATTR2]] 1649 // CHECK9-NEXT: br label [[OMP_OFFLOAD_CONT9]] 1650 // CHECK9: omp_offload.cont9: 1651 // CHECK9-NEXT: [[TMP67:%.*]] = load i32, ptr [[COMP]], align 4 1652 // CHECK9-NEXT: ret i32 [[TMP67]] 1653 // 1654 // 1655 // CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z18teams_template_argv_l116 1656 // CHECK9-SAME: (ptr noundef nonnull align 4 dereferenceable(8) [[GBLA:%.*]], ptr noundef nonnull align 4 dereferenceable(8) [[LA:%.*]], i64 noundef [[COMP:%.*]]) #[[ATTR1:[0-9]+]] { 1657 // CHECK9-NEXT: entry: 1658 // CHECK9-NEXT: [[GBLA_ADDR:%.*]] = alloca ptr, align 8 1659 // CHECK9-NEXT: [[LA_ADDR:%.*]] = alloca ptr, align 8 1660 // CHECK9-NEXT: [[COMP_ADDR:%.*]] = alloca i64, align 8 1661 // CHECK9-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]]) 1662 // CHECK9-NEXT: store ptr [[GBLA]], ptr [[GBLA_ADDR]], align 8 1663 // CHECK9-NEXT: store ptr [[LA]], ptr [[LA_ADDR]], align 8 1664 // CHECK9-NEXT: store i64 [[COMP]], ptr [[COMP_ADDR]], align 8 1665 // CHECK9-NEXT: [[TMP1:%.*]] = load ptr, ptr [[GBLA_ADDR]], align 8 1666 // CHECK9-NEXT: [[TMP2:%.*]] = load ptr, ptr [[LA_ADDR]], align 8 1667 // CHECK9-NEXT: [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_SS:%.*]], ptr [[TMP1]], i32 0, i32 0 1668 // CHECK9-NEXT: [[TMP3:%.*]] = load i32, ptr [[A]], align 4 1669 // CHECK9-NEXT: [[B:%.*]] = getelementptr inbounds nuw [[STRUCT_SS]], ptr [[TMP2]], i32 0, i32 1 1670 // CHECK9-NEXT: [[TMP4:%.*]] = load float, ptr [[B]], align 4 1671 // CHECK9-NEXT: [[CONV:%.*]] = fptosi float [[TMP4]] to i64 1672 // CHECK9-NEXT: [[TMP5:%.*]] = trunc i64 [[CONV]] to i32 1673 // CHECK9-NEXT: call void @__kmpc_push_num_teams(ptr @[[GLOB1]], i32 [[TMP0]], i32 [[TMP3]], i32 [[TMP5]]) 1674 // CHECK9-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB1]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z18teams_template_argv_l116.omp_outlined, ptr [[COMP_ADDR]]) 1675 // CHECK9-NEXT: ret void 1676 // 1677 // 1678 // CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z18teams_template_argv_l116.omp_outlined 1679 // CHECK9-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[COMP:%.*]]) #[[ATTR1]] { 1680 // CHECK9-NEXT: entry: 1681 // CHECK9-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8 1682 // CHECK9-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8 1683 // CHECK9-NEXT: [[COMP_ADDR:%.*]] = alloca ptr, align 8 1684 // CHECK9-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8 1685 // CHECK9-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8 1686 // CHECK9-NEXT: store ptr [[COMP]], ptr [[COMP_ADDR]], align 8 1687 // CHECK9-NEXT: [[TMP0:%.*]] = load ptr, ptr [[COMP_ADDR]], align 8 1688 // CHECK9-NEXT: [[TMP1:%.*]] = load i32, ptr [[TMP0]], align 4 1689 // CHECK9-NEXT: [[INC:%.*]] = add nsw i32 [[TMP1]], 1 1690 // CHECK9-NEXT: store i32 [[INC]], ptr [[TMP0]], align 4 1691 // CHECK9-NEXT: ret void 1692 // 1693 // 1694 // CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z18teams_template_argv_l125 1695 // CHECK9-SAME: (ptr noundef nonnull align 8 dereferenceable(16) [[LB:%.*]], ptr noundef nonnull align 8 dereferenceable(16) [[GBLB:%.*]], i64 noundef [[COMP:%.*]]) #[[ATTR1]] { 1696 // CHECK9-NEXT: entry: 1697 // CHECK9-NEXT: [[LB_ADDR:%.*]] = alloca ptr, align 8 1698 // CHECK9-NEXT: [[GBLB_ADDR:%.*]] = alloca ptr, align 8 1699 // CHECK9-NEXT: [[COMP_ADDR:%.*]] = alloca i64, align 8 1700 // CHECK9-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]]) 1701 // CHECK9-NEXT: store ptr [[LB]], ptr [[LB_ADDR]], align 8 1702 // CHECK9-NEXT: store ptr [[GBLB]], ptr [[GBLB_ADDR]], align 8 1703 // CHECK9-NEXT: store i64 [[COMP]], ptr [[COMP_ADDR]], align 8 1704 // CHECK9-NEXT: [[TMP1:%.*]] = load ptr, ptr [[LB_ADDR]], align 8 1705 // CHECK9-NEXT: [[TMP2:%.*]] = load ptr, ptr [[GBLB_ADDR]], align 8 1706 // CHECK9-NEXT: [[B:%.*]] = getelementptr inbounds nuw [[STRUCT_SS_0:%.*]], ptr [[TMP1]], i32 0, i32 1 1707 // CHECK9-NEXT: [[TMP3:%.*]] = load float, ptr [[B]], align 8 1708 // CHECK9-NEXT: [[CONV:%.*]] = fptosi float [[TMP3]] to i64 1709 // CHECK9-NEXT: [[TMP4:%.*]] = trunc i64 [[CONV]] to i32 1710 // CHECK9-NEXT: [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_SS_0]], ptr [[TMP2]], i32 0, i32 0 1711 // CHECK9-NEXT: [[TMP5:%.*]] = load i64, ptr [[A]], align 8 1712 // CHECK9-NEXT: [[TMP6:%.*]] = trunc i64 [[TMP5]] to i32 1713 // CHECK9-NEXT: call void @__kmpc_push_num_teams(ptr @[[GLOB1]], i32 [[TMP0]], i32 [[TMP4]], i32 [[TMP6]]) 1714 // CHECK9-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB1]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z18teams_template_argv_l125.omp_outlined, ptr [[COMP_ADDR]]) 1715 // CHECK9-NEXT: ret void 1716 // 1717 // 1718 // CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z18teams_template_argv_l125.omp_outlined 1719 // CHECK9-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[COMP:%.*]]) #[[ATTR1]] { 1720 // CHECK9-NEXT: entry: 1721 // CHECK9-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8 1722 // CHECK9-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8 1723 // CHECK9-NEXT: [[COMP_ADDR:%.*]] = alloca ptr, align 8 1724 // CHECK9-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8 1725 // CHECK9-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8 1726 // CHECK9-NEXT: store ptr [[COMP]], ptr [[COMP_ADDR]], align 8 1727 // CHECK9-NEXT: [[TMP0:%.*]] = load ptr, ptr [[COMP_ADDR]], align 8 1728 // CHECK9-NEXT: [[TMP1:%.*]] = load i32, ptr [[TMP0]], align 4 1729 // CHECK9-NEXT: [[INC:%.*]] = add nsw i32 [[TMP1]], 1 1730 // CHECK9-NEXT: store i32 [[INC]], ptr [[TMP0]], align 4 1731 // CHECK9-NEXT: ret void 1732 // 1733 // 1734 // CHECK11-LABEL: define {{[^@]+}}@_Z18teams_template_argv 1735 // CHECK11-SAME: () #[[ATTR0:[0-9]+]] { 1736 // CHECK11-NEXT: entry: 1737 // CHECK11-NEXT: [[COMP:%.*]] = alloca i32, align 4 1738 // CHECK11-NEXT: [[LA:%.*]] = alloca [[STRUCT_SS:%.*]], align 4 1739 // CHECK11-NEXT: [[LB:%.*]] = alloca [[STRUCT_SS_0:%.*]], align 4 1740 // CHECK11-NEXT: [[COMP_CASTED:%.*]] = alloca i32, align 4 1741 // CHECK11-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [3 x ptr], align 4 1742 // CHECK11-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [3 x ptr], align 4 1743 // CHECK11-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [3 x ptr], align 4 1744 // CHECK11-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8 1745 // CHECK11-NEXT: [[COMP_CASTED1:%.*]] = alloca i32, align 4 1746 // CHECK11-NEXT: [[DOTOFFLOAD_BASEPTRS2:%.*]] = alloca [3 x ptr], align 4 1747 // CHECK11-NEXT: [[DOTOFFLOAD_PTRS3:%.*]] = alloca [3 x ptr], align 4 1748 // CHECK11-NEXT: [[DOTOFFLOAD_MAPPERS4:%.*]] = alloca [3 x ptr], align 4 1749 // CHECK11-NEXT: [[KERNEL_ARGS7:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 8 1750 // CHECK11-NEXT: store i32 1, ptr [[COMP]], align 4 1751 // CHECK11-NEXT: [[TMP0:%.*]] = load i32, ptr [[COMP]], align 4 1752 // CHECK11-NEXT: store i32 [[TMP0]], ptr [[COMP_CASTED]], align 4 1753 // CHECK11-NEXT: [[TMP1:%.*]] = load i32, ptr [[COMP_CASTED]], align 4 1754 // CHECK11-NEXT: [[TMP2:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 1755 // CHECK11-NEXT: store ptr @Gbla, ptr [[TMP2]], align 4 1756 // CHECK11-NEXT: [[TMP3:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0 1757 // CHECK11-NEXT: store ptr @Gbla, ptr [[TMP3]], align 4 1758 // CHECK11-NEXT: [[TMP4:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 0 1759 // CHECK11-NEXT: store ptr null, ptr [[TMP4]], align 4 1760 // CHECK11-NEXT: [[TMP5:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 1 1761 // CHECK11-NEXT: store ptr [[LA]], ptr [[TMP5]], align 4 1762 // CHECK11-NEXT: [[TMP6:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 1 1763 // CHECK11-NEXT: store ptr [[LA]], ptr [[TMP6]], align 4 1764 // CHECK11-NEXT: [[TMP7:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 1 1765 // CHECK11-NEXT: store ptr null, ptr [[TMP7]], align 4 1766 // CHECK11-NEXT: [[TMP8:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 2 1767 // CHECK11-NEXT: store i32 [[TMP1]], ptr [[TMP8]], align 4 1768 // CHECK11-NEXT: [[TMP9:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 2 1769 // CHECK11-NEXT: store i32 [[TMP1]], ptr [[TMP9]], align 4 1770 // CHECK11-NEXT: [[TMP10:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 2 1771 // CHECK11-NEXT: store ptr null, ptr [[TMP10]], align 4 1772 // CHECK11-NEXT: [[TMP11:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 1773 // CHECK11-NEXT: [[TMP12:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0 1774 // CHECK11-NEXT: [[TMP13:%.*]] = load i32, ptr @Gbla, align 4 1775 // CHECK11-NEXT: [[B:%.*]] = getelementptr inbounds nuw [[STRUCT_SS]], ptr [[LA]], i32 0, i32 1 1776 // CHECK11-NEXT: [[TMP14:%.*]] = load float, ptr [[B]], align 4 1777 // CHECK11-NEXT: [[CONV:%.*]] = fptosi float [[TMP14]] to i64 1778 // CHECK11-NEXT: [[TMP15:%.*]] = trunc i64 [[CONV]] to i32 1779 // CHECK11-NEXT: [[TMP16:%.*]] = insertvalue [3 x i32] zeroinitializer, i32 [[TMP13]], 0 1780 // CHECK11-NEXT: [[TMP17:%.*]] = insertvalue [3 x i32] zeroinitializer, i32 [[TMP15]], 0 1781 // CHECK11-NEXT: [[TMP18:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0 1782 // CHECK11-NEXT: store i32 3, ptr [[TMP18]], align 4 1783 // CHECK11-NEXT: [[TMP19:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1 1784 // CHECK11-NEXT: store i32 3, ptr [[TMP19]], align 4 1785 // CHECK11-NEXT: [[TMP20:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2 1786 // CHECK11-NEXT: store ptr [[TMP11]], ptr [[TMP20]], align 4 1787 // CHECK11-NEXT: [[TMP21:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3 1788 // CHECK11-NEXT: store ptr [[TMP12]], ptr [[TMP21]], align 4 1789 // CHECK11-NEXT: [[TMP22:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4 1790 // CHECK11-NEXT: store ptr @.offload_sizes, ptr [[TMP22]], align 4 1791 // CHECK11-NEXT: [[TMP23:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5 1792 // CHECK11-NEXT: store ptr @.offload_maptypes, ptr [[TMP23]], align 4 1793 // CHECK11-NEXT: [[TMP24:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6 1794 // CHECK11-NEXT: store ptr null, ptr [[TMP24]], align 4 1795 // CHECK11-NEXT: [[TMP25:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7 1796 // CHECK11-NEXT: store ptr null, ptr [[TMP25]], align 4 1797 // CHECK11-NEXT: [[TMP26:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8 1798 // CHECK11-NEXT: store i64 0, ptr [[TMP26]], align 8 1799 // CHECK11-NEXT: [[TMP27:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9 1800 // CHECK11-NEXT: store i64 0, ptr [[TMP27]], align 8 1801 // CHECK11-NEXT: [[TMP28:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10 1802 // CHECK11-NEXT: store [3 x i32] [[TMP16]], ptr [[TMP28]], align 4 1803 // CHECK11-NEXT: [[TMP29:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11 1804 // CHECK11-NEXT: store [3 x i32] [[TMP17]], ptr [[TMP29]], align 4 1805 // CHECK11-NEXT: [[TMP30:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12 1806 // CHECK11-NEXT: store i32 0, ptr [[TMP30]], align 4 1807 // CHECK11-NEXT: [[TMP31:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1:[0-9]+]], i64 -1, i32 [[TMP13]], i32 [[TMP15]], ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z18teams_template_argv_l116.region_id, ptr [[KERNEL_ARGS]]) 1808 // CHECK11-NEXT: [[TMP32:%.*]] = icmp ne i32 [[TMP31]], 0 1809 // CHECK11-NEXT: br i1 [[TMP32]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]] 1810 // CHECK11: omp_offload.failed: 1811 // CHECK11-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z18teams_template_argv_l116(ptr @Gbla, ptr [[LA]], i32 [[TMP1]]) #[[ATTR2:[0-9]+]] 1812 // CHECK11-NEXT: br label [[OMP_OFFLOAD_CONT]] 1813 // CHECK11: omp_offload.cont: 1814 // CHECK11-NEXT: [[TMP33:%.*]] = load i32, ptr [[COMP]], align 4 1815 // CHECK11-NEXT: store i32 [[TMP33]], ptr [[COMP_CASTED1]], align 4 1816 // CHECK11-NEXT: [[TMP34:%.*]] = load i32, ptr [[COMP_CASTED1]], align 4 1817 // CHECK11-NEXT: [[TMP35:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS2]], i32 0, i32 0 1818 // CHECK11-NEXT: store ptr [[LB]], ptr [[TMP35]], align 4 1819 // CHECK11-NEXT: [[TMP36:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS3]], i32 0, i32 0 1820 // CHECK11-NEXT: store ptr [[LB]], ptr [[TMP36]], align 4 1821 // CHECK11-NEXT: [[TMP37:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS4]], i32 0, i32 0 1822 // CHECK11-NEXT: store ptr null, ptr [[TMP37]], align 4 1823 // CHECK11-NEXT: [[TMP38:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS2]], i32 0, i32 1 1824 // CHECK11-NEXT: store ptr @Gblb, ptr [[TMP38]], align 4 1825 // CHECK11-NEXT: [[TMP39:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS3]], i32 0, i32 1 1826 // CHECK11-NEXT: store ptr @Gblb, ptr [[TMP39]], align 4 1827 // CHECK11-NEXT: [[TMP40:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS4]], i32 0, i32 1 1828 // CHECK11-NEXT: store ptr null, ptr [[TMP40]], align 4 1829 // CHECK11-NEXT: [[TMP41:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS2]], i32 0, i32 2 1830 // CHECK11-NEXT: store i32 [[TMP34]], ptr [[TMP41]], align 4 1831 // CHECK11-NEXT: [[TMP42:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS3]], i32 0, i32 2 1832 // CHECK11-NEXT: store i32 [[TMP34]], ptr [[TMP42]], align 4 1833 // CHECK11-NEXT: [[TMP43:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS4]], i32 0, i32 2 1834 // CHECK11-NEXT: store ptr null, ptr [[TMP43]], align 4 1835 // CHECK11-NEXT: [[TMP44:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS2]], i32 0, i32 0 1836 // CHECK11-NEXT: [[TMP45:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS3]], i32 0, i32 0 1837 // CHECK11-NEXT: [[B5:%.*]] = getelementptr inbounds nuw [[STRUCT_SS_0]], ptr [[LB]], i32 0, i32 1 1838 // CHECK11-NEXT: [[TMP46:%.*]] = load float, ptr [[B5]], align 4 1839 // CHECK11-NEXT: [[CONV6:%.*]] = fptosi float [[TMP46]] to i64 1840 // CHECK11-NEXT: [[TMP47:%.*]] = trunc i64 [[CONV6]] to i32 1841 // CHECK11-NEXT: [[TMP48:%.*]] = load i64, ptr @Gblb, align 4 1842 // CHECK11-NEXT: [[TMP49:%.*]] = trunc i64 [[TMP48]] to i32 1843 // CHECK11-NEXT: [[TMP50:%.*]] = insertvalue [3 x i32] zeroinitializer, i32 [[TMP47]], 0 1844 // CHECK11-NEXT: [[TMP51:%.*]] = insertvalue [3 x i32] zeroinitializer, i32 [[TMP49]], 0 1845 // CHECK11-NEXT: [[TMP52:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 0 1846 // CHECK11-NEXT: store i32 3, ptr [[TMP52]], align 4 1847 // CHECK11-NEXT: [[TMP53:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 1 1848 // CHECK11-NEXT: store i32 3, ptr [[TMP53]], align 4 1849 // CHECK11-NEXT: [[TMP54:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 2 1850 // CHECK11-NEXT: store ptr [[TMP44]], ptr [[TMP54]], align 4 1851 // CHECK11-NEXT: [[TMP55:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 3 1852 // CHECK11-NEXT: store ptr [[TMP45]], ptr [[TMP55]], align 4 1853 // CHECK11-NEXT: [[TMP56:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 4 1854 // CHECK11-NEXT: store ptr @.offload_sizes.1, ptr [[TMP56]], align 4 1855 // CHECK11-NEXT: [[TMP57:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 5 1856 // CHECK11-NEXT: store ptr @.offload_maptypes.2, ptr [[TMP57]], align 4 1857 // CHECK11-NEXT: [[TMP58:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 6 1858 // CHECK11-NEXT: store ptr null, ptr [[TMP58]], align 4 1859 // CHECK11-NEXT: [[TMP59:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 7 1860 // CHECK11-NEXT: store ptr null, ptr [[TMP59]], align 4 1861 // CHECK11-NEXT: [[TMP60:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 8 1862 // CHECK11-NEXT: store i64 0, ptr [[TMP60]], align 8 1863 // CHECK11-NEXT: [[TMP61:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 9 1864 // CHECK11-NEXT: store i64 0, ptr [[TMP61]], align 8 1865 // CHECK11-NEXT: [[TMP62:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 10 1866 // CHECK11-NEXT: store [3 x i32] [[TMP50]], ptr [[TMP62]], align 4 1867 // CHECK11-NEXT: [[TMP63:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 11 1868 // CHECK11-NEXT: store [3 x i32] [[TMP51]], ptr [[TMP63]], align 4 1869 // CHECK11-NEXT: [[TMP64:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS7]], i32 0, i32 12 1870 // CHECK11-NEXT: store i32 0, ptr [[TMP64]], align 4 1871 // CHECK11-NEXT: [[TMP65:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 -1, i32 [[TMP47]], i32 [[TMP49]], ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z18teams_template_argv_l125.region_id, ptr [[KERNEL_ARGS7]]) 1872 // CHECK11-NEXT: [[TMP66:%.*]] = icmp ne i32 [[TMP65]], 0 1873 // CHECK11-NEXT: br i1 [[TMP66]], label [[OMP_OFFLOAD_FAILED8:%.*]], label [[OMP_OFFLOAD_CONT9:%.*]] 1874 // CHECK11: omp_offload.failed8: 1875 // CHECK11-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z18teams_template_argv_l125(ptr [[LB]], ptr @Gblb, i32 [[TMP34]]) #[[ATTR2]] 1876 // CHECK11-NEXT: br label [[OMP_OFFLOAD_CONT9]] 1877 // CHECK11: omp_offload.cont9: 1878 // CHECK11-NEXT: [[TMP67:%.*]] = load i32, ptr [[COMP]], align 4 1879 // CHECK11-NEXT: ret i32 [[TMP67]] 1880 // 1881 // 1882 // CHECK11-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z18teams_template_argv_l116 1883 // CHECK11-SAME: (ptr noundef nonnull align 4 dereferenceable(8) [[GBLA:%.*]], ptr noundef nonnull align 4 dereferenceable(8) [[LA:%.*]], i32 noundef [[COMP:%.*]]) #[[ATTR1:[0-9]+]] { 1884 // CHECK11-NEXT: entry: 1885 // CHECK11-NEXT: [[GBLA_ADDR:%.*]] = alloca ptr, align 4 1886 // CHECK11-NEXT: [[LA_ADDR:%.*]] = alloca ptr, align 4 1887 // CHECK11-NEXT: [[COMP_ADDR:%.*]] = alloca i32, align 4 1888 // CHECK11-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]]) 1889 // CHECK11-NEXT: store ptr [[GBLA]], ptr [[GBLA_ADDR]], align 4 1890 // CHECK11-NEXT: store ptr [[LA]], ptr [[LA_ADDR]], align 4 1891 // CHECK11-NEXT: store i32 [[COMP]], ptr [[COMP_ADDR]], align 4 1892 // CHECK11-NEXT: [[TMP1:%.*]] = load ptr, ptr [[GBLA_ADDR]], align 4 1893 // CHECK11-NEXT: [[TMP2:%.*]] = load ptr, ptr [[LA_ADDR]], align 4 1894 // CHECK11-NEXT: [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_SS:%.*]], ptr [[TMP1]], i32 0, i32 0 1895 // CHECK11-NEXT: [[TMP3:%.*]] = load i32, ptr [[A]], align 4 1896 // CHECK11-NEXT: [[B:%.*]] = getelementptr inbounds nuw [[STRUCT_SS]], ptr [[TMP2]], i32 0, i32 1 1897 // CHECK11-NEXT: [[TMP4:%.*]] = load float, ptr [[B]], align 4 1898 // CHECK11-NEXT: [[CONV:%.*]] = fptosi float [[TMP4]] to i64 1899 // CHECK11-NEXT: [[TMP5:%.*]] = trunc i64 [[CONV]] to i32 1900 // CHECK11-NEXT: call void @__kmpc_push_num_teams(ptr @[[GLOB1]], i32 [[TMP0]], i32 [[TMP3]], i32 [[TMP5]]) 1901 // CHECK11-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB1]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z18teams_template_argv_l116.omp_outlined, ptr [[COMP_ADDR]]) 1902 // CHECK11-NEXT: ret void 1903 // 1904 // 1905 // CHECK11-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z18teams_template_argv_l116.omp_outlined 1906 // CHECK11-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[COMP:%.*]]) #[[ATTR1]] { 1907 // CHECK11-NEXT: entry: 1908 // CHECK11-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 4 1909 // CHECK11-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 4 1910 // CHECK11-NEXT: [[COMP_ADDR:%.*]] = alloca ptr, align 4 1911 // CHECK11-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 4 1912 // CHECK11-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 4 1913 // CHECK11-NEXT: store ptr [[COMP]], ptr [[COMP_ADDR]], align 4 1914 // CHECK11-NEXT: [[TMP0:%.*]] = load ptr, ptr [[COMP_ADDR]], align 4 1915 // CHECK11-NEXT: [[TMP1:%.*]] = load i32, ptr [[TMP0]], align 4 1916 // CHECK11-NEXT: [[INC:%.*]] = add nsw i32 [[TMP1]], 1 1917 // CHECK11-NEXT: store i32 [[INC]], ptr [[TMP0]], align 4 1918 // CHECK11-NEXT: ret void 1919 // 1920 // 1921 // CHECK11-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z18teams_template_argv_l125 1922 // CHECK11-SAME: (ptr noundef nonnull align 4 dereferenceable(12) [[LB:%.*]], ptr noundef nonnull align 4 dereferenceable(12) [[GBLB:%.*]], i32 noundef [[COMP:%.*]]) #[[ATTR1]] { 1923 // CHECK11-NEXT: entry: 1924 // CHECK11-NEXT: [[LB_ADDR:%.*]] = alloca ptr, align 4 1925 // CHECK11-NEXT: [[GBLB_ADDR:%.*]] = alloca ptr, align 4 1926 // CHECK11-NEXT: [[COMP_ADDR:%.*]] = alloca i32, align 4 1927 // CHECK11-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]]) 1928 // CHECK11-NEXT: store ptr [[LB]], ptr [[LB_ADDR]], align 4 1929 // CHECK11-NEXT: store ptr [[GBLB]], ptr [[GBLB_ADDR]], align 4 1930 // CHECK11-NEXT: store i32 [[COMP]], ptr [[COMP_ADDR]], align 4 1931 // CHECK11-NEXT: [[TMP1:%.*]] = load ptr, ptr [[LB_ADDR]], align 4 1932 // CHECK11-NEXT: [[TMP2:%.*]] = load ptr, ptr [[GBLB_ADDR]], align 4 1933 // CHECK11-NEXT: [[B:%.*]] = getelementptr inbounds nuw [[STRUCT_SS_0:%.*]], ptr [[TMP1]], i32 0, i32 1 1934 // CHECK11-NEXT: [[TMP3:%.*]] = load float, ptr [[B]], align 4 1935 // CHECK11-NEXT: [[CONV:%.*]] = fptosi float [[TMP3]] to i64 1936 // CHECK11-NEXT: [[TMP4:%.*]] = trunc i64 [[CONV]] to i32 1937 // CHECK11-NEXT: [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_SS_0]], ptr [[TMP2]], i32 0, i32 0 1938 // CHECK11-NEXT: [[TMP5:%.*]] = load i64, ptr [[A]], align 4 1939 // CHECK11-NEXT: [[TMP6:%.*]] = trunc i64 [[TMP5]] to i32 1940 // CHECK11-NEXT: call void @__kmpc_push_num_teams(ptr @[[GLOB1]], i32 [[TMP0]], i32 [[TMP4]], i32 [[TMP6]]) 1941 // CHECK11-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB1]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z18teams_template_argv_l125.omp_outlined, ptr [[COMP_ADDR]]) 1942 // CHECK11-NEXT: ret void 1943 // 1944 // 1945 // CHECK11-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z18teams_template_argv_l125.omp_outlined 1946 // CHECK11-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[COMP:%.*]]) #[[ATTR1]] { 1947 // CHECK11-NEXT: entry: 1948 // CHECK11-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 4 1949 // CHECK11-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 4 1950 // CHECK11-NEXT: [[COMP_ADDR:%.*]] = alloca ptr, align 4 1951 // CHECK11-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 4 1952 // CHECK11-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 4 1953 // CHECK11-NEXT: store ptr [[COMP]], ptr [[COMP_ADDR]], align 4 1954 // CHECK11-NEXT: [[TMP0:%.*]] = load ptr, ptr [[COMP_ADDR]], align 4 1955 // CHECK11-NEXT: [[TMP1:%.*]] = load i32, ptr [[TMP0]], align 4 1956 // CHECK11-NEXT: [[INC:%.*]] = add nsw i32 [[TMP1]], 1 1957 // CHECK11-NEXT: store i32 [[INC]], ptr [[TMP0]], align 4 1958 // CHECK11-NEXT: ret void 1959 // 1960 // 1961 // CHECK17-LABEL: define {{[^@]+}}@_Z21teams_template_structv 1962 // CHECK17-SAME: () #[[ATTR0:[0-9]+]] { 1963 // CHECK17-NEXT: entry: 1964 // CHECK17-NEXT: [[V:%.*]] = alloca [[STRUCT_SS:%.*]], align 4 1965 // CHECK17-NEXT: [[CALL:%.*]] = call noundef signext i32 @_ZN2SSIiLi123ELx456EE3fooEv(ptr noundef nonnull align 4 dereferenceable(8) [[V]]) 1966 // CHECK17-NEXT: ret i32 [[CALL]] 1967 // 1968 // 1969 // CHECK17-LABEL: define {{[^@]+}}@_ZN2SSIiLi123ELx456EE3fooEv 1970 // CHECK17-SAME: (ptr noundef nonnull align 4 dereferenceable(8) [[THIS:%.*]]) #[[ATTR0]] comdat { 1971 // CHECK17-NEXT: entry: 1972 // CHECK17-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 8 1973 // CHECK17-NEXT: [[COMP:%.*]] = alloca i32, align 4 1974 // CHECK17-NEXT: [[COMP_CASTED:%.*]] = alloca i64, align 8 1975 // CHECK17-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [2 x ptr], align 8 1976 // CHECK17-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [2 x ptr], align 8 1977 // CHECK17-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [2 x ptr], align 8 1978 // CHECK17-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8 1979 // CHECK17-NEXT: [[COMP_CASTED3:%.*]] = alloca i64, align 8 1980 // CHECK17-NEXT: [[DOTOFFLOAD_BASEPTRS4:%.*]] = alloca [2 x ptr], align 8 1981 // CHECK17-NEXT: [[DOTOFFLOAD_PTRS5:%.*]] = alloca [2 x ptr], align 8 1982 // CHECK17-NEXT: [[DOTOFFLOAD_MAPPERS6:%.*]] = alloca [2 x ptr], align 8 1983 // CHECK17-NEXT: [[KERNEL_ARGS8:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 8 1984 // CHECK17-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 8 1985 // CHECK17-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8 1986 // CHECK17-NEXT: store i32 1, ptr [[COMP]], align 4 1987 // CHECK17-NEXT: [[TMP0:%.*]] = load i32, ptr [[COMP]], align 4 1988 // CHECK17-NEXT: store i32 [[TMP0]], ptr [[COMP_CASTED]], align 4 1989 // CHECK17-NEXT: [[TMP1:%.*]] = load i64, ptr [[COMP_CASTED]], align 8 1990 // CHECK17-NEXT: [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_SS:%.*]], ptr [[THIS1]], i32 0, i32 0 1991 // CHECK17-NEXT: [[TMP2:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 1992 // CHECK17-NEXT: store ptr [[THIS1]], ptr [[TMP2]], align 8 1993 // CHECK17-NEXT: [[TMP3:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0 1994 // CHECK17-NEXT: store ptr [[A]], ptr [[TMP3]], align 8 1995 // CHECK17-NEXT: [[TMP4:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 0 1996 // CHECK17-NEXT: store ptr null, ptr [[TMP4]], align 8 1997 // CHECK17-NEXT: [[TMP5:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 1 1998 // CHECK17-NEXT: store i64 [[TMP1]], ptr [[TMP5]], align 8 1999 // CHECK17-NEXT: [[TMP6:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 1 2000 // CHECK17-NEXT: store i64 [[TMP1]], ptr [[TMP6]], align 8 2001 // CHECK17-NEXT: [[TMP7:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 1 2002 // CHECK17-NEXT: store ptr null, ptr [[TMP7]], align 8 2003 // CHECK17-NEXT: [[TMP8:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 2004 // CHECK17-NEXT: [[TMP9:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0 2005 // CHECK17-NEXT: [[A2:%.*]] = getelementptr inbounds nuw [[STRUCT_SS]], ptr [[THIS1]], i32 0, i32 0 2006 // CHECK17-NEXT: [[TMP10:%.*]] = load i32, ptr [[A2]], align 4 2007 // CHECK17-NEXT: [[TMP11:%.*]] = insertvalue [3 x i32] zeroinitializer, i32 [[TMP10]], 0 2008 // CHECK17-NEXT: [[TMP12:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0 2009 // CHECK17-NEXT: store i32 3, ptr [[TMP12]], align 4 2010 // CHECK17-NEXT: [[TMP13:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1 2011 // CHECK17-NEXT: store i32 2, ptr [[TMP13]], align 4 2012 // CHECK17-NEXT: [[TMP14:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2 2013 // CHECK17-NEXT: store ptr [[TMP8]], ptr [[TMP14]], align 8 2014 // CHECK17-NEXT: [[TMP15:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3 2015 // CHECK17-NEXT: store ptr [[TMP9]], ptr [[TMP15]], align 8 2016 // CHECK17-NEXT: [[TMP16:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4 2017 // CHECK17-NEXT: store ptr @.offload_sizes, ptr [[TMP16]], align 8 2018 // CHECK17-NEXT: [[TMP17:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5 2019 // CHECK17-NEXT: store ptr @.offload_maptypes, ptr [[TMP17]], align 8 2020 // CHECK17-NEXT: [[TMP18:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6 2021 // CHECK17-NEXT: store ptr null, ptr [[TMP18]], align 8 2022 // CHECK17-NEXT: [[TMP19:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7 2023 // CHECK17-NEXT: store ptr null, ptr [[TMP19]], align 8 2024 // CHECK17-NEXT: [[TMP20:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8 2025 // CHECK17-NEXT: store i64 0, ptr [[TMP20]], align 8 2026 // CHECK17-NEXT: [[TMP21:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9 2027 // CHECK17-NEXT: store i64 0, ptr [[TMP21]], align 8 2028 // CHECK17-NEXT: [[TMP22:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10 2029 // CHECK17-NEXT: store [3 x i32] [[TMP11]], ptr [[TMP22]], align 4 2030 // CHECK17-NEXT: [[TMP23:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11 2031 // CHECK17-NEXT: store [3 x i32] [i32 123, i32 0, i32 0], ptr [[TMP23]], align 4 2032 // CHECK17-NEXT: [[TMP24:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12 2033 // CHECK17-NEXT: store i32 0, ptr [[TMP24]], align 4 2034 // CHECK17-NEXT: [[TMP25:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1:[0-9]+]], i64 -1, i32 [[TMP10]], i32 123, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l161.region_id, ptr [[KERNEL_ARGS]]) 2035 // CHECK17-NEXT: [[TMP26:%.*]] = icmp ne i32 [[TMP25]], 0 2036 // CHECK17-NEXT: br i1 [[TMP26]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]] 2037 // CHECK17: omp_offload.failed: 2038 // CHECK17-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l161(ptr [[THIS1]], i64 [[TMP1]]) #[[ATTR2:[0-9]+]] 2039 // CHECK17-NEXT: br label [[OMP_OFFLOAD_CONT]] 2040 // CHECK17: omp_offload.cont: 2041 // CHECK17-NEXT: [[TMP27:%.*]] = load i32, ptr [[COMP]], align 4 2042 // CHECK17-NEXT: store i32 [[TMP27]], ptr [[COMP_CASTED3]], align 4 2043 // CHECK17-NEXT: [[TMP28:%.*]] = load i64, ptr [[COMP_CASTED3]], align 8 2044 // CHECK17-NEXT: [[B:%.*]] = getelementptr inbounds nuw [[STRUCT_SS]], ptr [[THIS1]], i32 0, i32 1 2045 // CHECK17-NEXT: [[TMP29:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS4]], i32 0, i32 0 2046 // CHECK17-NEXT: store ptr [[THIS1]], ptr [[TMP29]], align 8 2047 // CHECK17-NEXT: [[TMP30:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS5]], i32 0, i32 0 2048 // CHECK17-NEXT: store ptr [[B]], ptr [[TMP30]], align 8 2049 // CHECK17-NEXT: [[TMP31:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS6]], i64 0, i64 0 2050 // CHECK17-NEXT: store ptr null, ptr [[TMP31]], align 8 2051 // CHECK17-NEXT: [[TMP32:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS4]], i32 0, i32 1 2052 // CHECK17-NEXT: store i64 [[TMP28]], ptr [[TMP32]], align 8 2053 // CHECK17-NEXT: [[TMP33:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS5]], i32 0, i32 1 2054 // CHECK17-NEXT: store i64 [[TMP28]], ptr [[TMP33]], align 8 2055 // CHECK17-NEXT: [[TMP34:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS6]], i64 0, i64 1 2056 // CHECK17-NEXT: store ptr null, ptr [[TMP34]], align 8 2057 // CHECK17-NEXT: [[TMP35:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS4]], i32 0, i32 0 2058 // CHECK17-NEXT: [[TMP36:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS5]], i32 0, i32 0 2059 // CHECK17-NEXT: [[B7:%.*]] = getelementptr inbounds nuw [[STRUCT_SS]], ptr [[THIS1]], i32 0, i32 1 2060 // CHECK17-NEXT: [[TMP37:%.*]] = load float, ptr [[B7]], align 4 2061 // CHECK17-NEXT: [[CONV:%.*]] = fptosi float [[TMP37]] to i32 2062 // CHECK17-NEXT: [[ADD:%.*]] = add nsw i32 [[CONV]], 123 2063 // CHECK17-NEXT: [[TMP38:%.*]] = insertvalue [3 x i32] zeroinitializer, i32 [[ADD]], 0 2064 // CHECK17-NEXT: [[TMP39:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS8]], i32 0, i32 0 2065 // CHECK17-NEXT: store i32 3, ptr [[TMP39]], align 4 2066 // CHECK17-NEXT: [[TMP40:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS8]], i32 0, i32 1 2067 // CHECK17-NEXT: store i32 2, ptr [[TMP40]], align 4 2068 // CHECK17-NEXT: [[TMP41:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS8]], i32 0, i32 2 2069 // CHECK17-NEXT: store ptr [[TMP35]], ptr [[TMP41]], align 8 2070 // CHECK17-NEXT: [[TMP42:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS8]], i32 0, i32 3 2071 // CHECK17-NEXT: store ptr [[TMP36]], ptr [[TMP42]], align 8 2072 // CHECK17-NEXT: [[TMP43:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS8]], i32 0, i32 4 2073 // CHECK17-NEXT: store ptr @.offload_sizes.1, ptr [[TMP43]], align 8 2074 // CHECK17-NEXT: [[TMP44:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS8]], i32 0, i32 5 2075 // CHECK17-NEXT: store ptr @.offload_maptypes.2, ptr [[TMP44]], align 8 2076 // CHECK17-NEXT: [[TMP45:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS8]], i32 0, i32 6 2077 // CHECK17-NEXT: store ptr null, ptr [[TMP45]], align 8 2078 // CHECK17-NEXT: [[TMP46:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS8]], i32 0, i32 7 2079 // CHECK17-NEXT: store ptr null, ptr [[TMP46]], align 8 2080 // CHECK17-NEXT: [[TMP47:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS8]], i32 0, i32 8 2081 // CHECK17-NEXT: store i64 0, ptr [[TMP47]], align 8 2082 // CHECK17-NEXT: [[TMP48:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS8]], i32 0, i32 9 2083 // CHECK17-NEXT: store i64 0, ptr [[TMP48]], align 8 2084 // CHECK17-NEXT: [[TMP49:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS8]], i32 0, i32 10 2085 // CHECK17-NEXT: store [3 x i32] [i32 456, i32 0, i32 0], ptr [[TMP49]], align 4 2086 // CHECK17-NEXT: [[TMP50:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS8]], i32 0, i32 11 2087 // CHECK17-NEXT: store [3 x i32] [[TMP38]], ptr [[TMP50]], align 4 2088 // CHECK17-NEXT: [[TMP51:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS8]], i32 0, i32 12 2089 // CHECK17-NEXT: store i32 0, ptr [[TMP51]], align 4 2090 // CHECK17-NEXT: [[TMP52:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 -1, i32 456, i32 [[ADD]], ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l169.region_id, ptr [[KERNEL_ARGS8]]) 2091 // CHECK17-NEXT: [[TMP53:%.*]] = icmp ne i32 [[TMP52]], 0 2092 // CHECK17-NEXT: br i1 [[TMP53]], label [[OMP_OFFLOAD_FAILED9:%.*]], label [[OMP_OFFLOAD_CONT10:%.*]] 2093 // CHECK17: omp_offload.failed9: 2094 // CHECK17-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l169(ptr [[THIS1]], i64 [[TMP28]]) #[[ATTR2]] 2095 // CHECK17-NEXT: br label [[OMP_OFFLOAD_CONT10]] 2096 // CHECK17: omp_offload.cont10: 2097 // CHECK17-NEXT: [[TMP54:%.*]] = load i32, ptr [[COMP]], align 4 2098 // CHECK17-NEXT: ret i32 [[TMP54]] 2099 // 2100 // 2101 // CHECK17-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l161 2102 // CHECK17-SAME: (ptr noundef [[THIS:%.*]], i64 noundef [[COMP:%.*]]) #[[ATTR1:[0-9]+]] { 2103 // CHECK17-NEXT: entry: 2104 // CHECK17-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 8 2105 // CHECK17-NEXT: [[COMP_ADDR:%.*]] = alloca i64, align 8 2106 // CHECK17-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]]) 2107 // CHECK17-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 8 2108 // CHECK17-NEXT: store i64 [[COMP]], ptr [[COMP_ADDR]], align 8 2109 // CHECK17-NEXT: [[TMP1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8 2110 // CHECK17-NEXT: [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_SS:%.*]], ptr [[TMP1]], i32 0, i32 0 2111 // CHECK17-NEXT: [[TMP2:%.*]] = load i32, ptr [[A]], align 4 2112 // CHECK17-NEXT: call void @__kmpc_push_num_teams(ptr @[[GLOB1]], i32 [[TMP0]], i32 [[TMP2]], i32 123) 2113 // CHECK17-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB1]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l161.omp_outlined, ptr [[COMP_ADDR]]) 2114 // CHECK17-NEXT: ret void 2115 // 2116 // 2117 // CHECK17-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l161.omp_outlined 2118 // CHECK17-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[COMP:%.*]]) #[[ATTR1]] { 2119 // CHECK17-NEXT: entry: 2120 // CHECK17-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8 2121 // CHECK17-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8 2122 // CHECK17-NEXT: [[COMP_ADDR:%.*]] = alloca ptr, align 8 2123 // CHECK17-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8 2124 // CHECK17-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8 2125 // CHECK17-NEXT: store ptr [[COMP]], ptr [[COMP_ADDR]], align 8 2126 // CHECK17-NEXT: [[TMP0:%.*]] = load ptr, ptr [[COMP_ADDR]], align 8 2127 // CHECK17-NEXT: [[TMP1:%.*]] = load i32, ptr [[TMP0]], align 4 2128 // CHECK17-NEXT: [[INC:%.*]] = add nsw i32 [[TMP1]], 1 2129 // CHECK17-NEXT: store i32 [[INC]], ptr [[TMP0]], align 4 2130 // CHECK17-NEXT: ret void 2131 // 2132 // 2133 // CHECK17-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l169 2134 // CHECK17-SAME: (ptr noundef [[THIS:%.*]], i64 noundef [[COMP:%.*]]) #[[ATTR1]] { 2135 // CHECK17-NEXT: entry: 2136 // CHECK17-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 8 2137 // CHECK17-NEXT: [[COMP_ADDR:%.*]] = alloca i64, align 8 2138 // CHECK17-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]]) 2139 // CHECK17-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 8 2140 // CHECK17-NEXT: store i64 [[COMP]], ptr [[COMP_ADDR]], align 8 2141 // CHECK17-NEXT: [[TMP1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8 2142 // CHECK17-NEXT: [[B:%.*]] = getelementptr inbounds nuw [[STRUCT_SS:%.*]], ptr [[TMP1]], i32 0, i32 1 2143 // CHECK17-NEXT: [[TMP2:%.*]] = load float, ptr [[B]], align 4 2144 // CHECK17-NEXT: [[CONV:%.*]] = fptosi float [[TMP2]] to i32 2145 // CHECK17-NEXT: [[ADD:%.*]] = add nsw i32 [[CONV]], 123 2146 // CHECK17-NEXT: call void @__kmpc_push_num_teams(ptr @[[GLOB1]], i32 [[TMP0]], i32 456, i32 [[ADD]]) 2147 // CHECK17-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB1]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l169.omp_outlined, ptr [[COMP_ADDR]]) 2148 // CHECK17-NEXT: ret void 2149 // 2150 // 2151 // CHECK17-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l169.omp_outlined 2152 // CHECK17-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[COMP:%.*]]) #[[ATTR1]] { 2153 // CHECK17-NEXT: entry: 2154 // CHECK17-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8 2155 // CHECK17-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8 2156 // CHECK17-NEXT: [[COMP_ADDR:%.*]] = alloca ptr, align 8 2157 // CHECK17-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8 2158 // CHECK17-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8 2159 // CHECK17-NEXT: store ptr [[COMP]], ptr [[COMP_ADDR]], align 8 2160 // CHECK17-NEXT: [[TMP0:%.*]] = load ptr, ptr [[COMP_ADDR]], align 8 2161 // CHECK17-NEXT: [[TMP1:%.*]] = load i32, ptr [[TMP0]], align 4 2162 // CHECK17-NEXT: [[INC:%.*]] = add nsw i32 [[TMP1]], 1 2163 // CHECK17-NEXT: store i32 [[INC]], ptr [[TMP0]], align 4 2164 // CHECK17-NEXT: ret void 2165 // 2166 // 2167 // CHECK19-LABEL: define {{[^@]+}}@_Z21teams_template_structv 2168 // CHECK19-SAME: () #[[ATTR0:[0-9]+]] { 2169 // CHECK19-NEXT: entry: 2170 // CHECK19-NEXT: [[V:%.*]] = alloca [[STRUCT_SS:%.*]], align 4 2171 // CHECK19-NEXT: [[CALL:%.*]] = call noundef i32 @_ZN2SSIiLi123ELx456EE3fooEv(ptr noundef nonnull align 4 dereferenceable(8) [[V]]) 2172 // CHECK19-NEXT: ret i32 [[CALL]] 2173 // 2174 // 2175 // CHECK19-LABEL: define {{[^@]+}}@_ZN2SSIiLi123ELx456EE3fooEv 2176 // CHECK19-SAME: (ptr noundef nonnull align 4 dereferenceable(8) [[THIS:%.*]]) #[[ATTR0]] comdat align 2 { 2177 // CHECK19-NEXT: entry: 2178 // CHECK19-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 4 2179 // CHECK19-NEXT: [[COMP:%.*]] = alloca i32, align 4 2180 // CHECK19-NEXT: [[COMP_CASTED:%.*]] = alloca i32, align 4 2181 // CHECK19-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [2 x ptr], align 4 2182 // CHECK19-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [2 x ptr], align 4 2183 // CHECK19-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [2 x ptr], align 4 2184 // CHECK19-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8 2185 // CHECK19-NEXT: [[COMP_CASTED3:%.*]] = alloca i32, align 4 2186 // CHECK19-NEXT: [[DOTOFFLOAD_BASEPTRS4:%.*]] = alloca [2 x ptr], align 4 2187 // CHECK19-NEXT: [[DOTOFFLOAD_PTRS5:%.*]] = alloca [2 x ptr], align 4 2188 // CHECK19-NEXT: [[DOTOFFLOAD_MAPPERS6:%.*]] = alloca [2 x ptr], align 4 2189 // CHECK19-NEXT: [[KERNEL_ARGS8:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 8 2190 // CHECK19-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 4 2191 // CHECK19-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 4 2192 // CHECK19-NEXT: store i32 1, ptr [[COMP]], align 4 2193 // CHECK19-NEXT: [[TMP0:%.*]] = load i32, ptr [[COMP]], align 4 2194 // CHECK19-NEXT: store i32 [[TMP0]], ptr [[COMP_CASTED]], align 4 2195 // CHECK19-NEXT: [[TMP1:%.*]] = load i32, ptr [[COMP_CASTED]], align 4 2196 // CHECK19-NEXT: [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_SS:%.*]], ptr [[THIS1]], i32 0, i32 0 2197 // CHECK19-NEXT: [[TMP2:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 2198 // CHECK19-NEXT: store ptr [[THIS1]], ptr [[TMP2]], align 4 2199 // CHECK19-NEXT: [[TMP3:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0 2200 // CHECK19-NEXT: store ptr [[A]], ptr [[TMP3]], align 4 2201 // CHECK19-NEXT: [[TMP4:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 0 2202 // CHECK19-NEXT: store ptr null, ptr [[TMP4]], align 4 2203 // CHECK19-NEXT: [[TMP5:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 1 2204 // CHECK19-NEXT: store i32 [[TMP1]], ptr [[TMP5]], align 4 2205 // CHECK19-NEXT: [[TMP6:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 1 2206 // CHECK19-NEXT: store i32 [[TMP1]], ptr [[TMP6]], align 4 2207 // CHECK19-NEXT: [[TMP7:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 1 2208 // CHECK19-NEXT: store ptr null, ptr [[TMP7]], align 4 2209 // CHECK19-NEXT: [[TMP8:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 2210 // CHECK19-NEXT: [[TMP9:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0 2211 // CHECK19-NEXT: [[A2:%.*]] = getelementptr inbounds nuw [[STRUCT_SS]], ptr [[THIS1]], i32 0, i32 0 2212 // CHECK19-NEXT: [[TMP10:%.*]] = load i32, ptr [[A2]], align 4 2213 // CHECK19-NEXT: [[TMP11:%.*]] = insertvalue [3 x i32] zeroinitializer, i32 [[TMP10]], 0 2214 // CHECK19-NEXT: [[TMP12:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0 2215 // CHECK19-NEXT: store i32 3, ptr [[TMP12]], align 4 2216 // CHECK19-NEXT: [[TMP13:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1 2217 // CHECK19-NEXT: store i32 2, ptr [[TMP13]], align 4 2218 // CHECK19-NEXT: [[TMP14:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2 2219 // CHECK19-NEXT: store ptr [[TMP8]], ptr [[TMP14]], align 4 2220 // CHECK19-NEXT: [[TMP15:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3 2221 // CHECK19-NEXT: store ptr [[TMP9]], ptr [[TMP15]], align 4 2222 // CHECK19-NEXT: [[TMP16:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4 2223 // CHECK19-NEXT: store ptr @.offload_sizes, ptr [[TMP16]], align 4 2224 // CHECK19-NEXT: [[TMP17:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5 2225 // CHECK19-NEXT: store ptr @.offload_maptypes, ptr [[TMP17]], align 4 2226 // CHECK19-NEXT: [[TMP18:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6 2227 // CHECK19-NEXT: store ptr null, ptr [[TMP18]], align 4 2228 // CHECK19-NEXT: [[TMP19:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7 2229 // CHECK19-NEXT: store ptr null, ptr [[TMP19]], align 4 2230 // CHECK19-NEXT: [[TMP20:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8 2231 // CHECK19-NEXT: store i64 0, ptr [[TMP20]], align 8 2232 // CHECK19-NEXT: [[TMP21:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9 2233 // CHECK19-NEXT: store i64 0, ptr [[TMP21]], align 8 2234 // CHECK19-NEXT: [[TMP22:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10 2235 // CHECK19-NEXT: store [3 x i32] [[TMP11]], ptr [[TMP22]], align 4 2236 // CHECK19-NEXT: [[TMP23:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11 2237 // CHECK19-NEXT: store [3 x i32] [i32 123, i32 0, i32 0], ptr [[TMP23]], align 4 2238 // CHECK19-NEXT: [[TMP24:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12 2239 // CHECK19-NEXT: store i32 0, ptr [[TMP24]], align 4 2240 // CHECK19-NEXT: [[TMP25:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1:[0-9]+]], i64 -1, i32 [[TMP10]], i32 123, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l161.region_id, ptr [[KERNEL_ARGS]]) 2241 // CHECK19-NEXT: [[TMP26:%.*]] = icmp ne i32 [[TMP25]], 0 2242 // CHECK19-NEXT: br i1 [[TMP26]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]] 2243 // CHECK19: omp_offload.failed: 2244 // CHECK19-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l161(ptr [[THIS1]], i32 [[TMP1]]) #[[ATTR2:[0-9]+]] 2245 // CHECK19-NEXT: br label [[OMP_OFFLOAD_CONT]] 2246 // CHECK19: omp_offload.cont: 2247 // CHECK19-NEXT: [[TMP27:%.*]] = load i32, ptr [[COMP]], align 4 2248 // CHECK19-NEXT: store i32 [[TMP27]], ptr [[COMP_CASTED3]], align 4 2249 // CHECK19-NEXT: [[TMP28:%.*]] = load i32, ptr [[COMP_CASTED3]], align 4 2250 // CHECK19-NEXT: [[B:%.*]] = getelementptr inbounds nuw [[STRUCT_SS]], ptr [[THIS1]], i32 0, i32 1 2251 // CHECK19-NEXT: [[TMP29:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS4]], i32 0, i32 0 2252 // CHECK19-NEXT: store ptr [[THIS1]], ptr [[TMP29]], align 4 2253 // CHECK19-NEXT: [[TMP30:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS5]], i32 0, i32 0 2254 // CHECK19-NEXT: store ptr [[B]], ptr [[TMP30]], align 4 2255 // CHECK19-NEXT: [[TMP31:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS6]], i32 0, i32 0 2256 // CHECK19-NEXT: store ptr null, ptr [[TMP31]], align 4 2257 // CHECK19-NEXT: [[TMP32:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS4]], i32 0, i32 1 2258 // CHECK19-NEXT: store i32 [[TMP28]], ptr [[TMP32]], align 4 2259 // CHECK19-NEXT: [[TMP33:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS5]], i32 0, i32 1 2260 // CHECK19-NEXT: store i32 [[TMP28]], ptr [[TMP33]], align 4 2261 // CHECK19-NEXT: [[TMP34:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS6]], i32 0, i32 1 2262 // CHECK19-NEXT: store ptr null, ptr [[TMP34]], align 4 2263 // CHECK19-NEXT: [[TMP35:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS4]], i32 0, i32 0 2264 // CHECK19-NEXT: [[TMP36:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS5]], i32 0, i32 0 2265 // CHECK19-NEXT: [[B7:%.*]] = getelementptr inbounds nuw [[STRUCT_SS]], ptr [[THIS1]], i32 0, i32 1 2266 // CHECK19-NEXT: [[TMP37:%.*]] = load float, ptr [[B7]], align 4 2267 // CHECK19-NEXT: [[CONV:%.*]] = fptosi float [[TMP37]] to i32 2268 // CHECK19-NEXT: [[ADD:%.*]] = add nsw i32 [[CONV]], 123 2269 // CHECK19-NEXT: [[TMP38:%.*]] = insertvalue [3 x i32] zeroinitializer, i32 [[ADD]], 0 2270 // CHECK19-NEXT: [[TMP39:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS8]], i32 0, i32 0 2271 // CHECK19-NEXT: store i32 3, ptr [[TMP39]], align 4 2272 // CHECK19-NEXT: [[TMP40:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS8]], i32 0, i32 1 2273 // CHECK19-NEXT: store i32 2, ptr [[TMP40]], align 4 2274 // CHECK19-NEXT: [[TMP41:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS8]], i32 0, i32 2 2275 // CHECK19-NEXT: store ptr [[TMP35]], ptr [[TMP41]], align 4 2276 // CHECK19-NEXT: [[TMP42:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS8]], i32 0, i32 3 2277 // CHECK19-NEXT: store ptr [[TMP36]], ptr [[TMP42]], align 4 2278 // CHECK19-NEXT: [[TMP43:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS8]], i32 0, i32 4 2279 // CHECK19-NEXT: store ptr @.offload_sizes.1, ptr [[TMP43]], align 4 2280 // CHECK19-NEXT: [[TMP44:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS8]], i32 0, i32 5 2281 // CHECK19-NEXT: store ptr @.offload_maptypes.2, ptr [[TMP44]], align 4 2282 // CHECK19-NEXT: [[TMP45:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS8]], i32 0, i32 6 2283 // CHECK19-NEXT: store ptr null, ptr [[TMP45]], align 4 2284 // CHECK19-NEXT: [[TMP46:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS8]], i32 0, i32 7 2285 // CHECK19-NEXT: store ptr null, ptr [[TMP46]], align 4 2286 // CHECK19-NEXT: [[TMP47:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS8]], i32 0, i32 8 2287 // CHECK19-NEXT: store i64 0, ptr [[TMP47]], align 8 2288 // CHECK19-NEXT: [[TMP48:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS8]], i32 0, i32 9 2289 // CHECK19-NEXT: store i64 0, ptr [[TMP48]], align 8 2290 // CHECK19-NEXT: [[TMP49:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS8]], i32 0, i32 10 2291 // CHECK19-NEXT: store [3 x i32] [i32 456, i32 0, i32 0], ptr [[TMP49]], align 4 2292 // CHECK19-NEXT: [[TMP50:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS8]], i32 0, i32 11 2293 // CHECK19-NEXT: store [3 x i32] [[TMP38]], ptr [[TMP50]], align 4 2294 // CHECK19-NEXT: [[TMP51:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS8]], i32 0, i32 12 2295 // CHECK19-NEXT: store i32 0, ptr [[TMP51]], align 4 2296 // CHECK19-NEXT: [[TMP52:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 -1, i32 456, i32 [[ADD]], ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l169.region_id, ptr [[KERNEL_ARGS8]]) 2297 // CHECK19-NEXT: [[TMP53:%.*]] = icmp ne i32 [[TMP52]], 0 2298 // CHECK19-NEXT: br i1 [[TMP53]], label [[OMP_OFFLOAD_FAILED9:%.*]], label [[OMP_OFFLOAD_CONT10:%.*]] 2299 // CHECK19: omp_offload.failed9: 2300 // CHECK19-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l169(ptr [[THIS1]], i32 [[TMP28]]) #[[ATTR2]] 2301 // CHECK19-NEXT: br label [[OMP_OFFLOAD_CONT10]] 2302 // CHECK19: omp_offload.cont10: 2303 // CHECK19-NEXT: [[TMP54:%.*]] = load i32, ptr [[COMP]], align 4 2304 // CHECK19-NEXT: ret i32 [[TMP54]] 2305 // 2306 // 2307 // CHECK19-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l161 2308 // CHECK19-SAME: (ptr noundef [[THIS:%.*]], i32 noundef [[COMP:%.*]]) #[[ATTR1:[0-9]+]] { 2309 // CHECK19-NEXT: entry: 2310 // CHECK19-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 4 2311 // CHECK19-NEXT: [[COMP_ADDR:%.*]] = alloca i32, align 4 2312 // CHECK19-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]]) 2313 // CHECK19-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 4 2314 // CHECK19-NEXT: store i32 [[COMP]], ptr [[COMP_ADDR]], align 4 2315 // CHECK19-NEXT: [[TMP1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 4 2316 // CHECK19-NEXT: [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_SS:%.*]], ptr [[TMP1]], i32 0, i32 0 2317 // CHECK19-NEXT: [[TMP2:%.*]] = load i32, ptr [[A]], align 4 2318 // CHECK19-NEXT: call void @__kmpc_push_num_teams(ptr @[[GLOB1]], i32 [[TMP0]], i32 [[TMP2]], i32 123) 2319 // CHECK19-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB1]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l161.omp_outlined, ptr [[COMP_ADDR]]) 2320 // CHECK19-NEXT: ret void 2321 // 2322 // 2323 // CHECK19-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l161.omp_outlined 2324 // CHECK19-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[COMP:%.*]]) #[[ATTR1]] { 2325 // CHECK19-NEXT: entry: 2326 // CHECK19-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 4 2327 // CHECK19-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 4 2328 // CHECK19-NEXT: [[COMP_ADDR:%.*]] = alloca ptr, align 4 2329 // CHECK19-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 4 2330 // CHECK19-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 4 2331 // CHECK19-NEXT: store ptr [[COMP]], ptr [[COMP_ADDR]], align 4 2332 // CHECK19-NEXT: [[TMP0:%.*]] = load ptr, ptr [[COMP_ADDR]], align 4 2333 // CHECK19-NEXT: [[TMP1:%.*]] = load i32, ptr [[TMP0]], align 4 2334 // CHECK19-NEXT: [[INC:%.*]] = add nsw i32 [[TMP1]], 1 2335 // CHECK19-NEXT: store i32 [[INC]], ptr [[TMP0]], align 4 2336 // CHECK19-NEXT: ret void 2337 // 2338 // 2339 // CHECK19-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l169 2340 // CHECK19-SAME: (ptr noundef [[THIS:%.*]], i32 noundef [[COMP:%.*]]) #[[ATTR1]] { 2341 // CHECK19-NEXT: entry: 2342 // CHECK19-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 4 2343 // CHECK19-NEXT: [[COMP_ADDR:%.*]] = alloca i32, align 4 2344 // CHECK19-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]]) 2345 // CHECK19-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 4 2346 // CHECK19-NEXT: store i32 [[COMP]], ptr [[COMP_ADDR]], align 4 2347 // CHECK19-NEXT: [[TMP1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 4 2348 // CHECK19-NEXT: [[B:%.*]] = getelementptr inbounds nuw [[STRUCT_SS:%.*]], ptr [[TMP1]], i32 0, i32 1 2349 // CHECK19-NEXT: [[TMP2:%.*]] = load float, ptr [[B]], align 4 2350 // CHECK19-NEXT: [[CONV:%.*]] = fptosi float [[TMP2]] to i32 2351 // CHECK19-NEXT: [[ADD:%.*]] = add nsw i32 [[CONV]], 123 2352 // CHECK19-NEXT: call void @__kmpc_push_num_teams(ptr @[[GLOB1]], i32 [[TMP0]], i32 456, i32 [[ADD]]) 2353 // CHECK19-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB1]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l169.omp_outlined, ptr [[COMP_ADDR]]) 2354 // CHECK19-NEXT: ret void 2355 // 2356 // 2357 // CHECK19-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2SSIiLi123ELx456EE3fooEv_l169.omp_outlined 2358 // CHECK19-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[COMP:%.*]]) #[[ATTR1]] { 2359 // CHECK19-NEXT: entry: 2360 // CHECK19-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 4 2361 // CHECK19-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 4 2362 // CHECK19-NEXT: [[COMP_ADDR:%.*]] = alloca ptr, align 4 2363 // CHECK19-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 4 2364 // CHECK19-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 4 2365 // CHECK19-NEXT: store ptr [[COMP]], ptr [[COMP_ADDR]], align 4 2366 // CHECK19-NEXT: [[TMP0:%.*]] = load ptr, ptr [[COMP_ADDR]], align 4 2367 // CHECK19-NEXT: [[TMP1:%.*]] = load i32, ptr [[TMP0]], align 4 2368 // CHECK19-NEXT: [[INC:%.*]] = add nsw i32 [[TMP1]], 1 2369 // CHECK19-NEXT: store i32 [[INC]], ptr [[TMP0]], align 4 2370 // CHECK19-NEXT: ret void 2371 // 2372 // 2373 // CHECK25-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l216 2374 // CHECK25-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], i64 noundef [[ARGC:%.*]]) #[[ATTR0:[0-9]+]] { 2375 // CHECK25-NEXT: entry: 2376 // CHECK25-NEXT: [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 8 2377 // CHECK25-NEXT: [[ARGC_ADDR:%.*]] = alloca i64, align 8 2378 // CHECK25-NEXT: store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 8 2379 // CHECK25-NEXT: store i64 [[ARGC]], ptr [[ARGC_ADDR]], align 8 2380 // CHECK25-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB1:[0-9]+]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l216.omp_outlined, ptr [[ARGC_ADDR]]) 2381 // CHECK25-NEXT: ret void 2382 // 2383 // 2384 // CHECK25-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l216.omp_outlined 2385 // CHECK25-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[ARGC:%.*]]) #[[ATTR0]] { 2386 // CHECK25-NEXT: entry: 2387 // CHECK25-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8 2388 // CHECK25-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8 2389 // CHECK25-NEXT: [[ARGC_ADDR:%.*]] = alloca ptr, align 8 2390 // CHECK25-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8 2391 // CHECK25-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8 2392 // CHECK25-NEXT: store ptr [[ARGC]], ptr [[ARGC_ADDR]], align 8 2393 // CHECK25-NEXT: [[TMP0:%.*]] = load ptr, ptr [[ARGC_ADDR]], align 8 2394 // CHECK25-NEXT: store i32 0, ptr [[TMP0]], align 4 2395 // CHECK25-NEXT: ret void 2396 // 2397 // 2398 // CHECK25-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIPPcEiT__l209 2399 // CHECK25-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], ptr noundef [[ARGC:%.*]]) #[[ATTR0]] { 2400 // CHECK25-NEXT: entry: 2401 // CHECK25-NEXT: [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 8 2402 // CHECK25-NEXT: [[ARGC_ADDR:%.*]] = alloca ptr, align 8 2403 // CHECK25-NEXT: store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 8 2404 // CHECK25-NEXT: store ptr [[ARGC]], ptr [[ARGC_ADDR]], align 8 2405 // CHECK25-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB1]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIPPcEiT__l209.omp_outlined, ptr [[ARGC_ADDR]]) 2406 // CHECK25-NEXT: ret void 2407 // 2408 // 2409 // CHECK25-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIPPcEiT__l209.omp_outlined 2410 // CHECK25-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 8 dereferenceable(8) [[ARGC:%.*]]) #[[ATTR0]] { 2411 // CHECK25-NEXT: entry: 2412 // CHECK25-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8 2413 // CHECK25-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8 2414 // CHECK25-NEXT: [[ARGC_ADDR:%.*]] = alloca ptr, align 8 2415 // CHECK25-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8 2416 // CHECK25-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8 2417 // CHECK25-NEXT: store ptr [[ARGC]], ptr [[ARGC_ADDR]], align 8 2418 // CHECK25-NEXT: [[TMP0:%.*]] = load ptr, ptr [[ARGC_ADDR]], align 8 2419 // CHECK25-NEXT: store ptr null, ptr [[TMP0]], align 8 2420 // CHECK25-NEXT: ret void 2421 // 2422 // 2423 // CHECK27-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l216 2424 // CHECK27-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], i32 noundef [[ARGC:%.*]]) #[[ATTR0:[0-9]+]] { 2425 // CHECK27-NEXT: entry: 2426 // CHECK27-NEXT: [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 4 2427 // CHECK27-NEXT: [[ARGC_ADDR:%.*]] = alloca i32, align 4 2428 // CHECK27-NEXT: store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 4 2429 // CHECK27-NEXT: store i32 [[ARGC]], ptr [[ARGC_ADDR]], align 4 2430 // CHECK27-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB1:[0-9]+]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l216.omp_outlined, ptr [[ARGC_ADDR]]) 2431 // CHECK27-NEXT: ret void 2432 // 2433 // 2434 // CHECK27-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l216.omp_outlined 2435 // CHECK27-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[ARGC:%.*]]) #[[ATTR0]] { 2436 // CHECK27-NEXT: entry: 2437 // CHECK27-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 4 2438 // CHECK27-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 4 2439 // CHECK27-NEXT: [[ARGC_ADDR:%.*]] = alloca ptr, align 4 2440 // CHECK27-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 4 2441 // CHECK27-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 4 2442 // CHECK27-NEXT: store ptr [[ARGC]], ptr [[ARGC_ADDR]], align 4 2443 // CHECK27-NEXT: [[TMP0:%.*]] = load ptr, ptr [[ARGC_ADDR]], align 4 2444 // CHECK27-NEXT: store i32 0, ptr [[TMP0]], align 4 2445 // CHECK27-NEXT: ret void 2446 // 2447 // 2448 // CHECK27-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIPPcEiT__l209 2449 // CHECK27-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], ptr noundef [[ARGC:%.*]]) #[[ATTR0]] { 2450 // CHECK27-NEXT: entry: 2451 // CHECK27-NEXT: [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 4 2452 // CHECK27-NEXT: [[ARGC_ADDR:%.*]] = alloca ptr, align 4 2453 // CHECK27-NEXT: store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 4 2454 // CHECK27-NEXT: store ptr [[ARGC]], ptr [[ARGC_ADDR]], align 4 2455 // CHECK27-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB1]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIPPcEiT__l209.omp_outlined, ptr [[ARGC_ADDR]]) 2456 // CHECK27-NEXT: ret void 2457 // 2458 // 2459 // CHECK27-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIPPcEiT__l209.omp_outlined 2460 // CHECK27-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[ARGC:%.*]]) #[[ATTR0]] { 2461 // CHECK27-NEXT: entry: 2462 // CHECK27-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 4 2463 // CHECK27-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 4 2464 // CHECK27-NEXT: [[ARGC_ADDR:%.*]] = alloca ptr, align 4 2465 // CHECK27-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 4 2466 // CHECK27-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 4 2467 // CHECK27-NEXT: store ptr [[ARGC]], ptr [[ARGC_ADDR]], align 4 2468 // CHECK27-NEXT: [[TMP0:%.*]] = load ptr, ptr [[ARGC_ADDR]], align 4 2469 // CHECK27-NEXT: store ptr null, ptr [[TMP0]], align 4 2470 // CHECK27-NEXT: ret void 2471 // 2472 // 2473 // CHECK33-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l265 2474 // CHECK33-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], i64 noundef [[A:%.*]], i64 noundef [[B:%.*]], i64 noundef [[ARGC:%.*]]) #[[ATTR0:[0-9]+]] { 2475 // CHECK33-NEXT: entry: 2476 // CHECK33-NEXT: [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 8 2477 // CHECK33-NEXT: [[A_ADDR:%.*]] = alloca i64, align 8 2478 // CHECK33-NEXT: [[B_ADDR:%.*]] = alloca i64, align 8 2479 // CHECK33-NEXT: [[ARGC_ADDR:%.*]] = alloca i64, align 8 2480 // CHECK33-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1:[0-9]+]]) 2481 // CHECK33-NEXT: store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 8 2482 // CHECK33-NEXT: store i64 [[A]], ptr [[A_ADDR]], align 8 2483 // CHECK33-NEXT: store i64 [[B]], ptr [[B_ADDR]], align 8 2484 // CHECK33-NEXT: store i64 [[ARGC]], ptr [[ARGC_ADDR]], align 8 2485 // CHECK33-NEXT: [[TMP1:%.*]] = load i32, ptr [[A_ADDR]], align 4 2486 // CHECK33-NEXT: [[TMP2:%.*]] = load i32, ptr [[B_ADDR]], align 4 2487 // CHECK33-NEXT: call void @__kmpc_push_num_teams(ptr @[[GLOB1]], i32 [[TMP0]], i32 [[TMP1]], i32 [[TMP2]]) 2488 // CHECK33-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB1]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l265.omp_outlined, ptr [[ARGC_ADDR]]) 2489 // CHECK33-NEXT: ret void 2490 // 2491 // 2492 // CHECK33-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l265.omp_outlined 2493 // CHECK33-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[ARGC:%.*]]) #[[ATTR0]] { 2494 // CHECK33-NEXT: entry: 2495 // CHECK33-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8 2496 // CHECK33-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8 2497 // CHECK33-NEXT: [[ARGC_ADDR:%.*]] = alloca ptr, align 8 2498 // CHECK33-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8 2499 // CHECK33-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8 2500 // CHECK33-NEXT: store ptr [[ARGC]], ptr [[ARGC_ADDR]], align 8 2501 // CHECK33-NEXT: [[TMP0:%.*]] = load ptr, ptr [[ARGC_ADDR]], align 8 2502 // CHECK33-NEXT: store i32 0, ptr [[TMP0]], align 4 2503 // CHECK33-NEXT: ret void 2504 // 2505 // 2506 // CHECK33-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIPPcEiT__l254 2507 // CHECK33-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], i64 noundef [[A:%.*]], i64 noundef [[B:%.*]], ptr noundef [[ARGC:%.*]]) #[[ATTR0]] { 2508 // CHECK33-NEXT: entry: 2509 // CHECK33-NEXT: [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 8 2510 // CHECK33-NEXT: [[A_ADDR:%.*]] = alloca i64, align 8 2511 // CHECK33-NEXT: [[B_ADDR:%.*]] = alloca i64, align 8 2512 // CHECK33-NEXT: [[ARGC_ADDR:%.*]] = alloca ptr, align 8 2513 // CHECK33-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]]) 2514 // CHECK33-NEXT: store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 8 2515 // CHECK33-NEXT: store i64 [[A]], ptr [[A_ADDR]], align 8 2516 // CHECK33-NEXT: store i64 [[B]], ptr [[B_ADDR]], align 8 2517 // CHECK33-NEXT: store ptr [[ARGC]], ptr [[ARGC_ADDR]], align 8 2518 // CHECK33-NEXT: [[TMP1:%.*]] = load i32, ptr [[A_ADDR]], align 4 2519 // CHECK33-NEXT: [[TMP2:%.*]] = load i32, ptr [[B_ADDR]], align 4 2520 // CHECK33-NEXT: call void @__kmpc_push_num_teams(ptr @[[GLOB1]], i32 [[TMP0]], i32 [[TMP1]], i32 [[TMP2]]) 2521 // CHECK33-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB1]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIPPcEiT__l254.omp_outlined, ptr [[ARGC_ADDR]]) 2522 // CHECK33-NEXT: ret void 2523 // 2524 // 2525 // CHECK33-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIPPcEiT__l254.omp_outlined 2526 // CHECK33-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 8 dereferenceable(8) [[ARGC:%.*]]) #[[ATTR0]] { 2527 // CHECK33-NEXT: entry: 2528 // CHECK33-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8 2529 // CHECK33-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8 2530 // CHECK33-NEXT: [[ARGC_ADDR:%.*]] = alloca ptr, align 8 2531 // CHECK33-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8 2532 // CHECK33-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8 2533 // CHECK33-NEXT: store ptr [[ARGC]], ptr [[ARGC_ADDR]], align 8 2534 // CHECK33-NEXT: [[TMP0:%.*]] = load ptr, ptr [[ARGC_ADDR]], align 8 2535 // CHECK33-NEXT: store ptr null, ptr [[TMP0]], align 8 2536 // CHECK33-NEXT: ret void 2537 // 2538 // 2539 // CHECK35-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l265 2540 // CHECK35-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], i32 noundef [[A:%.*]], i32 noundef [[B:%.*]], i32 noundef [[ARGC:%.*]]) #[[ATTR0:[0-9]+]] { 2541 // CHECK35-NEXT: entry: 2542 // CHECK35-NEXT: [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 4 2543 // CHECK35-NEXT: [[A_ADDR:%.*]] = alloca i32, align 4 2544 // CHECK35-NEXT: [[B_ADDR:%.*]] = alloca i32, align 4 2545 // CHECK35-NEXT: [[ARGC_ADDR:%.*]] = alloca i32, align 4 2546 // CHECK35-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1:[0-9]+]]) 2547 // CHECK35-NEXT: store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 4 2548 // CHECK35-NEXT: store i32 [[A]], ptr [[A_ADDR]], align 4 2549 // CHECK35-NEXT: store i32 [[B]], ptr [[B_ADDR]], align 4 2550 // CHECK35-NEXT: store i32 [[ARGC]], ptr [[ARGC_ADDR]], align 4 2551 // CHECK35-NEXT: [[TMP1:%.*]] = load i32, ptr [[A_ADDR]], align 4 2552 // CHECK35-NEXT: [[TMP2:%.*]] = load i32, ptr [[B_ADDR]], align 4 2553 // CHECK35-NEXT: call void @__kmpc_push_num_teams(ptr @[[GLOB1]], i32 [[TMP0]], i32 [[TMP1]], i32 [[TMP2]]) 2554 // CHECK35-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB1]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l265.omp_outlined, ptr [[ARGC_ADDR]]) 2555 // CHECK35-NEXT: ret void 2556 // 2557 // 2558 // CHECK35-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l265.omp_outlined 2559 // CHECK35-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[ARGC:%.*]]) #[[ATTR0]] { 2560 // CHECK35-NEXT: entry: 2561 // CHECK35-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 4 2562 // CHECK35-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 4 2563 // CHECK35-NEXT: [[ARGC_ADDR:%.*]] = alloca ptr, align 4 2564 // CHECK35-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 4 2565 // CHECK35-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 4 2566 // CHECK35-NEXT: store ptr [[ARGC]], ptr [[ARGC_ADDR]], align 4 2567 // CHECK35-NEXT: [[TMP0:%.*]] = load ptr, ptr [[ARGC_ADDR]], align 4 2568 // CHECK35-NEXT: store i32 0, ptr [[TMP0]], align 4 2569 // CHECK35-NEXT: ret void 2570 // 2571 // 2572 // CHECK35-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIPPcEiT__l254 2573 // CHECK35-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], i32 noundef [[A:%.*]], i32 noundef [[B:%.*]], ptr noundef [[ARGC:%.*]]) #[[ATTR0]] { 2574 // CHECK35-NEXT: entry: 2575 // CHECK35-NEXT: [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 4 2576 // CHECK35-NEXT: [[A_ADDR:%.*]] = alloca i32, align 4 2577 // CHECK35-NEXT: [[B_ADDR:%.*]] = alloca i32, align 4 2578 // CHECK35-NEXT: [[ARGC_ADDR:%.*]] = alloca ptr, align 4 2579 // CHECK35-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]]) 2580 // CHECK35-NEXT: store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 4 2581 // CHECK35-NEXT: store i32 [[A]], ptr [[A_ADDR]], align 4 2582 // CHECK35-NEXT: store i32 [[B]], ptr [[B_ADDR]], align 4 2583 // CHECK35-NEXT: store ptr [[ARGC]], ptr [[ARGC_ADDR]], align 4 2584 // CHECK35-NEXT: [[TMP1:%.*]] = load i32, ptr [[A_ADDR]], align 4 2585 // CHECK35-NEXT: [[TMP2:%.*]] = load i32, ptr [[B_ADDR]], align 4 2586 // CHECK35-NEXT: call void @__kmpc_push_num_teams(ptr @[[GLOB1]], i32 [[TMP0]], i32 [[TMP1]], i32 [[TMP2]]) 2587 // CHECK35-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB1]], i32 1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIPPcEiT__l254.omp_outlined, ptr [[ARGC_ADDR]]) 2588 // CHECK35-NEXT: ret void 2589 // 2590 // 2591 // CHECK35-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIPPcEiT__l254.omp_outlined 2592 // CHECK35-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[ARGC:%.*]]) #[[ATTR0]] { 2593 // CHECK35-NEXT: entry: 2594 // CHECK35-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 4 2595 // CHECK35-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 4 2596 // CHECK35-NEXT: [[ARGC_ADDR:%.*]] = alloca ptr, align 4 2597 // CHECK35-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 4 2598 // CHECK35-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 4 2599 // CHECK35-NEXT: store ptr [[ARGC]], ptr [[ARGC_ADDR]], align 4 2600 // CHECK35-NEXT: [[TMP0:%.*]] = load ptr, ptr [[ARGC_ADDR]], align 4 2601 // CHECK35-NEXT: store ptr null, ptr [[TMP0]], align 4 2602 // CHECK35-NEXT: ret void 2603 // 2604 // 2605 // CHECK41-LABEL: define {{[^@]+}}@_Z3foov 2606 // CHECK41-SAME: () #[[ATTR0:[0-9]+]] { 2607 // CHECK41-NEXT: entry: 2608 // CHECK41-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB1:[0-9]+]], i32 0, ptr @_Z3foov.omp_outlined) 2609 // CHECK41-NEXT: ret void 2610 // 2611 // 2612 // CHECK41-LABEL: define {{[^@]+}}@_Z3foov.omp_outlined 2613 // CHECK41-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR1:[0-9]+]] { 2614 // CHECK41-NEXT: entry: 2615 // CHECK41-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8 2616 // CHECK41-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8 2617 // CHECK41-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8 2618 // CHECK41-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8 2619 // CHECK41-NEXT: ret void 2620 // 2621 // 2622 // CHECK43-LABEL: define {{[^@]+}}@_Z3foov 2623 // CHECK43-SAME: () #[[ATTR0:[0-9]+]] { 2624 // CHECK43-NEXT: entry: 2625 // CHECK43-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @[[GLOB1:[0-9]+]], i32 0, ptr @_Z3foov.omp_outlined) 2626 // CHECK43-NEXT: ret void 2627 // 2628 // 2629 // CHECK43-LABEL: define {{[^@]+}}@_Z3foov.omp_outlined 2630 // CHECK43-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR1:[0-9]+]] { 2631 // CHECK43-NEXT: entry: 2632 // CHECK43-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 4 2633 // CHECK43-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 4 2634 // CHECK43-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 4 2635 // CHECK43-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 4 2636 // CHECK43-NEXT: ret void 2637 // 2638