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]+" --prefix-filecheck-ir-name _ 2 // RUN: %clang_cc1 -DCK1 -verify -fopenmp -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s 3 4 // RUN: %clang_cc1 -DCK1 -verify -fopenmp-simd -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s 5 // expected-no-diagnostics 6 7 struct ST { 8 int *a; 9 }; 10 typedef int arr[10]; 11 typedef ST STarr[10]; 12 struct SA { 13 const int da[5] = { 0 }; 14 ST g[10]; 15 STarr &rg = g; 16 int i; 17 int &j = i; 18 int *k = &j; 19 int *&z = k; 20 int aa[10]; 21 arr &raa = aa; 22 void func(int arg) { 23 #pragma omp target has_device_addr(k) 24 {k++;} 25 #pragma omp target has_device_addr(z) 26 {z++;} 27 #pragma omp target has_device_addr(aa) 28 {aa[0]=1;} 29 #pragma omp target has_device_addr(raa) 30 {raa[0] = 10;} 31 #pragma omp target has_device_addr(g) 32 {g[0].a= &i;} 33 #pragma omp target has_device_addr(da) 34 {int a = da[1];} 35 return; 36 } 37 }; 38 39 struct SB { 40 unsigned A; 41 unsigned B; 42 float Arr[100]; 43 float *Ptr; 44 float *foo() { 45 return &Arr[0]; 46 } 47 }; 48 49 struct SC { 50 unsigned A : 2; 51 unsigned B : 3; 52 unsigned C; 53 unsigned D; 54 float Arr[100]; 55 SB S; 56 SB ArrS[100]; 57 SB *PtrS; 58 SB *&RPtrS; 59 float *Ptr; 60 61 SC(SB *&_RPtrS) : RPtrS(_RPtrS) {} 62 }; 63 64 union SD { 65 unsigned A; 66 float B; 67 }; 68 69 struct S1; 70 extern S1 a; 71 class S2 { 72 mutable int a; 73 public: 74 S2():a(0) { } 75 S2(S2 &s2):a(s2.a) { } 76 static float S2s; 77 static const float S2sc; 78 }; 79 const float S2::S2sc = 0; 80 const S2 b; 81 const S2 ba[5]; 82 class S3 { 83 int a; 84 public: 85 S3():a(0) { } 86 S3(S3 &s3):a(s3.a) { } 87 }; 88 const S3 c; 89 const S3 ca[5]; 90 extern const int f; 91 class S4 { 92 int a; 93 S4(); 94 S4(const S4 &s4); 95 public: 96 S4(int v):a(v) { } 97 }; 98 class S5 { 99 int a; 100 S5():a(0) {} 101 S5(const S5 &s5):a(s5.a) { } 102 public: 103 S5(int v):a(v) { } 104 }; 105 106 S3 h; 107 #pragma omp threadprivate(h) 108 109 typedef struct { 110 int a; 111 } S6; 112 113 template <typename T> 114 T tmain(T argc) { 115 const T da[5] = { 0 }; 116 S6 h[10]; 117 auto &rh = h; 118 T i; 119 T &j = i; 120 T *k = &j; 121 T *&z = k; 122 T aa[10]; 123 #pragma omp target has_device_addr(k) 124 {k++;} 125 #pragma omp target has_device_addr(z) 126 {z++;} 127 #pragma omp target has_device_addr(aa) 128 {T a = aa[0];} 129 #pragma omp target has_device_addr(h) 130 {int a = h[0].a;} 131 return 0; 132 } 133 134 135 int main(int argc, char **argv) { 136 const int da[5] = { 0 }; 137 S6 h[10]; 138 auto &rh = h; 139 int i; 140 int &j = i; 141 int *k = &j; 142 int *&z = k; 143 int aa[10]; 144 auto &raa = aa; 145 #pragma omp target has_device_addr(k) 146 {k++;} 147 #pragma omp target has_device_addr(z) 148 {z++;} 149 #pragma omp target has_device_addr(aa) 150 {aa[0]=1;} 151 #pragma omp target has_device_addr(raa) 152 {int a = raa[0];} 153 #pragma omp target has_device_addr(h) 154 {int a = h[1].a;} 155 #pragma omp target has_device_addr(da[1:3]) 156 {int a = da[1];} 157 return tmain<int>(argc) + *tmain<int *>(&argc); 158 } 159 160 struct SomeKernel { 161 int targetDev; 162 float devPtr; 163 SomeKernel(); 164 ~SomeKernel(); 165 166 template<unsigned int nRHS> 167 void apply() { 168 #pragma omp target has_device_addr(devPtr) device(targetDev) 169 { 170 devPtr++; 171 targetDev++; 172 } 173 } 174 }; 175 176 void use_template() { 177 SomeKernel aKern; 178 aKern.apply<32>(); 179 } 180 // CHECK-LABEL: define {{[^@]+}}@__cxx_global_var_init 181 // CHECK-SAME: () #[[ATTR0:[0-9]+]] { 182 // CHECK-NEXT: entry: 183 // CHECK-NEXT: call void @_ZN2S2C1Ev(ptr noundef nonnull align 4 dereferenceable(4) @_ZL1b) 184 // CHECK-NEXT: ret void 185 // 186 // 187 // CHECK-LABEL: define {{[^@]+}}@_ZN2S2C1Ev 188 // CHECK-SAME: (ptr noundef nonnull align 4 dereferenceable(4) [[THIS:%.*]]) unnamed_addr #[[ATTR1:[0-9]+]] comdat { 189 // CHECK-NEXT: entry: 190 // CHECK-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 8 191 // CHECK-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 8 192 // CHECK-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8 193 // CHECK-NEXT: call void @_ZN2S2C2Ev(ptr noundef nonnull align 4 dereferenceable(4) [[THIS1]]) 194 // CHECK-NEXT: ret void 195 // 196 // 197 // CHECK-LABEL: define {{[^@]+}}@_ZN2S2C2Ev 198 // CHECK-SAME: (ptr noundef nonnull align 4 dereferenceable(4) [[THIS:%.*]]) unnamed_addr #[[ATTR1]] comdat { 199 // CHECK-NEXT: entry: 200 // CHECK-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 8 201 // CHECK-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 8 202 // CHECK-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8 203 // CHECK-NEXT: [[A:%.*]] = getelementptr inbounds nuw [[CLASS_S2:%.*]], ptr [[THIS1]], i32 0, i32 0 204 // CHECK-NEXT: store i32 0, ptr [[A]], align 4 205 // CHECK-NEXT: ret void 206 // 207 // 208 // CHECK-LABEL: define {{[^@]+}}@__cxx_global_var_init.1 209 // CHECK-SAME: () #[[ATTR0]] { 210 // CHECK-NEXT: entry: 211 // CHECK-NEXT: br label [[ARRAYCTOR_LOOP:%.*]] 212 // CHECK: arrayctor.loop: 213 // CHECK-NEXT: [[ARRAYCTOR_CUR:%.*]] = phi ptr [ @_ZL2ba, [[ENTRY:%.*]] ], [ [[ARRAYCTOR_NEXT:%.*]], [[ARRAYCTOR_LOOP]] ] 214 // CHECK-NEXT: call void @_ZN2S2C1Ev(ptr noundef nonnull align 4 dereferenceable(4) [[ARRAYCTOR_CUR]]) 215 // CHECK-NEXT: [[ARRAYCTOR_NEXT]] = getelementptr inbounds [[CLASS_S2:%.*]], ptr [[ARRAYCTOR_CUR]], i64 1 216 // CHECK-NEXT: [[ARRAYCTOR_DONE:%.*]] = icmp eq ptr [[ARRAYCTOR_NEXT]], getelementptr inbounds ([[CLASS_S2]], ptr @_ZL2ba, i64 5) 217 // CHECK-NEXT: br i1 [[ARRAYCTOR_DONE]], label [[ARRAYCTOR_CONT:%.*]], label [[ARRAYCTOR_LOOP]] 218 // CHECK: arrayctor.cont: 219 // CHECK-NEXT: ret void 220 // 221 // 222 // CHECK-LABEL: define {{[^@]+}}@__cxx_global_var_init.2 223 // CHECK-SAME: () #[[ATTR0]] { 224 // CHECK-NEXT: entry: 225 // CHECK-NEXT: call void @_ZN2S3C1Ev(ptr noundef nonnull align 4 dereferenceable(4) @_ZL1c) 226 // CHECK-NEXT: ret void 227 // 228 // 229 // CHECK-LABEL: define {{[^@]+}}@_ZN2S3C1Ev 230 // CHECK-SAME: (ptr noundef nonnull align 4 dereferenceable(4) [[THIS:%.*]]) unnamed_addr #[[ATTR1]] comdat { 231 // CHECK-NEXT: entry: 232 // CHECK-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 8 233 // CHECK-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 8 234 // CHECK-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8 235 // CHECK-NEXT: call void @_ZN2S3C2Ev(ptr noundef nonnull align 4 dereferenceable(4) [[THIS1]]) 236 // CHECK-NEXT: ret void 237 // 238 // 239 // CHECK-LABEL: define {{[^@]+}}@_ZN2S3C2Ev 240 // CHECK-SAME: (ptr noundef nonnull align 4 dereferenceable(4) [[THIS:%.*]]) unnamed_addr #[[ATTR1]] comdat { 241 // CHECK-NEXT: entry: 242 // CHECK-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 8 243 // CHECK-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 8 244 // CHECK-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8 245 // CHECK-NEXT: [[A:%.*]] = getelementptr inbounds nuw [[CLASS_S3:%.*]], ptr [[THIS1]], i32 0, i32 0 246 // CHECK-NEXT: store i32 0, ptr [[A]], align 4 247 // CHECK-NEXT: ret void 248 // 249 // 250 // CHECK-LABEL: define {{[^@]+}}@__cxx_global_var_init.3 251 // CHECK-SAME: () #[[ATTR0]] { 252 // CHECK-NEXT: entry: 253 // CHECK-NEXT: br label [[ARRAYCTOR_LOOP:%.*]] 254 // CHECK: arrayctor.loop: 255 // CHECK-NEXT: [[ARRAYCTOR_CUR:%.*]] = phi ptr [ @_ZL2ca, [[ENTRY:%.*]] ], [ [[ARRAYCTOR_NEXT:%.*]], [[ARRAYCTOR_LOOP]] ] 256 // CHECK-NEXT: call void @_ZN2S3C1Ev(ptr noundef nonnull align 4 dereferenceable(4) [[ARRAYCTOR_CUR]]) 257 // CHECK-NEXT: [[ARRAYCTOR_NEXT]] = getelementptr inbounds [[CLASS_S3:%.*]], ptr [[ARRAYCTOR_CUR]], i64 1 258 // CHECK-NEXT: [[ARRAYCTOR_DONE:%.*]] = icmp eq ptr [[ARRAYCTOR_NEXT]], getelementptr inbounds ([[CLASS_S3]], ptr @_ZL2ca, i64 5) 259 // CHECK-NEXT: br i1 [[ARRAYCTOR_DONE]], label [[ARRAYCTOR_CONT:%.*]], label [[ARRAYCTOR_LOOP]] 260 // CHECK: arrayctor.cont: 261 // CHECK-NEXT: ret void 262 // 263 // 264 // CHECK-LABEL: define {{[^@]+}}@__cxx_global_var_init.4 265 // CHECK-SAME: () #[[ATTR0]] { 266 // CHECK-NEXT: entry: 267 // CHECK-NEXT: call void @_ZN2S3C1Ev(ptr noundef nonnull align 4 dereferenceable(4) @h) 268 // CHECK-NEXT: ret void 269 // 270 // 271 // CHECK-LABEL: define {{[^@]+}}@main 272 // CHECK-SAME: (i32 noundef signext [[ARGC:%.*]], ptr noundef [[ARGV:%.*]]) #[[ATTR2:[0-9]+]] { 273 // CHECK-NEXT: entry: 274 // CHECK-NEXT: [[RETVAL:%.*]] = alloca i32, align 4 275 // CHECK-NEXT: [[ARGC_ADDR:%.*]] = alloca i32, align 4 276 // CHECK-NEXT: [[ARGV_ADDR:%.*]] = alloca ptr, align 8 277 // CHECK-NEXT: [[DA:%.*]] = alloca [5 x i32], align 4 278 // CHECK-NEXT: [[H:%.*]] = alloca [10 x %struct.S6], align 4 279 // CHECK-NEXT: [[RH:%.*]] = alloca ptr, align 8 280 // CHECK-NEXT: [[I:%.*]] = alloca i32, align 4 281 // CHECK-NEXT: [[J:%.*]] = alloca ptr, align 8 282 // CHECK-NEXT: [[K:%.*]] = alloca ptr, align 8 283 // CHECK-NEXT: [[Z:%.*]] = alloca ptr, align 8 284 // CHECK-NEXT: [[AA:%.*]] = alloca [10 x i32], align 4 285 // CHECK-NEXT: [[RAA:%.*]] = alloca ptr, align 8 286 // CHECK-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x ptr], align 8 287 // CHECK-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x ptr], align 8 288 // CHECK-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [1 x ptr], align 8 289 // CHECK-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8 290 // CHECK-NEXT: [[TMP:%.*]] = alloca ptr, align 8 291 // CHECK-NEXT: [[DOTOFFLOAD_BASEPTRS1:%.*]] = alloca [1 x ptr], align 8 292 // CHECK-NEXT: [[DOTOFFLOAD_PTRS2:%.*]] = alloca [1 x ptr], align 8 293 // CHECK-NEXT: [[DOTOFFLOAD_MAPPERS3:%.*]] = alloca [1 x ptr], align 8 294 // CHECK-NEXT: [[KERNEL_ARGS4:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 8 295 // CHECK-NEXT: [[DOTOFFLOAD_BASEPTRS7:%.*]] = alloca [1 x ptr], align 8 296 // CHECK-NEXT: [[DOTOFFLOAD_PTRS8:%.*]] = alloca [1 x ptr], align 8 297 // CHECK-NEXT: [[DOTOFFLOAD_MAPPERS9:%.*]] = alloca [1 x ptr], align 8 298 // CHECK-NEXT: [[KERNEL_ARGS10:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 8 299 // CHECK-NEXT: [[_TMP13:%.*]] = alloca ptr, align 8 300 // CHECK-NEXT: [[DOTOFFLOAD_BASEPTRS14:%.*]] = alloca [1 x ptr], align 8 301 // CHECK-NEXT: [[DOTOFFLOAD_PTRS15:%.*]] = alloca [1 x ptr], align 8 302 // CHECK-NEXT: [[DOTOFFLOAD_MAPPERS16:%.*]] = alloca [1 x ptr], align 8 303 // CHECK-NEXT: [[KERNEL_ARGS17:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 8 304 // CHECK-NEXT: [[DOTOFFLOAD_BASEPTRS20:%.*]] = alloca [1 x ptr], align 8 305 // CHECK-NEXT: [[DOTOFFLOAD_PTRS21:%.*]] = alloca [1 x ptr], align 8 306 // CHECK-NEXT: [[DOTOFFLOAD_MAPPERS22:%.*]] = alloca [1 x ptr], align 8 307 // CHECK-NEXT: [[KERNEL_ARGS23:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 8 308 // CHECK-NEXT: [[DOTOFFLOAD_BASEPTRS26:%.*]] = alloca [1 x ptr], align 8 309 // CHECK-NEXT: [[DOTOFFLOAD_PTRS27:%.*]] = alloca [1 x ptr], align 8 310 // CHECK-NEXT: [[DOTOFFLOAD_MAPPERS28:%.*]] = alloca [1 x ptr], align 8 311 // CHECK-NEXT: [[KERNEL_ARGS29:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 8 312 // CHECK-NEXT: store i32 0, ptr [[RETVAL]], align 4 313 // CHECK-NEXT: store i32 [[ARGC]], ptr [[ARGC_ADDR]], align 4 314 // CHECK-NEXT: store ptr [[ARGV]], ptr [[ARGV_ADDR]], align 8 315 // CHECK-NEXT: call void @llvm.memset.p0.i64(ptr align 4 [[DA]], i8 0, i64 20, i1 false) 316 // CHECK-NEXT: store ptr [[H]], ptr [[RH]], align 8 317 // CHECK-NEXT: store ptr [[I]], ptr [[J]], align 8 318 // CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[J]], align 8 319 // CHECK-NEXT: store ptr [[TMP0]], ptr [[K]], align 8 320 // CHECK-NEXT: store ptr [[K]], ptr [[Z]], align 8 321 // CHECK-NEXT: store ptr [[AA]], ptr [[RAA]], align 8 322 // CHECK-NEXT: [[TMP1:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 323 // CHECK-NEXT: store ptr [[K]], ptr [[TMP1]], align 8 324 // CHECK-NEXT: [[TMP2:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0 325 // CHECK-NEXT: store ptr [[K]], ptr [[TMP2]], align 8 326 // CHECK-NEXT: [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 0 327 // CHECK-NEXT: store ptr null, ptr [[TMP3]], align 8 328 // CHECK-NEXT: [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 329 // CHECK-NEXT: [[TMP5:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0 330 // CHECK-NEXT: [[TMP6:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0 331 // CHECK-NEXT: store i32 3, ptr [[TMP6]], align 4 332 // CHECK-NEXT: [[TMP7:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1 333 // CHECK-NEXT: store i32 1, ptr [[TMP7]], align 4 334 // CHECK-NEXT: [[TMP8:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2 335 // CHECK-NEXT: store ptr [[TMP4]], ptr [[TMP8]], align 8 336 // CHECK-NEXT: [[TMP9:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3 337 // CHECK-NEXT: store ptr [[TMP5]], ptr [[TMP9]], align 8 338 // CHECK-NEXT: [[TMP10:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4 339 // CHECK-NEXT: store ptr @.offload_sizes, ptr [[TMP10]], align 8 340 // CHECK-NEXT: [[TMP11:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5 341 // CHECK-NEXT: store ptr @.offload_maptypes, ptr [[TMP11]], align 8 342 // CHECK-NEXT: [[TMP12:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6 343 // CHECK-NEXT: store ptr null, ptr [[TMP12]], align 8 344 // CHECK-NEXT: [[TMP13:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7 345 // CHECK-NEXT: store ptr null, ptr [[TMP13]], align 8 346 // CHECK-NEXT: [[TMP14:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8 347 // CHECK-NEXT: store i64 0, ptr [[TMP14]], align 8 348 // CHECK-NEXT: [[TMP15:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9 349 // CHECK-NEXT: store i64 0, ptr [[TMP15]], align 8 350 // CHECK-NEXT: [[TMP16:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10 351 // CHECK-NEXT: store [3 x i32] [i32 -1, i32 0, i32 0], ptr [[TMP16]], align 4 352 // CHECK-NEXT: [[TMP17:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11 353 // CHECK-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP17]], align 4 354 // CHECK-NEXT: [[TMP18:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12 355 // CHECK-NEXT: store i32 0, ptr [[TMP18]], align 4 356 // CHECK-NEXT: [[TMP19:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1:[0-9]+]], i64 -1, i32 -1, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l145.region_id, ptr [[KERNEL_ARGS]]) 357 // CHECK-NEXT: [[TMP20:%.*]] = icmp ne i32 [[TMP19]], 0 358 // CHECK-NEXT: br i1 [[TMP20]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]] 359 // CHECK: omp_offload.failed: 360 // CHECK-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l145(ptr [[K]]) #[[ATTR5:[0-9]+]] 361 // CHECK-NEXT: br label [[OMP_OFFLOAD_CONT]] 362 // CHECK: omp_offload.cont: 363 // CHECK-NEXT: [[TMP21:%.*]] = load ptr, ptr [[Z]], align 8 364 // CHECK-NEXT: store ptr [[TMP21]], ptr [[TMP]], align 8 365 // CHECK-NEXT: [[TMP22:%.*]] = load ptr, ptr [[TMP]], align 8 366 // CHECK-NEXT: [[TMP23:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS1]], i32 0, i32 0 367 // CHECK-NEXT: store ptr [[TMP22]], ptr [[TMP23]], align 8 368 // CHECK-NEXT: [[TMP24:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS2]], i32 0, i32 0 369 // CHECK-NEXT: store ptr [[TMP22]], ptr [[TMP24]], align 8 370 // CHECK-NEXT: [[TMP25:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS3]], i64 0, i64 0 371 // CHECK-NEXT: store ptr null, ptr [[TMP25]], align 8 372 // CHECK-NEXT: [[TMP26:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS1]], i32 0, i32 0 373 // CHECK-NEXT: [[TMP27:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS2]], i32 0, i32 0 374 // CHECK-NEXT: [[TMP28:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS4]], i32 0, i32 0 375 // CHECK-NEXT: store i32 3, ptr [[TMP28]], align 4 376 // CHECK-NEXT: [[TMP29:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS4]], i32 0, i32 1 377 // CHECK-NEXT: store i32 1, ptr [[TMP29]], align 4 378 // CHECK-NEXT: [[TMP30:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS4]], i32 0, i32 2 379 // CHECK-NEXT: store ptr [[TMP26]], ptr [[TMP30]], align 8 380 // CHECK-NEXT: [[TMP31:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS4]], i32 0, i32 3 381 // CHECK-NEXT: store ptr [[TMP27]], ptr [[TMP31]], align 8 382 // CHECK-NEXT: [[TMP32:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS4]], i32 0, i32 4 383 // CHECK-NEXT: store ptr @.offload_sizes.5, ptr [[TMP32]], align 8 384 // CHECK-NEXT: [[TMP33:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS4]], i32 0, i32 5 385 // CHECK-NEXT: store ptr @.offload_maptypes.6, ptr [[TMP33]], align 8 386 // CHECK-NEXT: [[TMP34:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS4]], i32 0, i32 6 387 // CHECK-NEXT: store ptr null, ptr [[TMP34]], align 8 388 // CHECK-NEXT: [[TMP35:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS4]], i32 0, i32 7 389 // CHECK-NEXT: store ptr null, ptr [[TMP35]], align 8 390 // CHECK-NEXT: [[TMP36:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS4]], i32 0, i32 8 391 // CHECK-NEXT: store i64 0, ptr [[TMP36]], align 8 392 // CHECK-NEXT: [[TMP37:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS4]], i32 0, i32 9 393 // CHECK-NEXT: store i64 0, ptr [[TMP37]], align 8 394 // CHECK-NEXT: [[TMP38:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS4]], i32 0, i32 10 395 // CHECK-NEXT: store [3 x i32] [i32 -1, i32 0, i32 0], ptr [[TMP38]], align 4 396 // CHECK-NEXT: [[TMP39:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS4]], i32 0, i32 11 397 // CHECK-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP39]], align 4 398 // CHECK-NEXT: [[TMP40:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS4]], i32 0, i32 12 399 // CHECK-NEXT: store i32 0, ptr [[TMP40]], align 4 400 // CHECK-NEXT: [[TMP41:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 -1, i32 -1, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l147.region_id, ptr [[KERNEL_ARGS4]]) 401 // CHECK-NEXT: [[TMP42:%.*]] = icmp ne i32 [[TMP41]], 0 402 // CHECK-NEXT: br i1 [[TMP42]], label [[OMP_OFFLOAD_FAILED5:%.*]], label [[OMP_OFFLOAD_CONT6:%.*]] 403 // CHECK: omp_offload.failed5: 404 // CHECK-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l147(ptr [[TMP22]]) #[[ATTR5]] 405 // CHECK-NEXT: br label [[OMP_OFFLOAD_CONT6]] 406 // CHECK: omp_offload.cont6: 407 // CHECK-NEXT: [[TMP43:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS7]], i32 0, i32 0 408 // CHECK-NEXT: store ptr [[AA]], ptr [[TMP43]], align 8 409 // CHECK-NEXT: [[TMP44:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS8]], i32 0, i32 0 410 // CHECK-NEXT: store ptr [[AA]], ptr [[TMP44]], align 8 411 // CHECK-NEXT: [[TMP45:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS9]], i64 0, i64 0 412 // CHECK-NEXT: store ptr null, ptr [[TMP45]], align 8 413 // CHECK-NEXT: [[TMP46:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS7]], i32 0, i32 0 414 // CHECK-NEXT: [[TMP47:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS8]], i32 0, i32 0 415 // CHECK-NEXT: [[TMP48:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS10]], i32 0, i32 0 416 // CHECK-NEXT: store i32 3, ptr [[TMP48]], align 4 417 // CHECK-NEXT: [[TMP49:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS10]], i32 0, i32 1 418 // CHECK-NEXT: store i32 1, ptr [[TMP49]], align 4 419 // CHECK-NEXT: [[TMP50:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS10]], i32 0, i32 2 420 // CHECK-NEXT: store ptr [[TMP46]], ptr [[TMP50]], align 8 421 // CHECK-NEXT: [[TMP51:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS10]], i32 0, i32 3 422 // CHECK-NEXT: store ptr [[TMP47]], ptr [[TMP51]], align 8 423 // CHECK-NEXT: [[TMP52:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS10]], i32 0, i32 4 424 // CHECK-NEXT: store ptr @.offload_sizes.7, ptr [[TMP52]], align 8 425 // CHECK-NEXT: [[TMP53:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS10]], i32 0, i32 5 426 // CHECK-NEXT: store ptr @.offload_maptypes.8, ptr [[TMP53]], align 8 427 // CHECK-NEXT: [[TMP54:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS10]], i32 0, i32 6 428 // CHECK-NEXT: store ptr null, ptr [[TMP54]], align 8 429 // CHECK-NEXT: [[TMP55:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS10]], i32 0, i32 7 430 // CHECK-NEXT: store ptr null, ptr [[TMP55]], align 8 431 // CHECK-NEXT: [[TMP56:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS10]], i32 0, i32 8 432 // CHECK-NEXT: store i64 0, ptr [[TMP56]], align 8 433 // CHECK-NEXT: [[TMP57:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS10]], i32 0, i32 9 434 // CHECK-NEXT: store i64 0, ptr [[TMP57]], align 8 435 // CHECK-NEXT: [[TMP58:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS10]], i32 0, i32 10 436 // CHECK-NEXT: store [3 x i32] [i32 -1, i32 0, i32 0], ptr [[TMP58]], align 4 437 // CHECK-NEXT: [[TMP59:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS10]], i32 0, i32 11 438 // CHECK-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP59]], align 4 439 // CHECK-NEXT: [[TMP60:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS10]], i32 0, i32 12 440 // CHECK-NEXT: store i32 0, ptr [[TMP60]], align 4 441 // CHECK-NEXT: [[TMP61:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 -1, i32 -1, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l149.region_id, ptr [[KERNEL_ARGS10]]) 442 // CHECK-NEXT: [[TMP62:%.*]] = icmp ne i32 [[TMP61]], 0 443 // CHECK-NEXT: br i1 [[TMP62]], label [[OMP_OFFLOAD_FAILED11:%.*]], label [[OMP_OFFLOAD_CONT12:%.*]] 444 // CHECK: omp_offload.failed11: 445 // CHECK-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l149(ptr [[AA]]) #[[ATTR5]] 446 // CHECK-NEXT: br label [[OMP_OFFLOAD_CONT12]] 447 // CHECK: omp_offload.cont12: 448 // CHECK-NEXT: [[TMP63:%.*]] = load ptr, ptr [[RAA]], align 8 449 // CHECK-NEXT: store ptr [[TMP63]], ptr [[_TMP13]], align 8 450 // CHECK-NEXT: [[TMP64:%.*]] = load ptr, ptr [[_TMP13]], align 8 451 // CHECK-NEXT: [[TMP65:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS14]], i32 0, i32 0 452 // CHECK-NEXT: store ptr [[TMP64]], ptr [[TMP65]], align 8 453 // CHECK-NEXT: [[TMP66:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS15]], i32 0, i32 0 454 // CHECK-NEXT: store ptr [[TMP64]], ptr [[TMP66]], align 8 455 // CHECK-NEXT: [[TMP67:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS16]], i64 0, i64 0 456 // CHECK-NEXT: store ptr null, ptr [[TMP67]], align 8 457 // CHECK-NEXT: [[TMP68:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS14]], i32 0, i32 0 458 // CHECK-NEXT: [[TMP69:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS15]], i32 0, i32 0 459 // CHECK-NEXT: [[TMP70:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS17]], i32 0, i32 0 460 // CHECK-NEXT: store i32 3, ptr [[TMP70]], align 4 461 // CHECK-NEXT: [[TMP71:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS17]], i32 0, i32 1 462 // CHECK-NEXT: store i32 1, ptr [[TMP71]], align 4 463 // CHECK-NEXT: [[TMP72:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS17]], i32 0, i32 2 464 // CHECK-NEXT: store ptr [[TMP68]], ptr [[TMP72]], align 8 465 // CHECK-NEXT: [[TMP73:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS17]], i32 0, i32 3 466 // CHECK-NEXT: store ptr [[TMP69]], ptr [[TMP73]], align 8 467 // CHECK-NEXT: [[TMP74:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS17]], i32 0, i32 4 468 // CHECK-NEXT: store ptr @.offload_sizes.9, ptr [[TMP74]], align 8 469 // CHECK-NEXT: [[TMP75:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS17]], i32 0, i32 5 470 // CHECK-NEXT: store ptr @.offload_maptypes.10, ptr [[TMP75]], align 8 471 // CHECK-NEXT: [[TMP76:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS17]], i32 0, i32 6 472 // CHECK-NEXT: store ptr null, ptr [[TMP76]], align 8 473 // CHECK-NEXT: [[TMP77:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS17]], i32 0, i32 7 474 // CHECK-NEXT: store ptr null, ptr [[TMP77]], align 8 475 // CHECK-NEXT: [[TMP78:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS17]], i32 0, i32 8 476 // CHECK-NEXT: store i64 0, ptr [[TMP78]], align 8 477 // CHECK-NEXT: [[TMP79:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS17]], i32 0, i32 9 478 // CHECK-NEXT: store i64 0, ptr [[TMP79]], align 8 479 // CHECK-NEXT: [[TMP80:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS17]], i32 0, i32 10 480 // CHECK-NEXT: store [3 x i32] [i32 -1, i32 0, i32 0], ptr [[TMP80]], align 4 481 // CHECK-NEXT: [[TMP81:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS17]], i32 0, i32 11 482 // CHECK-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP81]], align 4 483 // CHECK-NEXT: [[TMP82:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS17]], i32 0, i32 12 484 // CHECK-NEXT: store i32 0, ptr [[TMP82]], align 4 485 // CHECK-NEXT: [[TMP83:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 -1, i32 -1, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l151.region_id, ptr [[KERNEL_ARGS17]]) 486 // CHECK-NEXT: [[TMP84:%.*]] = icmp ne i32 [[TMP83]], 0 487 // CHECK-NEXT: br i1 [[TMP84]], label [[OMP_OFFLOAD_FAILED18:%.*]], label [[OMP_OFFLOAD_CONT19:%.*]] 488 // CHECK: omp_offload.failed18: 489 // CHECK-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l151(ptr [[TMP64]]) #[[ATTR5]] 490 // CHECK-NEXT: br label [[OMP_OFFLOAD_CONT19]] 491 // CHECK: omp_offload.cont19: 492 // CHECK-NEXT: [[TMP85:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS20]], i32 0, i32 0 493 // CHECK-NEXT: store ptr [[H]], ptr [[TMP85]], align 8 494 // CHECK-NEXT: [[TMP86:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS21]], i32 0, i32 0 495 // CHECK-NEXT: store ptr [[H]], ptr [[TMP86]], align 8 496 // CHECK-NEXT: [[TMP87:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS22]], i64 0, i64 0 497 // CHECK-NEXT: store ptr null, ptr [[TMP87]], align 8 498 // CHECK-NEXT: [[TMP88:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS20]], i32 0, i32 0 499 // CHECK-NEXT: [[TMP89:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS21]], i32 0, i32 0 500 // CHECK-NEXT: [[TMP90:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS23]], i32 0, i32 0 501 // CHECK-NEXT: store i32 3, ptr [[TMP90]], align 4 502 // CHECK-NEXT: [[TMP91:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS23]], i32 0, i32 1 503 // CHECK-NEXT: store i32 1, ptr [[TMP91]], align 4 504 // CHECK-NEXT: [[TMP92:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS23]], i32 0, i32 2 505 // CHECK-NEXT: store ptr [[TMP88]], ptr [[TMP92]], align 8 506 // CHECK-NEXT: [[TMP93:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS23]], i32 0, i32 3 507 // CHECK-NEXT: store ptr [[TMP89]], ptr [[TMP93]], align 8 508 // CHECK-NEXT: [[TMP94:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS23]], i32 0, i32 4 509 // CHECK-NEXT: store ptr @.offload_sizes.11, ptr [[TMP94]], align 8 510 // CHECK-NEXT: [[TMP95:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS23]], i32 0, i32 5 511 // CHECK-NEXT: store ptr @.offload_maptypes.12, ptr [[TMP95]], align 8 512 // CHECK-NEXT: [[TMP96:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS23]], i32 0, i32 6 513 // CHECK-NEXT: store ptr null, ptr [[TMP96]], align 8 514 // CHECK-NEXT: [[TMP97:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS23]], i32 0, i32 7 515 // CHECK-NEXT: store ptr null, ptr [[TMP97]], align 8 516 // CHECK-NEXT: [[TMP98:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS23]], i32 0, i32 8 517 // CHECK-NEXT: store i64 0, ptr [[TMP98]], align 8 518 // CHECK-NEXT: [[TMP99:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS23]], i32 0, i32 9 519 // CHECK-NEXT: store i64 0, ptr [[TMP99]], align 8 520 // CHECK-NEXT: [[TMP100:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS23]], i32 0, i32 10 521 // CHECK-NEXT: store [3 x i32] [i32 -1, i32 0, i32 0], ptr [[TMP100]], align 4 522 // CHECK-NEXT: [[TMP101:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS23]], i32 0, i32 11 523 // CHECK-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP101]], align 4 524 // CHECK-NEXT: [[TMP102:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS23]], i32 0, i32 12 525 // CHECK-NEXT: store i32 0, ptr [[TMP102]], align 4 526 // CHECK-NEXT: [[TMP103:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 -1, i32 -1, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l153.region_id, ptr [[KERNEL_ARGS23]]) 527 // CHECK-NEXT: [[TMP104:%.*]] = icmp ne i32 [[TMP103]], 0 528 // CHECK-NEXT: br i1 [[TMP104]], label [[OMP_OFFLOAD_FAILED24:%.*]], label [[OMP_OFFLOAD_CONT25:%.*]] 529 // CHECK: omp_offload.failed24: 530 // CHECK-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l153(ptr [[H]]) #[[ATTR5]] 531 // CHECK-NEXT: br label [[OMP_OFFLOAD_CONT25]] 532 // CHECK: omp_offload.cont25: 533 // CHECK-NEXT: [[TMP105:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS26]], i32 0, i32 0 534 // CHECK-NEXT: store ptr [[DA]], ptr [[TMP105]], align 8 535 // CHECK-NEXT: [[TMP106:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS27]], i32 0, i32 0 536 // CHECK-NEXT: store ptr [[DA]], ptr [[TMP106]], align 8 537 // CHECK-NEXT: [[TMP107:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS28]], i64 0, i64 0 538 // CHECK-NEXT: store ptr null, ptr [[TMP107]], align 8 539 // CHECK-NEXT: [[TMP108:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS26]], i32 0, i32 0 540 // CHECK-NEXT: [[TMP109:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS27]], i32 0, i32 0 541 // CHECK-NEXT: [[TMP110:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS29]], i32 0, i32 0 542 // CHECK-NEXT: store i32 3, ptr [[TMP110]], align 4 543 // CHECK-NEXT: [[TMP111:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS29]], i32 0, i32 1 544 // CHECK-NEXT: store i32 1, ptr [[TMP111]], align 4 545 // CHECK-NEXT: [[TMP112:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS29]], i32 0, i32 2 546 // CHECK-NEXT: store ptr [[TMP108]], ptr [[TMP112]], align 8 547 // CHECK-NEXT: [[TMP113:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS29]], i32 0, i32 3 548 // CHECK-NEXT: store ptr [[TMP109]], ptr [[TMP113]], align 8 549 // CHECK-NEXT: [[TMP114:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS29]], i32 0, i32 4 550 // CHECK-NEXT: store ptr @.offload_sizes.13, ptr [[TMP114]], align 8 551 // CHECK-NEXT: [[TMP115:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS29]], i32 0, i32 5 552 // CHECK-NEXT: store ptr @.offload_maptypes.14, ptr [[TMP115]], align 8 553 // CHECK-NEXT: [[TMP116:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS29]], i32 0, i32 6 554 // CHECK-NEXT: store ptr null, ptr [[TMP116]], align 8 555 // CHECK-NEXT: [[TMP117:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS29]], i32 0, i32 7 556 // CHECK-NEXT: store ptr null, ptr [[TMP117]], align 8 557 // CHECK-NEXT: [[TMP118:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS29]], i32 0, i32 8 558 // CHECK-NEXT: store i64 0, ptr [[TMP118]], align 8 559 // CHECK-NEXT: [[TMP119:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS29]], i32 0, i32 9 560 // CHECK-NEXT: store i64 0, ptr [[TMP119]], align 8 561 // CHECK-NEXT: [[TMP120:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS29]], i32 0, i32 10 562 // CHECK-NEXT: store [3 x i32] [i32 -1, i32 0, i32 0], ptr [[TMP120]], align 4 563 // CHECK-NEXT: [[TMP121:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS29]], i32 0, i32 11 564 // CHECK-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP121]], align 4 565 // CHECK-NEXT: [[TMP122:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS29]], i32 0, i32 12 566 // CHECK-NEXT: store i32 0, ptr [[TMP122]], align 4 567 // CHECK-NEXT: [[TMP123:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 -1, i32 -1, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l155.region_id, ptr [[KERNEL_ARGS29]]) 568 // CHECK-NEXT: [[TMP124:%.*]] = icmp ne i32 [[TMP123]], 0 569 // CHECK-NEXT: br i1 [[TMP124]], label [[OMP_OFFLOAD_FAILED30:%.*]], label [[OMP_OFFLOAD_CONT31:%.*]] 570 // CHECK: omp_offload.failed30: 571 // CHECK-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l155(ptr [[DA]]) #[[ATTR5]] 572 // CHECK-NEXT: br label [[OMP_OFFLOAD_CONT31]] 573 // CHECK: omp_offload.cont31: 574 // CHECK-NEXT: [[TMP125:%.*]] = load i32, ptr [[ARGC_ADDR]], align 4 575 // CHECK-NEXT: [[CALL:%.*]] = call noundef signext i32 @_Z5tmainIiET_S0_(i32 noundef signext [[TMP125]]) 576 // CHECK-NEXT: [[CALL32:%.*]] = call noundef ptr @_Z5tmainIPiET_S1_(ptr noundef [[ARGC_ADDR]]) 577 // CHECK-NEXT: [[TMP126:%.*]] = load i32, ptr [[CALL32]], align 4 578 // CHECK-NEXT: [[ADD:%.*]] = add nsw i32 [[CALL]], [[TMP126]] 579 // CHECK-NEXT: ret i32 [[ADD]] 580 // 581 // 582 // CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l145 583 // CHECK-SAME: (ptr noundef nonnull align 8 dereferenceable(8) [[K:%.*]]) #[[ATTR4:[0-9]+]] { 584 // CHECK-NEXT: entry: 585 // CHECK-NEXT: [[K_ADDR:%.*]] = alloca ptr, align 8 586 // CHECK-NEXT: store ptr [[K]], ptr [[K_ADDR]], align 8 587 // CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[K_ADDR]], align 8 588 // CHECK-NEXT: [[TMP1:%.*]] = load ptr, ptr [[TMP0]], align 8 589 // CHECK-NEXT: [[INCDEC_PTR:%.*]] = getelementptr inbounds nuw i32, ptr [[TMP1]], i32 1 590 // CHECK-NEXT: store ptr [[INCDEC_PTR]], ptr [[TMP0]], align 8 591 // CHECK-NEXT: ret void 592 // 593 // 594 // CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l147 595 // CHECK-SAME: (ptr noundef nonnull align 8 dereferenceable(8) [[Z:%.*]]) #[[ATTR4]] { 596 // CHECK-NEXT: entry: 597 // CHECK-NEXT: [[Z_ADDR:%.*]] = alloca ptr, align 8 598 // CHECK-NEXT: [[TMP:%.*]] = alloca ptr, align 8 599 // CHECK-NEXT: store ptr [[Z]], ptr [[Z_ADDR]], align 8 600 // CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[Z_ADDR]], align 8 601 // CHECK-NEXT: store ptr [[TMP0]], ptr [[TMP]], align 8 602 // CHECK-NEXT: [[TMP1:%.*]] = load ptr, ptr [[TMP]], align 8 603 // CHECK-NEXT: [[TMP2:%.*]] = load ptr, ptr [[TMP1]], align 8 604 // CHECK-NEXT: [[INCDEC_PTR:%.*]] = getelementptr inbounds nuw i32, ptr [[TMP2]], i32 1 605 // CHECK-NEXT: store ptr [[INCDEC_PTR]], ptr [[TMP1]], align 8 606 // CHECK-NEXT: ret void 607 // 608 // 609 // CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l149 610 // CHECK-SAME: (ptr noundef nonnull align 4 dereferenceable(40) [[AA:%.*]]) #[[ATTR4]] { 611 // CHECK-NEXT: entry: 612 // CHECK-NEXT: [[AA_ADDR:%.*]] = alloca ptr, align 8 613 // CHECK-NEXT: store ptr [[AA]], ptr [[AA_ADDR]], align 8 614 // CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[AA_ADDR]], align 8 615 // CHECK-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x i32], ptr [[TMP0]], i64 0, i64 0 616 // CHECK-NEXT: store i32 1, ptr [[ARRAYIDX]], align 4 617 // CHECK-NEXT: ret void 618 // 619 // 620 // CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l151 621 // CHECK-SAME: (ptr noundef nonnull align 4 dereferenceable(40) [[RAA:%.*]]) #[[ATTR4]] { 622 // CHECK-NEXT: entry: 623 // CHECK-NEXT: [[RAA_ADDR:%.*]] = alloca ptr, align 8 624 // CHECK-NEXT: [[TMP:%.*]] = alloca ptr, align 8 625 // CHECK-NEXT: [[A:%.*]] = alloca i32, align 4 626 // CHECK-NEXT: store ptr [[RAA]], ptr [[RAA_ADDR]], align 8 627 // CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[RAA_ADDR]], align 8 628 // CHECK-NEXT: store ptr [[TMP0]], ptr [[TMP]], align 8 629 // CHECK-NEXT: [[TMP1:%.*]] = load ptr, ptr [[TMP]], align 8 630 // CHECK-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x i32], ptr [[TMP1]], i64 0, i64 0 631 // CHECK-NEXT: [[TMP2:%.*]] = load i32, ptr [[ARRAYIDX]], align 4 632 // CHECK-NEXT: store i32 [[TMP2]], ptr [[A]], align 4 633 // CHECK-NEXT: ret void 634 // 635 // 636 // CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l153 637 // CHECK-SAME: (ptr noundef nonnull align 4 dereferenceable(40) [[H:%.*]]) #[[ATTR4]] { 638 // CHECK-NEXT: entry: 639 // CHECK-NEXT: [[H_ADDR:%.*]] = alloca ptr, align 8 640 // CHECK-NEXT: [[A:%.*]] = alloca i32, align 4 641 // CHECK-NEXT: store ptr [[H]], ptr [[H_ADDR]], align 8 642 // CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[H_ADDR]], align 8 643 // CHECK-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x %struct.S6], ptr [[TMP0]], i64 0, i64 1 644 // CHECK-NEXT: [[A1:%.*]] = getelementptr inbounds nuw [[STRUCT_S6:%.*]], ptr [[ARRAYIDX]], i32 0, i32 0 645 // CHECK-NEXT: [[TMP1:%.*]] = load i32, ptr [[A1]], align 4 646 // CHECK-NEXT: store i32 [[TMP1]], ptr [[A]], align 4 647 // CHECK-NEXT: ret void 648 // 649 // 650 // CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l155 651 // CHECK-SAME: (ptr noundef nonnull align 4 dereferenceable(20) [[DA:%.*]]) #[[ATTR4]] { 652 // CHECK-NEXT: entry: 653 // CHECK-NEXT: [[DA_ADDR:%.*]] = alloca ptr, align 8 654 // CHECK-NEXT: [[A:%.*]] = alloca i32, align 4 655 // CHECK-NEXT: store ptr [[DA]], ptr [[DA_ADDR]], align 8 656 // CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[DA_ADDR]], align 8 657 // CHECK-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [5 x i32], ptr [[TMP0]], i64 0, i64 1 658 // CHECK-NEXT: [[TMP1:%.*]] = load i32, ptr [[ARRAYIDX]], align 4 659 // CHECK-NEXT: store i32 [[TMP1]], ptr [[A]], align 4 660 // CHECK-NEXT: ret void 661 // 662 // 663 // CHECK-LABEL: define {{[^@]+}}@_Z5tmainIiET_S0_ 664 // CHECK-SAME: (i32 noundef signext [[ARGC:%.*]]) #[[ATTR1]] comdat { 665 // CHECK-NEXT: entry: 666 // CHECK-NEXT: [[ARGC_ADDR:%.*]] = alloca i32, align 4 667 // CHECK-NEXT: [[DA:%.*]] = alloca [5 x i32], align 4 668 // CHECK-NEXT: [[H:%.*]] = alloca [10 x %struct.S6], align 4 669 // CHECK-NEXT: [[RH:%.*]] = alloca ptr, align 8 670 // CHECK-NEXT: [[I:%.*]] = alloca i32, align 4 671 // CHECK-NEXT: [[J:%.*]] = alloca ptr, align 8 672 // CHECK-NEXT: [[K:%.*]] = alloca ptr, align 8 673 // CHECK-NEXT: [[Z:%.*]] = alloca ptr, align 8 674 // CHECK-NEXT: [[AA:%.*]] = alloca [10 x i32], align 4 675 // CHECK-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x ptr], align 8 676 // CHECK-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x ptr], align 8 677 // CHECK-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [1 x ptr], align 8 678 // CHECK-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8 679 // CHECK-NEXT: [[TMP:%.*]] = alloca ptr, align 8 680 // CHECK-NEXT: [[DOTOFFLOAD_BASEPTRS1:%.*]] = alloca [1 x ptr], align 8 681 // CHECK-NEXT: [[DOTOFFLOAD_PTRS2:%.*]] = alloca [1 x ptr], align 8 682 // CHECK-NEXT: [[DOTOFFLOAD_MAPPERS3:%.*]] = alloca [1 x ptr], align 8 683 // CHECK-NEXT: [[KERNEL_ARGS4:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 8 684 // CHECK-NEXT: [[DOTOFFLOAD_BASEPTRS7:%.*]] = alloca [1 x ptr], align 8 685 // CHECK-NEXT: [[DOTOFFLOAD_PTRS8:%.*]] = alloca [1 x ptr], align 8 686 // CHECK-NEXT: [[DOTOFFLOAD_MAPPERS9:%.*]] = alloca [1 x ptr], align 8 687 // CHECK-NEXT: [[KERNEL_ARGS10:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 8 688 // CHECK-NEXT: [[DOTOFFLOAD_BASEPTRS13:%.*]] = alloca [1 x ptr], align 8 689 // CHECK-NEXT: [[DOTOFFLOAD_PTRS14:%.*]] = alloca [1 x ptr], align 8 690 // CHECK-NEXT: [[DOTOFFLOAD_MAPPERS15:%.*]] = alloca [1 x ptr], align 8 691 // CHECK-NEXT: [[KERNEL_ARGS16:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 8 692 // CHECK-NEXT: store i32 [[ARGC]], ptr [[ARGC_ADDR]], align 4 693 // CHECK-NEXT: call void @llvm.memset.p0.i64(ptr align 4 [[DA]], i8 0, i64 20, i1 false) 694 // CHECK-NEXT: store ptr [[H]], ptr [[RH]], align 8 695 // CHECK-NEXT: store ptr [[I]], ptr [[J]], align 8 696 // CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[J]], align 8 697 // CHECK-NEXT: store ptr [[TMP0]], ptr [[K]], align 8 698 // CHECK-NEXT: store ptr [[K]], ptr [[Z]], align 8 699 // CHECK-NEXT: [[TMP1:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 700 // CHECK-NEXT: store ptr [[K]], ptr [[TMP1]], align 8 701 // CHECK-NEXT: [[TMP2:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0 702 // CHECK-NEXT: store ptr [[K]], ptr [[TMP2]], align 8 703 // CHECK-NEXT: [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 0 704 // CHECK-NEXT: store ptr null, ptr [[TMP3]], align 8 705 // CHECK-NEXT: [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 706 // CHECK-NEXT: [[TMP5:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0 707 // CHECK-NEXT: [[TMP6:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0 708 // CHECK-NEXT: store i32 3, ptr [[TMP6]], align 4 709 // CHECK-NEXT: [[TMP7:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1 710 // CHECK-NEXT: store i32 1, ptr [[TMP7]], align 4 711 // CHECK-NEXT: [[TMP8:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2 712 // CHECK-NEXT: store ptr [[TMP4]], ptr [[TMP8]], align 8 713 // CHECK-NEXT: [[TMP9:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3 714 // CHECK-NEXT: store ptr [[TMP5]], ptr [[TMP9]], align 8 715 // CHECK-NEXT: [[TMP10:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4 716 // CHECK-NEXT: store ptr @.offload_sizes.15, ptr [[TMP10]], align 8 717 // CHECK-NEXT: [[TMP11:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5 718 // CHECK-NEXT: store ptr @.offload_maptypes.16, ptr [[TMP11]], align 8 719 // CHECK-NEXT: [[TMP12:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6 720 // CHECK-NEXT: store ptr null, ptr [[TMP12]], align 8 721 // CHECK-NEXT: [[TMP13:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7 722 // CHECK-NEXT: store ptr null, ptr [[TMP13]], align 8 723 // CHECK-NEXT: [[TMP14:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8 724 // CHECK-NEXT: store i64 0, ptr [[TMP14]], align 8 725 // CHECK-NEXT: [[TMP15:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9 726 // CHECK-NEXT: store i64 0, ptr [[TMP15]], align 8 727 // CHECK-NEXT: [[TMP16:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10 728 // CHECK-NEXT: store [3 x i32] [i32 -1, i32 0, i32 0], ptr [[TMP16]], align 4 729 // CHECK-NEXT: [[TMP17:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11 730 // CHECK-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP17]], align 4 731 // CHECK-NEXT: [[TMP18:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12 732 // CHECK-NEXT: store i32 0, ptr [[TMP18]], align 4 733 // CHECK-NEXT: [[TMP19:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 -1, i32 -1, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_S0__l123.region_id, ptr [[KERNEL_ARGS]]) 734 // CHECK-NEXT: [[TMP20:%.*]] = icmp ne i32 [[TMP19]], 0 735 // CHECK-NEXT: br i1 [[TMP20]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]] 736 // CHECK: omp_offload.failed: 737 // CHECK-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_S0__l123(ptr [[K]]) #[[ATTR5]] 738 // CHECK-NEXT: br label [[OMP_OFFLOAD_CONT]] 739 // CHECK: omp_offload.cont: 740 // CHECK-NEXT: [[TMP21:%.*]] = load ptr, ptr [[Z]], align 8 741 // CHECK-NEXT: store ptr [[TMP21]], ptr [[TMP]], align 8 742 // CHECK-NEXT: [[TMP22:%.*]] = load ptr, ptr [[TMP]], align 8 743 // CHECK-NEXT: [[TMP23:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS1]], i32 0, i32 0 744 // CHECK-NEXT: store ptr [[TMP22]], ptr [[TMP23]], align 8 745 // CHECK-NEXT: [[TMP24:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS2]], i32 0, i32 0 746 // CHECK-NEXT: store ptr [[TMP22]], ptr [[TMP24]], align 8 747 // CHECK-NEXT: [[TMP25:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS3]], i64 0, i64 0 748 // CHECK-NEXT: store ptr null, ptr [[TMP25]], align 8 749 // CHECK-NEXT: [[TMP26:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS1]], i32 0, i32 0 750 // CHECK-NEXT: [[TMP27:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS2]], i32 0, i32 0 751 // CHECK-NEXT: [[TMP28:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS4]], i32 0, i32 0 752 // CHECK-NEXT: store i32 3, ptr [[TMP28]], align 4 753 // CHECK-NEXT: [[TMP29:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS4]], i32 0, i32 1 754 // CHECK-NEXT: store i32 1, ptr [[TMP29]], align 4 755 // CHECK-NEXT: [[TMP30:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS4]], i32 0, i32 2 756 // CHECK-NEXT: store ptr [[TMP26]], ptr [[TMP30]], align 8 757 // CHECK-NEXT: [[TMP31:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS4]], i32 0, i32 3 758 // CHECK-NEXT: store ptr [[TMP27]], ptr [[TMP31]], align 8 759 // CHECK-NEXT: [[TMP32:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS4]], i32 0, i32 4 760 // CHECK-NEXT: store ptr @.offload_sizes.17, ptr [[TMP32]], align 8 761 // CHECK-NEXT: [[TMP33:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS4]], i32 0, i32 5 762 // CHECK-NEXT: store ptr @.offload_maptypes.18, ptr [[TMP33]], align 8 763 // CHECK-NEXT: [[TMP34:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS4]], i32 0, i32 6 764 // CHECK-NEXT: store ptr null, ptr [[TMP34]], align 8 765 // CHECK-NEXT: [[TMP35:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS4]], i32 0, i32 7 766 // CHECK-NEXT: store ptr null, ptr [[TMP35]], align 8 767 // CHECK-NEXT: [[TMP36:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS4]], i32 0, i32 8 768 // CHECK-NEXT: store i64 0, ptr [[TMP36]], align 8 769 // CHECK-NEXT: [[TMP37:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS4]], i32 0, i32 9 770 // CHECK-NEXT: store i64 0, ptr [[TMP37]], align 8 771 // CHECK-NEXT: [[TMP38:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS4]], i32 0, i32 10 772 // CHECK-NEXT: store [3 x i32] [i32 -1, i32 0, i32 0], ptr [[TMP38]], align 4 773 // CHECK-NEXT: [[TMP39:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS4]], i32 0, i32 11 774 // CHECK-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP39]], align 4 775 // CHECK-NEXT: [[TMP40:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS4]], i32 0, i32 12 776 // CHECK-NEXT: store i32 0, ptr [[TMP40]], align 4 777 // CHECK-NEXT: [[TMP41:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 -1, i32 -1, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_S0__l125.region_id, ptr [[KERNEL_ARGS4]]) 778 // CHECK-NEXT: [[TMP42:%.*]] = icmp ne i32 [[TMP41]], 0 779 // CHECK-NEXT: br i1 [[TMP42]], label [[OMP_OFFLOAD_FAILED5:%.*]], label [[OMP_OFFLOAD_CONT6:%.*]] 780 // CHECK: omp_offload.failed5: 781 // CHECK-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_S0__l125(ptr [[TMP22]]) #[[ATTR5]] 782 // CHECK-NEXT: br label [[OMP_OFFLOAD_CONT6]] 783 // CHECK: omp_offload.cont6: 784 // CHECK-NEXT: [[TMP43:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS7]], i32 0, i32 0 785 // CHECK-NEXT: store ptr [[AA]], ptr [[TMP43]], align 8 786 // CHECK-NEXT: [[TMP44:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS8]], i32 0, i32 0 787 // CHECK-NEXT: store ptr [[AA]], ptr [[TMP44]], align 8 788 // CHECK-NEXT: [[TMP45:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS9]], i64 0, i64 0 789 // CHECK-NEXT: store ptr null, ptr [[TMP45]], align 8 790 // CHECK-NEXT: [[TMP46:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS7]], i32 0, i32 0 791 // CHECK-NEXT: [[TMP47:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS8]], i32 0, i32 0 792 // CHECK-NEXT: [[TMP48:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS10]], i32 0, i32 0 793 // CHECK-NEXT: store i32 3, ptr [[TMP48]], align 4 794 // CHECK-NEXT: [[TMP49:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS10]], i32 0, i32 1 795 // CHECK-NEXT: store i32 1, ptr [[TMP49]], align 4 796 // CHECK-NEXT: [[TMP50:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS10]], i32 0, i32 2 797 // CHECK-NEXT: store ptr [[TMP46]], ptr [[TMP50]], align 8 798 // CHECK-NEXT: [[TMP51:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS10]], i32 0, i32 3 799 // CHECK-NEXT: store ptr [[TMP47]], ptr [[TMP51]], align 8 800 // CHECK-NEXT: [[TMP52:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS10]], i32 0, i32 4 801 // CHECK-NEXT: store ptr @.offload_sizes.19, ptr [[TMP52]], align 8 802 // CHECK-NEXT: [[TMP53:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS10]], i32 0, i32 5 803 // CHECK-NEXT: store ptr @.offload_maptypes.20, ptr [[TMP53]], align 8 804 // CHECK-NEXT: [[TMP54:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS10]], i32 0, i32 6 805 // CHECK-NEXT: store ptr null, ptr [[TMP54]], align 8 806 // CHECK-NEXT: [[TMP55:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS10]], i32 0, i32 7 807 // CHECK-NEXT: store ptr null, ptr [[TMP55]], align 8 808 // CHECK-NEXT: [[TMP56:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS10]], i32 0, i32 8 809 // CHECK-NEXT: store i64 0, ptr [[TMP56]], align 8 810 // CHECK-NEXT: [[TMP57:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS10]], i32 0, i32 9 811 // CHECK-NEXT: store i64 0, ptr [[TMP57]], align 8 812 // CHECK-NEXT: [[TMP58:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS10]], i32 0, i32 10 813 // CHECK-NEXT: store [3 x i32] [i32 -1, i32 0, i32 0], ptr [[TMP58]], align 4 814 // CHECK-NEXT: [[TMP59:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS10]], i32 0, i32 11 815 // CHECK-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP59]], align 4 816 // CHECK-NEXT: [[TMP60:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS10]], i32 0, i32 12 817 // CHECK-NEXT: store i32 0, ptr [[TMP60]], align 4 818 // CHECK-NEXT: [[TMP61:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 -1, i32 -1, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_S0__l127.region_id, ptr [[KERNEL_ARGS10]]) 819 // CHECK-NEXT: [[TMP62:%.*]] = icmp ne i32 [[TMP61]], 0 820 // CHECK-NEXT: br i1 [[TMP62]], label [[OMP_OFFLOAD_FAILED11:%.*]], label [[OMP_OFFLOAD_CONT12:%.*]] 821 // CHECK: omp_offload.failed11: 822 // CHECK-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_S0__l127(ptr [[AA]]) #[[ATTR5]] 823 // CHECK-NEXT: br label [[OMP_OFFLOAD_CONT12]] 824 // CHECK: omp_offload.cont12: 825 // CHECK-NEXT: [[TMP63:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS13]], i32 0, i32 0 826 // CHECK-NEXT: store ptr [[H]], ptr [[TMP63]], align 8 827 // CHECK-NEXT: [[TMP64:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS14]], i32 0, i32 0 828 // CHECK-NEXT: store ptr [[H]], ptr [[TMP64]], align 8 829 // CHECK-NEXT: [[TMP65:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS15]], i64 0, i64 0 830 // CHECK-NEXT: store ptr null, ptr [[TMP65]], align 8 831 // CHECK-NEXT: [[TMP66:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS13]], i32 0, i32 0 832 // CHECK-NEXT: [[TMP67:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS14]], i32 0, i32 0 833 // CHECK-NEXT: [[TMP68:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS16]], i32 0, i32 0 834 // CHECK-NEXT: store i32 3, ptr [[TMP68]], align 4 835 // CHECK-NEXT: [[TMP69:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS16]], i32 0, i32 1 836 // CHECK-NEXT: store i32 1, ptr [[TMP69]], align 4 837 // CHECK-NEXT: [[TMP70:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS16]], i32 0, i32 2 838 // CHECK-NEXT: store ptr [[TMP66]], ptr [[TMP70]], align 8 839 // CHECK-NEXT: [[TMP71:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS16]], i32 0, i32 3 840 // CHECK-NEXT: store ptr [[TMP67]], ptr [[TMP71]], align 8 841 // CHECK-NEXT: [[TMP72:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS16]], i32 0, i32 4 842 // CHECK-NEXT: store ptr @.offload_sizes.21, ptr [[TMP72]], align 8 843 // CHECK-NEXT: [[TMP73:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS16]], i32 0, i32 5 844 // CHECK-NEXT: store ptr @.offload_maptypes.22, ptr [[TMP73]], align 8 845 // CHECK-NEXT: [[TMP74:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS16]], i32 0, i32 6 846 // CHECK-NEXT: store ptr null, ptr [[TMP74]], align 8 847 // CHECK-NEXT: [[TMP75:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS16]], i32 0, i32 7 848 // CHECK-NEXT: store ptr null, ptr [[TMP75]], align 8 849 // CHECK-NEXT: [[TMP76:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS16]], i32 0, i32 8 850 // CHECK-NEXT: store i64 0, ptr [[TMP76]], align 8 851 // CHECK-NEXT: [[TMP77:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS16]], i32 0, i32 9 852 // CHECK-NEXT: store i64 0, ptr [[TMP77]], align 8 853 // CHECK-NEXT: [[TMP78:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS16]], i32 0, i32 10 854 // CHECK-NEXT: store [3 x i32] [i32 -1, i32 0, i32 0], ptr [[TMP78]], align 4 855 // CHECK-NEXT: [[TMP79:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS16]], i32 0, i32 11 856 // CHECK-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP79]], align 4 857 // CHECK-NEXT: [[TMP80:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS16]], i32 0, i32 12 858 // CHECK-NEXT: store i32 0, ptr [[TMP80]], align 4 859 // CHECK-NEXT: [[TMP81:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 -1, i32 -1, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_S0__l129.region_id, ptr [[KERNEL_ARGS16]]) 860 // CHECK-NEXT: [[TMP82:%.*]] = icmp ne i32 [[TMP81]], 0 861 // CHECK-NEXT: br i1 [[TMP82]], label [[OMP_OFFLOAD_FAILED17:%.*]], label [[OMP_OFFLOAD_CONT18:%.*]] 862 // CHECK: omp_offload.failed17: 863 // CHECK-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_S0__l129(ptr [[H]]) #[[ATTR5]] 864 // CHECK-NEXT: br label [[OMP_OFFLOAD_CONT18]] 865 // CHECK: omp_offload.cont18: 866 // CHECK-NEXT: ret i32 0 867 // 868 // 869 // CHECK-LABEL: define {{[^@]+}}@_Z5tmainIPiET_S1_ 870 // CHECK-SAME: (ptr noundef [[ARGC:%.*]]) #[[ATTR1]] comdat { 871 // CHECK-NEXT: entry: 872 // CHECK-NEXT: [[ARGC_ADDR:%.*]] = alloca ptr, align 8 873 // CHECK-NEXT: [[DA:%.*]] = alloca [5 x ptr], align 8 874 // CHECK-NEXT: [[H:%.*]] = alloca [10 x %struct.S6], align 4 875 // CHECK-NEXT: [[RH:%.*]] = alloca ptr, align 8 876 // CHECK-NEXT: [[I:%.*]] = alloca ptr, align 8 877 // CHECK-NEXT: [[J:%.*]] = alloca ptr, align 8 878 // CHECK-NEXT: [[K:%.*]] = alloca ptr, align 8 879 // CHECK-NEXT: [[Z:%.*]] = alloca ptr, align 8 880 // CHECK-NEXT: [[AA:%.*]] = alloca [10 x ptr], align 8 881 // CHECK-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x ptr], align 8 882 // CHECK-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x ptr], align 8 883 // CHECK-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [1 x ptr], align 8 884 // CHECK-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8 885 // CHECK-NEXT: [[TMP:%.*]] = alloca ptr, align 8 886 // CHECK-NEXT: [[DOTOFFLOAD_BASEPTRS1:%.*]] = alloca [1 x ptr], align 8 887 // CHECK-NEXT: [[DOTOFFLOAD_PTRS2:%.*]] = alloca [1 x ptr], align 8 888 // CHECK-NEXT: [[DOTOFFLOAD_MAPPERS3:%.*]] = alloca [1 x ptr], align 8 889 // CHECK-NEXT: [[KERNEL_ARGS4:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 8 890 // CHECK-NEXT: [[DOTOFFLOAD_BASEPTRS7:%.*]] = alloca [1 x ptr], align 8 891 // CHECK-NEXT: [[DOTOFFLOAD_PTRS8:%.*]] = alloca [1 x ptr], align 8 892 // CHECK-NEXT: [[DOTOFFLOAD_MAPPERS9:%.*]] = alloca [1 x ptr], align 8 893 // CHECK-NEXT: [[KERNEL_ARGS10:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 8 894 // CHECK-NEXT: [[DOTOFFLOAD_BASEPTRS13:%.*]] = alloca [1 x ptr], align 8 895 // CHECK-NEXT: [[DOTOFFLOAD_PTRS14:%.*]] = alloca [1 x ptr], align 8 896 // CHECK-NEXT: [[DOTOFFLOAD_MAPPERS15:%.*]] = alloca [1 x ptr], align 8 897 // CHECK-NEXT: [[KERNEL_ARGS16:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 8 898 // CHECK-NEXT: store ptr [[ARGC]], ptr [[ARGC_ADDR]], align 8 899 // CHECK-NEXT: call void @llvm.memset.p0.i64(ptr align 8 [[DA]], i8 0, i64 40, i1 false) 900 // CHECK-NEXT: store ptr [[H]], ptr [[RH]], align 8 901 // CHECK-NEXT: store ptr [[I]], ptr [[J]], align 8 902 // CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[J]], align 8 903 // CHECK-NEXT: store ptr [[TMP0]], ptr [[K]], align 8 904 // CHECK-NEXT: store ptr [[K]], ptr [[Z]], align 8 905 // CHECK-NEXT: [[TMP1:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 906 // CHECK-NEXT: store ptr [[K]], ptr [[TMP1]], align 8 907 // CHECK-NEXT: [[TMP2:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0 908 // CHECK-NEXT: store ptr [[K]], ptr [[TMP2]], align 8 909 // CHECK-NEXT: [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 0 910 // CHECK-NEXT: store ptr null, ptr [[TMP3]], align 8 911 // CHECK-NEXT: [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 912 // CHECK-NEXT: [[TMP5:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0 913 // CHECK-NEXT: [[TMP6:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0 914 // CHECK-NEXT: store i32 3, ptr [[TMP6]], align 4 915 // CHECK-NEXT: [[TMP7:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1 916 // CHECK-NEXT: store i32 1, ptr [[TMP7]], align 4 917 // CHECK-NEXT: [[TMP8:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2 918 // CHECK-NEXT: store ptr [[TMP4]], ptr [[TMP8]], align 8 919 // CHECK-NEXT: [[TMP9:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3 920 // CHECK-NEXT: store ptr [[TMP5]], ptr [[TMP9]], align 8 921 // CHECK-NEXT: [[TMP10:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4 922 // CHECK-NEXT: store ptr @.offload_sizes.23, ptr [[TMP10]], align 8 923 // CHECK-NEXT: [[TMP11:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5 924 // CHECK-NEXT: store ptr @.offload_maptypes.24, ptr [[TMP11]], align 8 925 // CHECK-NEXT: [[TMP12:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6 926 // CHECK-NEXT: store ptr null, ptr [[TMP12]], align 8 927 // CHECK-NEXT: [[TMP13:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7 928 // CHECK-NEXT: store ptr null, ptr [[TMP13]], align 8 929 // CHECK-NEXT: [[TMP14:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8 930 // CHECK-NEXT: store i64 0, ptr [[TMP14]], align 8 931 // CHECK-NEXT: [[TMP15:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9 932 // CHECK-NEXT: store i64 0, ptr [[TMP15]], align 8 933 // CHECK-NEXT: [[TMP16:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10 934 // CHECK-NEXT: store [3 x i32] [i32 -1, i32 0, i32 0], ptr [[TMP16]], align 4 935 // CHECK-NEXT: [[TMP17:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11 936 // CHECK-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP17]], align 4 937 // CHECK-NEXT: [[TMP18:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12 938 // CHECK-NEXT: store i32 0, ptr [[TMP18]], align 4 939 // CHECK-NEXT: [[TMP19:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 -1, i32 -1, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIPiET_S1__l123.region_id, ptr [[KERNEL_ARGS]]) 940 // CHECK-NEXT: [[TMP20:%.*]] = icmp ne i32 [[TMP19]], 0 941 // CHECK-NEXT: br i1 [[TMP20]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]] 942 // CHECK: omp_offload.failed: 943 // CHECK-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIPiET_S1__l123(ptr [[K]]) #[[ATTR5]] 944 // CHECK-NEXT: br label [[OMP_OFFLOAD_CONT]] 945 // CHECK: omp_offload.cont: 946 // CHECK-NEXT: [[TMP21:%.*]] = load ptr, ptr [[Z]], align 8 947 // CHECK-NEXT: store ptr [[TMP21]], ptr [[TMP]], align 8 948 // CHECK-NEXT: [[TMP22:%.*]] = load ptr, ptr [[TMP]], align 8 949 // CHECK-NEXT: [[TMP23:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS1]], i32 0, i32 0 950 // CHECK-NEXT: store ptr [[TMP22]], ptr [[TMP23]], align 8 951 // CHECK-NEXT: [[TMP24:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS2]], i32 0, i32 0 952 // CHECK-NEXT: store ptr [[TMP22]], ptr [[TMP24]], align 8 953 // CHECK-NEXT: [[TMP25:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS3]], i64 0, i64 0 954 // CHECK-NEXT: store ptr null, ptr [[TMP25]], align 8 955 // CHECK-NEXT: [[TMP26:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS1]], i32 0, i32 0 956 // CHECK-NEXT: [[TMP27:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS2]], i32 0, i32 0 957 // CHECK-NEXT: [[TMP28:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS4]], i32 0, i32 0 958 // CHECK-NEXT: store i32 3, ptr [[TMP28]], align 4 959 // CHECK-NEXT: [[TMP29:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS4]], i32 0, i32 1 960 // CHECK-NEXT: store i32 1, ptr [[TMP29]], align 4 961 // CHECK-NEXT: [[TMP30:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS4]], i32 0, i32 2 962 // CHECK-NEXT: store ptr [[TMP26]], ptr [[TMP30]], align 8 963 // CHECK-NEXT: [[TMP31:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS4]], i32 0, i32 3 964 // CHECK-NEXT: store ptr [[TMP27]], ptr [[TMP31]], align 8 965 // CHECK-NEXT: [[TMP32:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS4]], i32 0, i32 4 966 // CHECK-NEXT: store ptr @.offload_sizes.25, ptr [[TMP32]], align 8 967 // CHECK-NEXT: [[TMP33:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS4]], i32 0, i32 5 968 // CHECK-NEXT: store ptr @.offload_maptypes.26, ptr [[TMP33]], align 8 969 // CHECK-NEXT: [[TMP34:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS4]], i32 0, i32 6 970 // CHECK-NEXT: store ptr null, ptr [[TMP34]], align 8 971 // CHECK-NEXT: [[TMP35:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS4]], i32 0, i32 7 972 // CHECK-NEXT: store ptr null, ptr [[TMP35]], align 8 973 // CHECK-NEXT: [[TMP36:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS4]], i32 0, i32 8 974 // CHECK-NEXT: store i64 0, ptr [[TMP36]], align 8 975 // CHECK-NEXT: [[TMP37:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS4]], i32 0, i32 9 976 // CHECK-NEXT: store i64 0, ptr [[TMP37]], align 8 977 // CHECK-NEXT: [[TMP38:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS4]], i32 0, i32 10 978 // CHECK-NEXT: store [3 x i32] [i32 -1, i32 0, i32 0], ptr [[TMP38]], align 4 979 // CHECK-NEXT: [[TMP39:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS4]], i32 0, i32 11 980 // CHECK-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP39]], align 4 981 // CHECK-NEXT: [[TMP40:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS4]], i32 0, i32 12 982 // CHECK-NEXT: store i32 0, ptr [[TMP40]], align 4 983 // CHECK-NEXT: [[TMP41:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 -1, i32 -1, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIPiET_S1__l125.region_id, ptr [[KERNEL_ARGS4]]) 984 // CHECK-NEXT: [[TMP42:%.*]] = icmp ne i32 [[TMP41]], 0 985 // CHECK-NEXT: br i1 [[TMP42]], label [[OMP_OFFLOAD_FAILED5:%.*]], label [[OMP_OFFLOAD_CONT6:%.*]] 986 // CHECK: omp_offload.failed5: 987 // CHECK-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIPiET_S1__l125(ptr [[TMP22]]) #[[ATTR5]] 988 // CHECK-NEXT: br label [[OMP_OFFLOAD_CONT6]] 989 // CHECK: omp_offload.cont6: 990 // CHECK-NEXT: [[TMP43:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS7]], i32 0, i32 0 991 // CHECK-NEXT: store ptr [[AA]], ptr [[TMP43]], align 8 992 // CHECK-NEXT: [[TMP44:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS8]], i32 0, i32 0 993 // CHECK-NEXT: store ptr [[AA]], ptr [[TMP44]], align 8 994 // CHECK-NEXT: [[TMP45:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS9]], i64 0, i64 0 995 // CHECK-NEXT: store ptr null, ptr [[TMP45]], align 8 996 // CHECK-NEXT: [[TMP46:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS7]], i32 0, i32 0 997 // CHECK-NEXT: [[TMP47:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS8]], i32 0, i32 0 998 // CHECK-NEXT: [[TMP48:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS10]], i32 0, i32 0 999 // CHECK-NEXT: store i32 3, ptr [[TMP48]], align 4 1000 // CHECK-NEXT: [[TMP49:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS10]], i32 0, i32 1 1001 // CHECK-NEXT: store i32 1, ptr [[TMP49]], align 4 1002 // CHECK-NEXT: [[TMP50:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS10]], i32 0, i32 2 1003 // CHECK-NEXT: store ptr [[TMP46]], ptr [[TMP50]], align 8 1004 // CHECK-NEXT: [[TMP51:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS10]], i32 0, i32 3 1005 // CHECK-NEXT: store ptr [[TMP47]], ptr [[TMP51]], align 8 1006 // CHECK-NEXT: [[TMP52:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS10]], i32 0, i32 4 1007 // CHECK-NEXT: store ptr @.offload_sizes.27, ptr [[TMP52]], align 8 1008 // CHECK-NEXT: [[TMP53:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS10]], i32 0, i32 5 1009 // CHECK-NEXT: store ptr @.offload_maptypes.28, ptr [[TMP53]], align 8 1010 // CHECK-NEXT: [[TMP54:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS10]], i32 0, i32 6 1011 // CHECK-NEXT: store ptr null, ptr [[TMP54]], align 8 1012 // CHECK-NEXT: [[TMP55:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS10]], i32 0, i32 7 1013 // CHECK-NEXT: store ptr null, ptr [[TMP55]], align 8 1014 // CHECK-NEXT: [[TMP56:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS10]], i32 0, i32 8 1015 // CHECK-NEXT: store i64 0, ptr [[TMP56]], align 8 1016 // CHECK-NEXT: [[TMP57:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS10]], i32 0, i32 9 1017 // CHECK-NEXT: store i64 0, ptr [[TMP57]], align 8 1018 // CHECK-NEXT: [[TMP58:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS10]], i32 0, i32 10 1019 // CHECK-NEXT: store [3 x i32] [i32 -1, i32 0, i32 0], ptr [[TMP58]], align 4 1020 // CHECK-NEXT: [[TMP59:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS10]], i32 0, i32 11 1021 // CHECK-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP59]], align 4 1022 // CHECK-NEXT: [[TMP60:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS10]], i32 0, i32 12 1023 // CHECK-NEXT: store i32 0, ptr [[TMP60]], align 4 1024 // CHECK-NEXT: [[TMP61:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 -1, i32 -1, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIPiET_S1__l127.region_id, ptr [[KERNEL_ARGS10]]) 1025 // CHECK-NEXT: [[TMP62:%.*]] = icmp ne i32 [[TMP61]], 0 1026 // CHECK-NEXT: br i1 [[TMP62]], label [[OMP_OFFLOAD_FAILED11:%.*]], label [[OMP_OFFLOAD_CONT12:%.*]] 1027 // CHECK: omp_offload.failed11: 1028 // CHECK-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIPiET_S1__l127(ptr [[AA]]) #[[ATTR5]] 1029 // CHECK-NEXT: br label [[OMP_OFFLOAD_CONT12]] 1030 // CHECK: omp_offload.cont12: 1031 // CHECK-NEXT: [[TMP63:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS13]], i32 0, i32 0 1032 // CHECK-NEXT: store ptr [[H]], ptr [[TMP63]], align 8 1033 // CHECK-NEXT: [[TMP64:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS14]], i32 0, i32 0 1034 // CHECK-NEXT: store ptr [[H]], ptr [[TMP64]], align 8 1035 // CHECK-NEXT: [[TMP65:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS15]], i64 0, i64 0 1036 // CHECK-NEXT: store ptr null, ptr [[TMP65]], align 8 1037 // CHECK-NEXT: [[TMP66:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS13]], i32 0, i32 0 1038 // CHECK-NEXT: [[TMP67:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS14]], i32 0, i32 0 1039 // CHECK-NEXT: [[TMP68:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS16]], i32 0, i32 0 1040 // CHECK-NEXT: store i32 3, ptr [[TMP68]], align 4 1041 // CHECK-NEXT: [[TMP69:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS16]], i32 0, i32 1 1042 // CHECK-NEXT: store i32 1, ptr [[TMP69]], align 4 1043 // CHECK-NEXT: [[TMP70:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS16]], i32 0, i32 2 1044 // CHECK-NEXT: store ptr [[TMP66]], ptr [[TMP70]], align 8 1045 // CHECK-NEXT: [[TMP71:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS16]], i32 0, i32 3 1046 // CHECK-NEXT: store ptr [[TMP67]], ptr [[TMP71]], align 8 1047 // CHECK-NEXT: [[TMP72:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS16]], i32 0, i32 4 1048 // CHECK-NEXT: store ptr @.offload_sizes.29, ptr [[TMP72]], align 8 1049 // CHECK-NEXT: [[TMP73:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS16]], i32 0, i32 5 1050 // CHECK-NEXT: store ptr @.offload_maptypes.30, ptr [[TMP73]], align 8 1051 // CHECK-NEXT: [[TMP74:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS16]], i32 0, i32 6 1052 // CHECK-NEXT: store ptr null, ptr [[TMP74]], align 8 1053 // CHECK-NEXT: [[TMP75:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS16]], i32 0, i32 7 1054 // CHECK-NEXT: store ptr null, ptr [[TMP75]], align 8 1055 // CHECK-NEXT: [[TMP76:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS16]], i32 0, i32 8 1056 // CHECK-NEXT: store i64 0, ptr [[TMP76]], align 8 1057 // CHECK-NEXT: [[TMP77:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS16]], i32 0, i32 9 1058 // CHECK-NEXT: store i64 0, ptr [[TMP77]], align 8 1059 // CHECK-NEXT: [[TMP78:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS16]], i32 0, i32 10 1060 // CHECK-NEXT: store [3 x i32] [i32 -1, i32 0, i32 0], ptr [[TMP78]], align 4 1061 // CHECK-NEXT: [[TMP79:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS16]], i32 0, i32 11 1062 // CHECK-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP79]], align 4 1063 // CHECK-NEXT: [[TMP80:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS16]], i32 0, i32 12 1064 // CHECK-NEXT: store i32 0, ptr [[TMP80]], align 4 1065 // CHECK-NEXT: [[TMP81:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 -1, i32 -1, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIPiET_S1__l129.region_id, ptr [[KERNEL_ARGS16]]) 1066 // CHECK-NEXT: [[TMP82:%.*]] = icmp ne i32 [[TMP81]], 0 1067 // CHECK-NEXT: br i1 [[TMP82]], label [[OMP_OFFLOAD_FAILED17:%.*]], label [[OMP_OFFLOAD_CONT18:%.*]] 1068 // CHECK: omp_offload.failed17: 1069 // CHECK-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIPiET_S1__l129(ptr [[H]]) #[[ATTR5]] 1070 // CHECK-NEXT: br label [[OMP_OFFLOAD_CONT18]] 1071 // CHECK: omp_offload.cont18: 1072 // CHECK-NEXT: ret ptr null 1073 // 1074 // 1075 // CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_S0__l123 1076 // CHECK-SAME: (ptr noundef nonnull align 8 dereferenceable(8) [[K:%.*]]) #[[ATTR4]] { 1077 // CHECK-NEXT: entry: 1078 // CHECK-NEXT: [[K_ADDR:%.*]] = alloca ptr, align 8 1079 // CHECK-NEXT: store ptr [[K]], ptr [[K_ADDR]], align 8 1080 // CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[K_ADDR]], align 8 1081 // CHECK-NEXT: [[TMP1:%.*]] = load ptr, ptr [[TMP0]], align 8 1082 // CHECK-NEXT: [[INCDEC_PTR:%.*]] = getelementptr inbounds nuw i32, ptr [[TMP1]], i32 1 1083 // CHECK-NEXT: store ptr [[INCDEC_PTR]], ptr [[TMP0]], align 8 1084 // CHECK-NEXT: ret void 1085 // 1086 // 1087 // CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_S0__l125 1088 // CHECK-SAME: (ptr noundef nonnull align 8 dereferenceable(8) [[Z:%.*]]) #[[ATTR4]] { 1089 // CHECK-NEXT: entry: 1090 // CHECK-NEXT: [[Z_ADDR:%.*]] = alloca ptr, align 8 1091 // CHECK-NEXT: [[TMP:%.*]] = alloca ptr, align 8 1092 // CHECK-NEXT: store ptr [[Z]], ptr [[Z_ADDR]], align 8 1093 // CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[Z_ADDR]], align 8 1094 // CHECK-NEXT: store ptr [[TMP0]], ptr [[TMP]], align 8 1095 // CHECK-NEXT: [[TMP1:%.*]] = load ptr, ptr [[TMP]], align 8 1096 // CHECK-NEXT: [[TMP2:%.*]] = load ptr, ptr [[TMP1]], align 8 1097 // CHECK-NEXT: [[INCDEC_PTR:%.*]] = getelementptr inbounds nuw i32, ptr [[TMP2]], i32 1 1098 // CHECK-NEXT: store ptr [[INCDEC_PTR]], ptr [[TMP1]], align 8 1099 // CHECK-NEXT: ret void 1100 // 1101 // 1102 // CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_S0__l127 1103 // CHECK-SAME: (ptr noundef nonnull align 4 dereferenceable(40) [[AA:%.*]]) #[[ATTR4]] { 1104 // CHECK-NEXT: entry: 1105 // CHECK-NEXT: [[AA_ADDR:%.*]] = alloca ptr, align 8 1106 // CHECK-NEXT: [[A:%.*]] = alloca i32, align 4 1107 // CHECK-NEXT: store ptr [[AA]], ptr [[AA_ADDR]], align 8 1108 // CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[AA_ADDR]], align 8 1109 // CHECK-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x i32], ptr [[TMP0]], i64 0, i64 0 1110 // CHECK-NEXT: [[TMP1:%.*]] = load i32, ptr [[ARRAYIDX]], align 4 1111 // CHECK-NEXT: store i32 [[TMP1]], ptr [[A]], align 4 1112 // CHECK-NEXT: ret void 1113 // 1114 // 1115 // CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIiET_S0__l129 1116 // CHECK-SAME: (ptr noundef nonnull align 4 dereferenceable(40) [[H:%.*]]) #[[ATTR4]] { 1117 // CHECK-NEXT: entry: 1118 // CHECK-NEXT: [[H_ADDR:%.*]] = alloca ptr, align 8 1119 // CHECK-NEXT: [[A:%.*]] = alloca i32, align 4 1120 // CHECK-NEXT: store ptr [[H]], ptr [[H_ADDR]], align 8 1121 // CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[H_ADDR]], align 8 1122 // CHECK-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x %struct.S6], ptr [[TMP0]], i64 0, i64 0 1123 // CHECK-NEXT: [[A1:%.*]] = getelementptr inbounds nuw [[STRUCT_S6:%.*]], ptr [[ARRAYIDX]], i32 0, i32 0 1124 // CHECK-NEXT: [[TMP1:%.*]] = load i32, ptr [[A1]], align 4 1125 // CHECK-NEXT: store i32 [[TMP1]], ptr [[A]], align 4 1126 // CHECK-NEXT: ret void 1127 // 1128 // 1129 // CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIPiET_S1__l123 1130 // CHECK-SAME: (ptr noundef nonnull align 8 dereferenceable(8) [[K:%.*]]) #[[ATTR4]] { 1131 // CHECK-NEXT: entry: 1132 // CHECK-NEXT: [[K_ADDR:%.*]] = alloca ptr, align 8 1133 // CHECK-NEXT: store ptr [[K]], ptr [[K_ADDR]], align 8 1134 // CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[K_ADDR]], align 8 1135 // CHECK-NEXT: [[TMP1:%.*]] = load ptr, ptr [[TMP0]], align 8 1136 // CHECK-NEXT: [[INCDEC_PTR:%.*]] = getelementptr inbounds nuw ptr, ptr [[TMP1]], i32 1 1137 // CHECK-NEXT: store ptr [[INCDEC_PTR]], ptr [[TMP0]], align 8 1138 // CHECK-NEXT: ret void 1139 // 1140 // 1141 // CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIPiET_S1__l125 1142 // CHECK-SAME: (ptr noundef nonnull align 8 dereferenceable(8) [[Z:%.*]]) #[[ATTR4]] { 1143 // CHECK-NEXT: entry: 1144 // CHECK-NEXT: [[Z_ADDR:%.*]] = alloca ptr, align 8 1145 // CHECK-NEXT: [[TMP:%.*]] = alloca ptr, align 8 1146 // CHECK-NEXT: store ptr [[Z]], ptr [[Z_ADDR]], align 8 1147 // CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[Z_ADDR]], align 8 1148 // CHECK-NEXT: store ptr [[TMP0]], ptr [[TMP]], align 8 1149 // CHECK-NEXT: [[TMP1:%.*]] = load ptr, ptr [[TMP]], align 8 1150 // CHECK-NEXT: [[TMP2:%.*]] = load ptr, ptr [[TMP1]], align 8 1151 // CHECK-NEXT: [[INCDEC_PTR:%.*]] = getelementptr inbounds nuw ptr, ptr [[TMP2]], i32 1 1152 // CHECK-NEXT: store ptr [[INCDEC_PTR]], ptr [[TMP1]], align 8 1153 // CHECK-NEXT: ret void 1154 // 1155 // 1156 // CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIPiET_S1__l127 1157 // CHECK-SAME: (ptr noundef nonnull align 8 dereferenceable(80) [[AA:%.*]]) #[[ATTR4]] { 1158 // CHECK-NEXT: entry: 1159 // CHECK-NEXT: [[AA_ADDR:%.*]] = alloca ptr, align 8 1160 // CHECK-NEXT: [[A:%.*]] = alloca ptr, align 8 1161 // CHECK-NEXT: store ptr [[AA]], ptr [[AA_ADDR]], align 8 1162 // CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[AA_ADDR]], align 8 1163 // CHECK-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x ptr], ptr [[TMP0]], i64 0, i64 0 1164 // CHECK-NEXT: [[TMP1:%.*]] = load ptr, ptr [[ARRAYIDX]], align 8 1165 // CHECK-NEXT: store ptr [[TMP1]], ptr [[A]], align 8 1166 // CHECK-NEXT: ret void 1167 // 1168 // 1169 // CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z5tmainIPiET_S1__l129 1170 // CHECK-SAME: (ptr noundef nonnull align 4 dereferenceable(40) [[H:%.*]]) #[[ATTR4]] { 1171 // CHECK-NEXT: entry: 1172 // CHECK-NEXT: [[H_ADDR:%.*]] = alloca ptr, align 8 1173 // CHECK-NEXT: [[A:%.*]] = alloca i32, align 4 1174 // CHECK-NEXT: store ptr [[H]], ptr [[H_ADDR]], align 8 1175 // CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[H_ADDR]], align 8 1176 // CHECK-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x %struct.S6], ptr [[TMP0]], i64 0, i64 0 1177 // CHECK-NEXT: [[A1:%.*]] = getelementptr inbounds nuw [[STRUCT_S6:%.*]], ptr [[ARRAYIDX]], i32 0, i32 0 1178 // CHECK-NEXT: [[TMP1:%.*]] = load i32, ptr [[A1]], align 4 1179 // CHECK-NEXT: store i32 [[TMP1]], ptr [[A]], align 4 1180 // CHECK-NEXT: ret void 1181 // 1182 // 1183 // CHECK-LABEL: define {{[^@]+}}@_Z12use_templatev 1184 // CHECK-SAME: () #[[ATTR1]] { 1185 // CHECK-NEXT: entry: 1186 // CHECK-NEXT: [[AKERN:%.*]] = alloca [[STRUCT_SOMEKERNEL:%.*]], align 4 1187 // CHECK-NEXT: call void @_ZN10SomeKernelC1Ev(ptr noundef nonnull align 4 dereferenceable(8) [[AKERN]]) 1188 // CHECK-NEXT: call void @_ZN10SomeKernel5applyILj32EEEvv(ptr noundef nonnull align 4 dereferenceable(8) [[AKERN]]) 1189 // CHECK-NEXT: call void @_ZN10SomeKernelD1Ev(ptr noundef nonnull align 4 dereferenceable(8) [[AKERN]]) #[[ATTR5]] 1190 // CHECK-NEXT: ret void 1191 // 1192 // 1193 // CHECK-LABEL: define {{[^@]+}}@_ZN10SomeKernel5applyILj32EEEvv 1194 // CHECK-SAME: (ptr noundef nonnull align 4 dereferenceable(8) [[THIS:%.*]]) #[[ATTR1]] comdat { 1195 // CHECK-NEXT: entry: 1196 // CHECK-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 8 1197 // CHECK-NEXT: [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 4 1198 // CHECK-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [3 x ptr], align 8 1199 // CHECK-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [3 x ptr], align 8 1200 // CHECK-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [3 x ptr], align 8 1201 // CHECK-NEXT: [[DOTOFFLOAD_SIZES:%.*]] = alloca [3 x i64], align 8 1202 // CHECK-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8 1203 // CHECK-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 8 1204 // CHECK-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8 1205 // CHECK-NEXT: [[TARGETDEV:%.*]] = getelementptr inbounds nuw [[STRUCT_SOMEKERNEL:%.*]], ptr [[THIS1]], i32 0, i32 0 1206 // CHECK-NEXT: [[TMP0:%.*]] = load i32, ptr [[TARGETDEV]], align 4 1207 // CHECK-NEXT: store i32 [[TMP0]], ptr [[DOTCAPTURE_EXPR_]], align 4 1208 // CHECK-NEXT: [[DEVPTR:%.*]] = getelementptr inbounds nuw [[STRUCT_SOMEKERNEL]], ptr [[THIS1]], i32 0, i32 1 1209 // CHECK-NEXT: [[TARGETDEV2:%.*]] = getelementptr inbounds nuw [[STRUCT_SOMEKERNEL]], ptr [[THIS1]], i32 0, i32 0 1210 // CHECK-NEXT: [[TMP1:%.*]] = getelementptr float, ptr [[DEVPTR]], i32 1 1211 // CHECK-NEXT: [[TMP2:%.*]] = ptrtoint ptr [[TMP1]] to i64 1212 // CHECK-NEXT: [[TMP3:%.*]] = ptrtoint ptr [[TARGETDEV2]] to i64 1213 // CHECK-NEXT: [[TMP4:%.*]] = sub i64 [[TMP2]], [[TMP3]] 1214 // CHECK-NEXT: [[TMP5:%.*]] = sdiv exact i64 [[TMP4]], ptrtoint (ptr getelementptr (i8, ptr null, i32 1) to i64) 1215 // CHECK-NEXT: call void @llvm.memcpy.p0.p0.i64(ptr align 8 [[DOTOFFLOAD_SIZES]], ptr align 8 @.offload_sizes.31, i64 24, i1 false) 1216 // CHECK-NEXT: [[TMP6:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 1217 // CHECK-NEXT: store ptr [[THIS1]], ptr [[TMP6]], align 8 1218 // CHECK-NEXT: [[TMP7:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0 1219 // CHECK-NEXT: store ptr [[TARGETDEV2]], ptr [[TMP7]], align 8 1220 // CHECK-NEXT: [[TMP8:%.*]] = getelementptr inbounds [3 x i64], ptr [[DOTOFFLOAD_SIZES]], i32 0, i32 0 1221 // CHECK-NEXT: store i64 [[TMP5]], ptr [[TMP8]], align 8 1222 // CHECK-NEXT: [[TMP9:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 0 1223 // CHECK-NEXT: store ptr null, ptr [[TMP9]], align 8 1224 // CHECK-NEXT: [[TMP10:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 1 1225 // CHECK-NEXT: store ptr [[THIS1]], ptr [[TMP10]], align 8 1226 // CHECK-NEXT: [[TMP11:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 1 1227 // CHECK-NEXT: store ptr [[DEVPTR]], ptr [[TMP11]], align 8 1228 // CHECK-NEXT: [[TMP12:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 1 1229 // CHECK-NEXT: store ptr null, ptr [[TMP12]], align 8 1230 // CHECK-NEXT: [[TMP13:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 2 1231 // CHECK-NEXT: store ptr [[THIS1]], ptr [[TMP13]], align 8 1232 // CHECK-NEXT: [[TMP14:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 2 1233 // CHECK-NEXT: store ptr [[TARGETDEV2]], ptr [[TMP14]], align 8 1234 // CHECK-NEXT: [[TMP15:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 2 1235 // CHECK-NEXT: store ptr null, ptr [[TMP15]], align 8 1236 // CHECK-NEXT: [[TMP16:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 1237 // CHECK-NEXT: [[TMP17:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0 1238 // CHECK-NEXT: [[TMP18:%.*]] = getelementptr inbounds [3 x i64], ptr [[DOTOFFLOAD_SIZES]], i32 0, i32 0 1239 // CHECK-NEXT: [[TMP19:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4 1240 // CHECK-NEXT: [[TMP20:%.*]] = sext i32 [[TMP19]] to i64 1241 // CHECK-NEXT: [[TMP21:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0 1242 // CHECK-NEXT: store i32 3, ptr [[TMP21]], align 4 1243 // CHECK-NEXT: [[TMP22:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1 1244 // CHECK-NEXT: store i32 3, ptr [[TMP22]], align 4 1245 // CHECK-NEXT: [[TMP23:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2 1246 // CHECK-NEXT: store ptr [[TMP16]], ptr [[TMP23]], align 8 1247 // CHECK-NEXT: [[TMP24:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3 1248 // CHECK-NEXT: store ptr [[TMP17]], ptr [[TMP24]], align 8 1249 // CHECK-NEXT: [[TMP25:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4 1250 // CHECK-NEXT: store ptr [[TMP18]], ptr [[TMP25]], align 8 1251 // CHECK-NEXT: [[TMP26:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5 1252 // CHECK-NEXT: store ptr @.offload_maptypes.32, ptr [[TMP26]], align 8 1253 // CHECK-NEXT: [[TMP27:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6 1254 // CHECK-NEXT: store ptr null, ptr [[TMP27]], align 8 1255 // CHECK-NEXT: [[TMP28:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7 1256 // CHECK-NEXT: store ptr null, ptr [[TMP28]], align 8 1257 // CHECK-NEXT: [[TMP29:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8 1258 // CHECK-NEXT: store i64 0, ptr [[TMP29]], align 8 1259 // CHECK-NEXT: [[TMP30:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9 1260 // CHECK-NEXT: store i64 0, ptr [[TMP30]], align 8 1261 // CHECK-NEXT: [[TMP31:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10 1262 // CHECK-NEXT: store [3 x i32] [i32 -1, i32 0, i32 0], ptr [[TMP31]], align 4 1263 // CHECK-NEXT: [[TMP32:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11 1264 // CHECK-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP32]], align 4 1265 // CHECK-NEXT: [[TMP33:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12 1266 // CHECK-NEXT: store i32 0, ptr [[TMP33]], align 4 1267 // CHECK-NEXT: [[TMP34:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 [[TMP20]], i32 -1, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN10SomeKernel5applyILj32EEEvv_l168.region_id, ptr [[KERNEL_ARGS]]) 1268 // CHECK-NEXT: [[TMP35:%.*]] = icmp ne i32 [[TMP34]], 0 1269 // CHECK-NEXT: br i1 [[TMP35]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]] 1270 // CHECK: omp_offload.failed: 1271 // CHECK-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN10SomeKernel5applyILj32EEEvv_l168(ptr [[THIS1]]) #[[ATTR5]] 1272 // CHECK-NEXT: br label [[OMP_OFFLOAD_CONT]] 1273 // CHECK: omp_offload.cont: 1274 // CHECK-NEXT: ret void 1275 // 1276 // 1277 // CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN10SomeKernel5applyILj32EEEvv_l168 1278 // CHECK-SAME: (ptr noundef [[THIS:%.*]]) #[[ATTR4]] { 1279 // CHECK-NEXT: entry: 1280 // CHECK-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 8 1281 // CHECK-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 8 1282 // CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8 1283 // CHECK-NEXT: [[DEVPTR:%.*]] = getelementptr inbounds nuw [[STRUCT_SOMEKERNEL:%.*]], ptr [[TMP0]], i32 0, i32 1 1284 // CHECK-NEXT: [[TMP1:%.*]] = load float, ptr [[DEVPTR]], align 4 1285 // CHECK-NEXT: [[INC:%.*]] = fadd float [[TMP1]], 1.000000e+00 1286 // CHECK-NEXT: store float [[INC]], ptr [[DEVPTR]], align 4 1287 // CHECK-NEXT: [[TARGETDEV:%.*]] = getelementptr inbounds nuw [[STRUCT_SOMEKERNEL]], ptr [[TMP0]], i32 0, i32 0 1288 // CHECK-NEXT: [[TMP2:%.*]] = load i32, ptr [[TARGETDEV]], align 4 1289 // CHECK-NEXT: [[INC1:%.*]] = add nsw i32 [[TMP2]], 1 1290 // CHECK-NEXT: store i32 [[INC1]], ptr [[TARGETDEV]], align 4 1291 // CHECK-NEXT: ret void 1292 // 1293 // 1294 // CHECK-LABEL: define {{[^@]+}}@_GLOBAL__sub_I_target_has_device_addr_codegen.cpp 1295 // CHECK-SAME: () #[[ATTR0]] { 1296 // CHECK-NEXT: entry: 1297 // CHECK-NEXT: call void @__cxx_global_var_init() 1298 // CHECK-NEXT: call void @__cxx_global_var_init.1() 1299 // CHECK-NEXT: call void @__cxx_global_var_init.2() 1300 // CHECK-NEXT: call void @__cxx_global_var_init.3() 1301 // CHECK-NEXT: ret void 1302 // 1303 // 1304 // CHECK-LABEL: define {{[^@]+}}@__tls_init 1305 // CHECK-SAME: () #[[ATTR0]] { 1306 // CHECK-NEXT: entry: 1307 // CHECK-NEXT: [[TMP0:%.*]] = load i8, ptr @__tls_guard, align 1 1308 // CHECK-NEXT: [[GUARD_UNINITIALIZED:%.*]] = icmp eq i8 [[TMP0]], 0 1309 // CHECK-NEXT: br i1 [[GUARD_UNINITIALIZED]], label [[INIT:%.*]], label [[EXIT:%.*]], !prof [[PROF18:![0-9]+]] 1310 // CHECK: init: 1311 // CHECK-NEXT: store i8 1, ptr @__tls_guard, align 1 1312 // CHECK-NEXT: call void @__cxx_global_var_init.4() 1313 // CHECK-NEXT: br label [[EXIT]] 1314 // CHECK: exit: 1315 // CHECK-NEXT: ret void 1316 // 1317 // 1318 // CHECK-LABEL: define {{[^@]+}}@_ZTW1h 1319 // CHECK-SAME: () #[[ATTR0]] comdat { 1320 // CHECK-NEXT: call void @_ZTH1h() 1321 // CHECK-NEXT: [[TMP1:%.*]] = call align 4 ptr @llvm.threadlocal.address.p0(ptr align 4 @h) 1322 // CHECK-NEXT: ret ptr [[TMP1]] 1323 // 1324 // 1325 // SIMD-ONLY0-LABEL: define {{[^@]+}}@__cxx_global_var_init 1326 // SIMD-ONLY0-SAME: () #[[ATTR0:[0-9]+]] { 1327 // SIMD-ONLY0-NEXT: entry: 1328 // SIMD-ONLY0-NEXT: call void @_ZN2S2C1Ev(ptr noundef nonnull align 4 dereferenceable(4) @_ZL1b) 1329 // SIMD-ONLY0-NEXT: ret void 1330 // 1331 // 1332 // SIMD-ONLY0-LABEL: define {{[^@]+}}@_ZN2S2C1Ev 1333 // SIMD-ONLY0-SAME: (ptr noundef nonnull align 4 dereferenceable(4) [[THIS:%.*]]) unnamed_addr #[[ATTR1:[0-9]+]] comdat { 1334 // SIMD-ONLY0-NEXT: entry: 1335 // SIMD-ONLY0-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 8 1336 // SIMD-ONLY0-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 8 1337 // SIMD-ONLY0-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8 1338 // SIMD-ONLY0-NEXT: call void @_ZN2S2C2Ev(ptr noundef nonnull align 4 dereferenceable(4) [[THIS1]]) 1339 // SIMD-ONLY0-NEXT: ret void 1340 // 1341 // 1342 // SIMD-ONLY0-LABEL: define {{[^@]+}}@__cxx_global_var_init.1 1343 // SIMD-ONLY0-SAME: () #[[ATTR0]] { 1344 // SIMD-ONLY0-NEXT: entry: 1345 // SIMD-ONLY0-NEXT: br label [[ARRAYCTOR_LOOP:%.*]] 1346 // SIMD-ONLY0: arrayctor.loop: 1347 // SIMD-ONLY0-NEXT: [[ARRAYCTOR_CUR:%.*]] = phi ptr [ @_ZL2ba, [[ENTRY:%.*]] ], [ [[ARRAYCTOR_NEXT:%.*]], [[ARRAYCTOR_LOOP]] ] 1348 // SIMD-ONLY0-NEXT: call void @_ZN2S2C1Ev(ptr noundef nonnull align 4 dereferenceable(4) [[ARRAYCTOR_CUR]]) 1349 // SIMD-ONLY0-NEXT: [[ARRAYCTOR_NEXT]] = getelementptr inbounds [[CLASS_S2:%.*]], ptr [[ARRAYCTOR_CUR]], i64 1 1350 // SIMD-ONLY0-NEXT: [[ARRAYCTOR_DONE:%.*]] = icmp eq ptr [[ARRAYCTOR_NEXT]], getelementptr inbounds ([[CLASS_S2]], ptr @_ZL2ba, i64 5) 1351 // SIMD-ONLY0-NEXT: br i1 [[ARRAYCTOR_DONE]], label [[ARRAYCTOR_CONT:%.*]], label [[ARRAYCTOR_LOOP]] 1352 // SIMD-ONLY0: arrayctor.cont: 1353 // SIMD-ONLY0-NEXT: ret void 1354 // 1355 // 1356 // SIMD-ONLY0-LABEL: define {{[^@]+}}@__cxx_global_var_init.2 1357 // SIMD-ONLY0-SAME: () #[[ATTR0]] { 1358 // SIMD-ONLY0-NEXT: entry: 1359 // SIMD-ONLY0-NEXT: call void @_ZN2S3C1Ev(ptr noundef nonnull align 4 dereferenceable(4) @_ZL1c) 1360 // SIMD-ONLY0-NEXT: ret void 1361 // 1362 // 1363 // SIMD-ONLY0-LABEL: define {{[^@]+}}@_ZN2S3C1Ev 1364 // SIMD-ONLY0-SAME: (ptr noundef nonnull align 4 dereferenceable(4) [[THIS:%.*]]) unnamed_addr #[[ATTR1]] comdat { 1365 // SIMD-ONLY0-NEXT: entry: 1366 // SIMD-ONLY0-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 8 1367 // SIMD-ONLY0-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 8 1368 // SIMD-ONLY0-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8 1369 // SIMD-ONLY0-NEXT: call void @_ZN2S3C2Ev(ptr noundef nonnull align 4 dereferenceable(4) [[THIS1]]) 1370 // SIMD-ONLY0-NEXT: ret void 1371 // 1372 // 1373 // SIMD-ONLY0-LABEL: define {{[^@]+}}@__cxx_global_var_init.3 1374 // SIMD-ONLY0-SAME: () #[[ATTR0]] { 1375 // SIMD-ONLY0-NEXT: entry: 1376 // SIMD-ONLY0-NEXT: br label [[ARRAYCTOR_LOOP:%.*]] 1377 // SIMD-ONLY0: arrayctor.loop: 1378 // SIMD-ONLY0-NEXT: [[ARRAYCTOR_CUR:%.*]] = phi ptr [ @_ZL2ca, [[ENTRY:%.*]] ], [ [[ARRAYCTOR_NEXT:%.*]], [[ARRAYCTOR_LOOP]] ] 1379 // SIMD-ONLY0-NEXT: call void @_ZN2S3C1Ev(ptr noundef nonnull align 4 dereferenceable(4) [[ARRAYCTOR_CUR]]) 1380 // SIMD-ONLY0-NEXT: [[ARRAYCTOR_NEXT]] = getelementptr inbounds [[CLASS_S3:%.*]], ptr [[ARRAYCTOR_CUR]], i64 1 1381 // SIMD-ONLY0-NEXT: [[ARRAYCTOR_DONE:%.*]] = icmp eq ptr [[ARRAYCTOR_NEXT]], getelementptr inbounds ([[CLASS_S3]], ptr @_ZL2ca, i64 5) 1382 // SIMD-ONLY0-NEXT: br i1 [[ARRAYCTOR_DONE]], label [[ARRAYCTOR_CONT:%.*]], label [[ARRAYCTOR_LOOP]] 1383 // SIMD-ONLY0: arrayctor.cont: 1384 // SIMD-ONLY0-NEXT: ret void 1385 // 1386 // 1387 // SIMD-ONLY0-LABEL: define {{[^@]+}}@__cxx_global_var_init.4 1388 // SIMD-ONLY0-SAME: () #[[ATTR0]] { 1389 // SIMD-ONLY0-NEXT: entry: 1390 // SIMD-ONLY0-NEXT: call void @_ZN2S3C1Ev(ptr noundef nonnull align 4 dereferenceable(4) @h) 1391 // SIMD-ONLY0-NEXT: ret void 1392 // 1393 // 1394 // SIMD-ONLY0-LABEL: define {{[^@]+}}@main 1395 // SIMD-ONLY0-SAME: (i32 noundef signext [[ARGC:%.*]], ptr noundef [[ARGV:%.*]]) #[[ATTR2:[0-9]+]] { 1396 // SIMD-ONLY0-NEXT: entry: 1397 // SIMD-ONLY0-NEXT: [[RETVAL:%.*]] = alloca i32, align 4 1398 // SIMD-ONLY0-NEXT: [[ARGC_ADDR:%.*]] = alloca i32, align 4 1399 // SIMD-ONLY0-NEXT: [[ARGV_ADDR:%.*]] = alloca ptr, align 8 1400 // SIMD-ONLY0-NEXT: [[DA:%.*]] = alloca [5 x i32], align 4 1401 // SIMD-ONLY0-NEXT: [[H:%.*]] = alloca [10 x %struct.S6], align 4 1402 // SIMD-ONLY0-NEXT: [[RH:%.*]] = alloca ptr, align 8 1403 // SIMD-ONLY0-NEXT: [[I:%.*]] = alloca i32, align 4 1404 // SIMD-ONLY0-NEXT: [[J:%.*]] = alloca ptr, align 8 1405 // SIMD-ONLY0-NEXT: [[K:%.*]] = alloca ptr, align 8 1406 // SIMD-ONLY0-NEXT: [[Z:%.*]] = alloca ptr, align 8 1407 // SIMD-ONLY0-NEXT: [[AA:%.*]] = alloca [10 x i32], align 4 1408 // SIMD-ONLY0-NEXT: [[RAA:%.*]] = alloca ptr, align 8 1409 // SIMD-ONLY0-NEXT: [[TMP:%.*]] = alloca ptr, align 8 1410 // SIMD-ONLY0-NEXT: [[_TMP2:%.*]] = alloca ptr, align 8 1411 // SIMD-ONLY0-NEXT: [[A:%.*]] = alloca i32, align 4 1412 // SIMD-ONLY0-NEXT: [[A4:%.*]] = alloca i32, align 4 1413 // SIMD-ONLY0-NEXT: [[A7:%.*]] = alloca i32, align 4 1414 // SIMD-ONLY0-NEXT: store i32 0, ptr [[RETVAL]], align 4 1415 // SIMD-ONLY0-NEXT: store i32 [[ARGC]], ptr [[ARGC_ADDR]], align 4 1416 // SIMD-ONLY0-NEXT: store ptr [[ARGV]], ptr [[ARGV_ADDR]], align 8 1417 // SIMD-ONLY0-NEXT: call void @llvm.memset.p0.i64(ptr align 4 [[DA]], i8 0, i64 20, i1 false) 1418 // SIMD-ONLY0-NEXT: store ptr [[H]], ptr [[RH]], align 8 1419 // SIMD-ONLY0-NEXT: store ptr [[I]], ptr [[J]], align 8 1420 // SIMD-ONLY0-NEXT: [[TMP0:%.*]] = load ptr, ptr [[J]], align 8 1421 // SIMD-ONLY0-NEXT: store ptr [[TMP0]], ptr [[K]], align 8 1422 // SIMD-ONLY0-NEXT: store ptr [[K]], ptr [[Z]], align 8 1423 // SIMD-ONLY0-NEXT: store ptr [[AA]], ptr [[RAA]], align 8 1424 // SIMD-ONLY0-NEXT: [[TMP1:%.*]] = load ptr, ptr [[K]], align 8 1425 // SIMD-ONLY0-NEXT: [[INCDEC_PTR:%.*]] = getelementptr inbounds nuw i32, ptr [[TMP1]], i32 1 1426 // SIMD-ONLY0-NEXT: store ptr [[INCDEC_PTR]], ptr [[K]], align 8 1427 // SIMD-ONLY0-NEXT: [[TMP2:%.*]] = load ptr, ptr [[Z]], align 8 1428 // SIMD-ONLY0-NEXT: store ptr [[TMP2]], ptr [[TMP]], align 8 1429 // SIMD-ONLY0-NEXT: [[TMP3:%.*]] = load ptr, ptr [[Z]], align 8 1430 // SIMD-ONLY0-NEXT: [[TMP4:%.*]] = load ptr, ptr [[TMP]], align 8 1431 // SIMD-ONLY0-NEXT: [[TMP5:%.*]] = load ptr, ptr [[TMP4]], align 8 1432 // SIMD-ONLY0-NEXT: [[INCDEC_PTR1:%.*]] = getelementptr inbounds nuw i32, ptr [[TMP5]], i32 1 1433 // SIMD-ONLY0-NEXT: store ptr [[INCDEC_PTR1]], ptr [[TMP4]], align 8 1434 // SIMD-ONLY0-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x i32], ptr [[AA]], i64 0, i64 0 1435 // SIMD-ONLY0-NEXT: store i32 1, ptr [[ARRAYIDX]], align 4 1436 // SIMD-ONLY0-NEXT: [[TMP6:%.*]] = load ptr, ptr [[RAA]], align 8 1437 // SIMD-ONLY0-NEXT: store ptr [[TMP6]], ptr [[_TMP2]], align 8 1438 // SIMD-ONLY0-NEXT: [[TMP7:%.*]] = load ptr, ptr [[RAA]], align 8 1439 // SIMD-ONLY0-NEXT: [[TMP8:%.*]] = load ptr, ptr [[_TMP2]], align 8 1440 // SIMD-ONLY0-NEXT: [[ARRAYIDX3:%.*]] = getelementptr inbounds [10 x i32], ptr [[TMP8]], i64 0, i64 0 1441 // SIMD-ONLY0-NEXT: [[TMP9:%.*]] = load i32, ptr [[ARRAYIDX3]], align 4 1442 // SIMD-ONLY0-NEXT: store i32 [[TMP9]], ptr [[A]], align 4 1443 // SIMD-ONLY0-NEXT: [[ARRAYIDX5:%.*]] = getelementptr inbounds [10 x %struct.S6], ptr [[H]], i64 0, i64 1 1444 // SIMD-ONLY0-NEXT: [[A6:%.*]] = getelementptr inbounds nuw [[STRUCT_S6:%.*]], ptr [[ARRAYIDX5]], i32 0, i32 0 1445 // SIMD-ONLY0-NEXT: [[TMP10:%.*]] = load i32, ptr [[A6]], align 4 1446 // SIMD-ONLY0-NEXT: store i32 [[TMP10]], ptr [[A4]], align 4 1447 // SIMD-ONLY0-NEXT: [[ARRAYIDX8:%.*]] = getelementptr inbounds [5 x i32], ptr [[DA]], i64 0, i64 1 1448 // SIMD-ONLY0-NEXT: [[TMP11:%.*]] = load i32, ptr [[ARRAYIDX8]], align 4 1449 // SIMD-ONLY0-NEXT: store i32 [[TMP11]], ptr [[A7]], align 4 1450 // SIMD-ONLY0-NEXT: [[TMP12:%.*]] = load i32, ptr [[ARGC_ADDR]], align 4 1451 // SIMD-ONLY0-NEXT: [[CALL:%.*]] = call noundef signext i32 @_Z5tmainIiET_S0_(i32 noundef signext [[TMP12]]) 1452 // SIMD-ONLY0-NEXT: [[CALL9:%.*]] = call noundef ptr @_Z5tmainIPiET_S1_(ptr noundef [[ARGC_ADDR]]) 1453 // SIMD-ONLY0-NEXT: [[TMP13:%.*]] = load i32, ptr [[CALL9]], align 4 1454 // SIMD-ONLY0-NEXT: [[ADD:%.*]] = add nsw i32 [[CALL]], [[TMP13]] 1455 // SIMD-ONLY0-NEXT: ret i32 [[ADD]] 1456 // 1457 // 1458 // SIMD-ONLY0-LABEL: define {{[^@]+}}@_Z5tmainIiET_S0_ 1459 // SIMD-ONLY0-SAME: (i32 noundef signext [[ARGC:%.*]]) #[[ATTR1]] comdat { 1460 // SIMD-ONLY0-NEXT: entry: 1461 // SIMD-ONLY0-NEXT: [[ARGC_ADDR:%.*]] = alloca i32, align 4 1462 // SIMD-ONLY0-NEXT: [[DA:%.*]] = alloca [5 x i32], align 4 1463 // SIMD-ONLY0-NEXT: [[H:%.*]] = alloca [10 x %struct.S6], align 4 1464 // SIMD-ONLY0-NEXT: [[RH:%.*]] = alloca ptr, align 8 1465 // SIMD-ONLY0-NEXT: [[I:%.*]] = alloca i32, align 4 1466 // SIMD-ONLY0-NEXT: [[J:%.*]] = alloca ptr, align 8 1467 // SIMD-ONLY0-NEXT: [[K:%.*]] = alloca ptr, align 8 1468 // SIMD-ONLY0-NEXT: [[Z:%.*]] = alloca ptr, align 8 1469 // SIMD-ONLY0-NEXT: [[AA:%.*]] = alloca [10 x i32], align 4 1470 // SIMD-ONLY0-NEXT: [[TMP:%.*]] = alloca ptr, align 8 1471 // SIMD-ONLY0-NEXT: [[A:%.*]] = alloca i32, align 4 1472 // SIMD-ONLY0-NEXT: [[A2:%.*]] = alloca i32, align 4 1473 // SIMD-ONLY0-NEXT: store i32 [[ARGC]], ptr [[ARGC_ADDR]], align 4 1474 // SIMD-ONLY0-NEXT: call void @llvm.memset.p0.i64(ptr align 4 [[DA]], i8 0, i64 20, i1 false) 1475 // SIMD-ONLY0-NEXT: store ptr [[H]], ptr [[RH]], align 8 1476 // SIMD-ONLY0-NEXT: store ptr [[I]], ptr [[J]], align 8 1477 // SIMD-ONLY0-NEXT: [[TMP0:%.*]] = load ptr, ptr [[J]], align 8 1478 // SIMD-ONLY0-NEXT: store ptr [[TMP0]], ptr [[K]], align 8 1479 // SIMD-ONLY0-NEXT: store ptr [[K]], ptr [[Z]], align 8 1480 // SIMD-ONLY0-NEXT: [[TMP1:%.*]] = load ptr, ptr [[K]], align 8 1481 // SIMD-ONLY0-NEXT: [[INCDEC_PTR:%.*]] = getelementptr inbounds nuw i32, ptr [[TMP1]], i32 1 1482 // SIMD-ONLY0-NEXT: store ptr [[INCDEC_PTR]], ptr [[K]], align 8 1483 // SIMD-ONLY0-NEXT: [[TMP2:%.*]] = load ptr, ptr [[Z]], align 8 1484 // SIMD-ONLY0-NEXT: store ptr [[TMP2]], ptr [[TMP]], align 8 1485 // SIMD-ONLY0-NEXT: [[TMP3:%.*]] = load ptr, ptr [[Z]], align 8 1486 // SIMD-ONLY0-NEXT: [[TMP4:%.*]] = load ptr, ptr [[TMP]], align 8 1487 // SIMD-ONLY0-NEXT: [[TMP5:%.*]] = load ptr, ptr [[TMP4]], align 8 1488 // SIMD-ONLY0-NEXT: [[INCDEC_PTR1:%.*]] = getelementptr inbounds nuw i32, ptr [[TMP5]], i32 1 1489 // SIMD-ONLY0-NEXT: store ptr [[INCDEC_PTR1]], ptr [[TMP4]], align 8 1490 // SIMD-ONLY0-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x i32], ptr [[AA]], i64 0, i64 0 1491 // SIMD-ONLY0-NEXT: [[TMP6:%.*]] = load i32, ptr [[ARRAYIDX]], align 4 1492 // SIMD-ONLY0-NEXT: store i32 [[TMP6]], ptr [[A]], align 4 1493 // SIMD-ONLY0-NEXT: [[ARRAYIDX3:%.*]] = getelementptr inbounds [10 x %struct.S6], ptr [[H]], i64 0, i64 0 1494 // SIMD-ONLY0-NEXT: [[A4:%.*]] = getelementptr inbounds nuw [[STRUCT_S6:%.*]], ptr [[ARRAYIDX3]], i32 0, i32 0 1495 // SIMD-ONLY0-NEXT: [[TMP7:%.*]] = load i32, ptr [[A4]], align 4 1496 // SIMD-ONLY0-NEXT: store i32 [[TMP7]], ptr [[A2]], align 4 1497 // SIMD-ONLY0-NEXT: ret i32 0 1498 // 1499 // 1500 // SIMD-ONLY0-LABEL: define {{[^@]+}}@_Z5tmainIPiET_S1_ 1501 // SIMD-ONLY0-SAME: (ptr noundef [[ARGC:%.*]]) #[[ATTR1]] comdat { 1502 // SIMD-ONLY0-NEXT: entry: 1503 // SIMD-ONLY0-NEXT: [[ARGC_ADDR:%.*]] = alloca ptr, align 8 1504 // SIMD-ONLY0-NEXT: [[DA:%.*]] = alloca [5 x ptr], align 8 1505 // SIMD-ONLY0-NEXT: [[H:%.*]] = alloca [10 x %struct.S6], align 4 1506 // SIMD-ONLY0-NEXT: [[RH:%.*]] = alloca ptr, align 8 1507 // SIMD-ONLY0-NEXT: [[I:%.*]] = alloca ptr, align 8 1508 // SIMD-ONLY0-NEXT: [[J:%.*]] = alloca ptr, align 8 1509 // SIMD-ONLY0-NEXT: [[K:%.*]] = alloca ptr, align 8 1510 // SIMD-ONLY0-NEXT: [[Z:%.*]] = alloca ptr, align 8 1511 // SIMD-ONLY0-NEXT: [[AA:%.*]] = alloca [10 x ptr], align 8 1512 // SIMD-ONLY0-NEXT: [[TMP:%.*]] = alloca ptr, align 8 1513 // SIMD-ONLY0-NEXT: [[A:%.*]] = alloca ptr, align 8 1514 // SIMD-ONLY0-NEXT: [[A2:%.*]] = alloca i32, align 4 1515 // SIMD-ONLY0-NEXT: store ptr [[ARGC]], ptr [[ARGC_ADDR]], align 8 1516 // SIMD-ONLY0-NEXT: call void @llvm.memset.p0.i64(ptr align 8 [[DA]], i8 0, i64 40, i1 false) 1517 // SIMD-ONLY0-NEXT: store ptr [[H]], ptr [[RH]], align 8 1518 // SIMD-ONLY0-NEXT: store ptr [[I]], ptr [[J]], align 8 1519 // SIMD-ONLY0-NEXT: [[TMP0:%.*]] = load ptr, ptr [[J]], align 8 1520 // SIMD-ONLY0-NEXT: store ptr [[TMP0]], ptr [[K]], align 8 1521 // SIMD-ONLY0-NEXT: store ptr [[K]], ptr [[Z]], align 8 1522 // SIMD-ONLY0-NEXT: [[TMP1:%.*]] = load ptr, ptr [[K]], align 8 1523 // SIMD-ONLY0-NEXT: [[INCDEC_PTR:%.*]] = getelementptr inbounds nuw ptr, ptr [[TMP1]], i32 1 1524 // SIMD-ONLY0-NEXT: store ptr [[INCDEC_PTR]], ptr [[K]], align 8 1525 // SIMD-ONLY0-NEXT: [[TMP2:%.*]] = load ptr, ptr [[Z]], align 8 1526 // SIMD-ONLY0-NEXT: store ptr [[TMP2]], ptr [[TMP]], align 8 1527 // SIMD-ONLY0-NEXT: [[TMP3:%.*]] = load ptr, ptr [[Z]], align 8 1528 // SIMD-ONLY0-NEXT: [[TMP4:%.*]] = load ptr, ptr [[TMP]], align 8 1529 // SIMD-ONLY0-NEXT: [[TMP5:%.*]] = load ptr, ptr [[TMP4]], align 8 1530 // SIMD-ONLY0-NEXT: [[INCDEC_PTR1:%.*]] = getelementptr inbounds nuw ptr, ptr [[TMP5]], i32 1 1531 // SIMD-ONLY0-NEXT: store ptr [[INCDEC_PTR1]], ptr [[TMP4]], align 8 1532 // SIMD-ONLY0-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x ptr], ptr [[AA]], i64 0, i64 0 1533 // SIMD-ONLY0-NEXT: [[TMP6:%.*]] = load ptr, ptr [[ARRAYIDX]], align 8 1534 // SIMD-ONLY0-NEXT: store ptr [[TMP6]], ptr [[A]], align 8 1535 // SIMD-ONLY0-NEXT: [[ARRAYIDX3:%.*]] = getelementptr inbounds [10 x %struct.S6], ptr [[H]], i64 0, i64 0 1536 // SIMD-ONLY0-NEXT: [[A4:%.*]] = getelementptr inbounds nuw [[STRUCT_S6:%.*]], ptr [[ARRAYIDX3]], i32 0, i32 0 1537 // SIMD-ONLY0-NEXT: [[TMP7:%.*]] = load i32, ptr [[A4]], align 4 1538 // SIMD-ONLY0-NEXT: store i32 [[TMP7]], ptr [[A2]], align 4 1539 // SIMD-ONLY0-NEXT: ret ptr null 1540 // 1541 // 1542 // SIMD-ONLY0-LABEL: define {{[^@]+}}@_Z12use_templatev 1543 // SIMD-ONLY0-SAME: () #[[ATTR1]] { 1544 // SIMD-ONLY0-NEXT: entry: 1545 // SIMD-ONLY0-NEXT: [[AKERN:%.*]] = alloca [[STRUCT_SOMEKERNEL:%.*]], align 4 1546 // SIMD-ONLY0-NEXT: call void @_ZN10SomeKernelC1Ev(ptr noundef nonnull align 4 dereferenceable(8) [[AKERN]]) 1547 // SIMD-ONLY0-NEXT: call void @_ZN10SomeKernel5applyILj32EEEvv(ptr noundef nonnull align 4 dereferenceable(8) [[AKERN]]) 1548 // SIMD-ONLY0-NEXT: call void @_ZN10SomeKernelD1Ev(ptr noundef nonnull align 4 dereferenceable(8) [[AKERN]]) #[[ATTR6:[0-9]+]] 1549 // SIMD-ONLY0-NEXT: ret void 1550 // 1551 // 1552 // SIMD-ONLY0-LABEL: define {{[^@]+}}@_ZN10SomeKernel5applyILj32EEEvv 1553 // SIMD-ONLY0-SAME: (ptr noundef nonnull align 4 dereferenceable(8) [[THIS:%.*]]) #[[ATTR1]] comdat { 1554 // SIMD-ONLY0-NEXT: entry: 1555 // SIMD-ONLY0-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 8 1556 // SIMD-ONLY0-NEXT: [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 4 1557 // SIMD-ONLY0-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 8 1558 // SIMD-ONLY0-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8 1559 // SIMD-ONLY0-NEXT: [[TARGETDEV:%.*]] = getelementptr inbounds nuw [[STRUCT_SOMEKERNEL:%.*]], ptr [[THIS1]], i32 0, i32 0 1560 // SIMD-ONLY0-NEXT: [[TMP0:%.*]] = load i32, ptr [[TARGETDEV]], align 4 1561 // SIMD-ONLY0-NEXT: store i32 [[TMP0]], ptr [[DOTCAPTURE_EXPR_]], align 4 1562 // SIMD-ONLY0-NEXT: [[DEVPTR:%.*]] = getelementptr inbounds nuw [[STRUCT_SOMEKERNEL]], ptr [[THIS1]], i32 0, i32 1 1563 // SIMD-ONLY0-NEXT: [[TMP1:%.*]] = load float, ptr [[DEVPTR]], align 4 1564 // SIMD-ONLY0-NEXT: [[INC:%.*]] = fadd float [[TMP1]], 1.000000e+00 1565 // SIMD-ONLY0-NEXT: store float [[INC]], ptr [[DEVPTR]], align 4 1566 // SIMD-ONLY0-NEXT: [[TARGETDEV2:%.*]] = getelementptr inbounds nuw [[STRUCT_SOMEKERNEL]], ptr [[THIS1]], i32 0, i32 0 1567 // SIMD-ONLY0-NEXT: [[TMP2:%.*]] = load i32, ptr [[TARGETDEV2]], align 4 1568 // SIMD-ONLY0-NEXT: [[INC3:%.*]] = add nsw i32 [[TMP2]], 1 1569 // SIMD-ONLY0-NEXT: store i32 [[INC3]], ptr [[TARGETDEV2]], align 4 1570 // SIMD-ONLY0-NEXT: ret void 1571 // 1572 // 1573 // SIMD-ONLY0-LABEL: define {{[^@]+}}@_ZN2S2C2Ev 1574 // SIMD-ONLY0-SAME: (ptr noundef nonnull align 4 dereferenceable(4) [[THIS:%.*]]) unnamed_addr #[[ATTR1]] comdat { 1575 // SIMD-ONLY0-NEXT: entry: 1576 // SIMD-ONLY0-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 8 1577 // SIMD-ONLY0-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 8 1578 // SIMD-ONLY0-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8 1579 // SIMD-ONLY0-NEXT: [[A:%.*]] = getelementptr inbounds nuw [[CLASS_S2:%.*]], ptr [[THIS1]], i32 0, i32 0 1580 // SIMD-ONLY0-NEXT: store i32 0, ptr [[A]], align 4 1581 // SIMD-ONLY0-NEXT: ret void 1582 // 1583 // 1584 // SIMD-ONLY0-LABEL: define {{[^@]+}}@_ZN2S3C2Ev 1585 // SIMD-ONLY0-SAME: (ptr noundef nonnull align 4 dereferenceable(4) [[THIS:%.*]]) unnamed_addr #[[ATTR1]] comdat { 1586 // SIMD-ONLY0-NEXT: entry: 1587 // SIMD-ONLY0-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 8 1588 // SIMD-ONLY0-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 8 1589 // SIMD-ONLY0-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8 1590 // SIMD-ONLY0-NEXT: [[A:%.*]] = getelementptr inbounds nuw [[CLASS_S3:%.*]], ptr [[THIS1]], i32 0, i32 0 1591 // SIMD-ONLY0-NEXT: store i32 0, ptr [[A]], align 4 1592 // SIMD-ONLY0-NEXT: ret void 1593 // 1594 // 1595 // SIMD-ONLY0-LABEL: define {{[^@]+}}@_GLOBAL__sub_I_target_has_device_addr_codegen.cpp 1596 // SIMD-ONLY0-SAME: () #[[ATTR0]] { 1597 // SIMD-ONLY0-NEXT: entry: 1598 // SIMD-ONLY0-NEXT: call void @__cxx_global_var_init() 1599 // SIMD-ONLY0-NEXT: call void @__cxx_global_var_init.1() 1600 // SIMD-ONLY0-NEXT: call void @__cxx_global_var_init.2() 1601 // SIMD-ONLY0-NEXT: call void @__cxx_global_var_init.3() 1602 // SIMD-ONLY0-NEXT: call void @__cxx_global_var_init.4() 1603 // SIMD-ONLY0-NEXT: ret void 1604 // 1605