765 lines · cpp
1// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --include-generated-funcs --replace-value-regex "__omp_offloading_[0-9a-z]+_[0-9a-z]+" "reduction_size[.].+[.]" "pl_cond[.].+[.|,]" --prefix-filecheck-ir-name _ --version 52// expected-no-diagnostics3#ifndef HEADER4#define HEADER5 6///==========================================================================///7// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK1 -verify -Wno-vla -fopenmp -fopenmp-version=60 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK1-648// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK1 -fopenmp -fopenmp-version=60 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s9// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-version=60 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK1-6410// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK1 -verify -Wno-vla -fopenmp -fopenmp-version=60 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK1-3211// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK1 -fopenmp -fopenmp-version=60 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s12// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-version=60 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK1-3213 14// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK1 -verify -Wno-vla -fopenmp-simd -fopenmp-version=60 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY1-64 %s15// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK1 -fopenmp-simd -fopenmp-version=60 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s16// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-version=60 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY1-64 %s17// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK1 -verify -Wno-vla -fopenmp-simd -fopenmp-version=60 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY1-32 %s18// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK1 -fopenmp-simd -fopenmp-version=60 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s19// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-version=60 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY1-32 %s20#ifdef CK121void foo1(int a){22 double d = (double)a;23 24 #pragma omp target defaultmap(private : scalar)25 {26 d += 1.0;27 }28}29#endif30 31///==========================================================================///32// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK2 -verify -Wno-vla -fopenmp -fopenmp-version=60 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK2-6433// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK2 -fopenmp -fopenmp-version=60 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s34// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-version=60 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK2-6435// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK2 -verify -Wno-vla -fopenmp -fopenmp-version=60 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK2-3236// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK2 -fopenmp -fopenmp-version=60 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s37// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-version=60 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK2-3238 39// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK2 -verify -Wno-vla -fopenmp-simd -fopenmp-version=60 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY2-64 %s40// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK2 -fopenmp-simd -fopenmp-version=60 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s41// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-version=60 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY2-64 %s42// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK2 -verify -Wno-vla -fopenmp-simd -fopenmp-version=60 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY2-32 %s43// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK2 -fopenmp-simd -fopenmp-version=60 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s44// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-version=60 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY2-32 %s45 46#ifdef CK247void foo2(){48 int pvtArr[10];49 50 #pragma omp target defaultmap(private : aggregate)51 {52 pvtArr[5]++;53 }54}55#endif56 57///==========================================================================///58// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK3 -verify -Wno-vla -fopenmp -fopenmp-version=60 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK3-6459// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK3 -fopenmp -fopenmp-version=60 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s60// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-version=60 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK3-6461// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK3 -verify -Wno-vla -fopenmp -fopenmp-version=60 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK3-3262// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK3 -fopenmp -fopenmp-version=60 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s63// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-version=60 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK3-3264 65// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK3 -verify -Wno-vla -fopenmp-simd -fopenmp-version=60 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY3-64 %s66// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK3 -fopenmp-simd -fopenmp-version=60 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s67// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-version=60 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY3-64 %s68// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK3 -verify -Wno-vla -fopenmp-simd -fopenmp-version=60 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY3-32 %s69// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK3 -fopenmp-simd -fopenmp-version=60 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s70// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-version=60 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY3-32 %s71#ifdef CK372void foo3(){73 int *pa;74 75 #pragma omp target defaultmap(private : pointer)76 {77 pa[50]++;78 }79}80#endif81 82///==========================================================================///83// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK4 -verify -Wno-vla -fopenmp -fopenmp-version=60 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK4-6484// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK4 -fopenmp -fopenmp-version=60 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s85// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-version=60 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK4-6486// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK4 -verify -Wno-vla -fopenmp -fopenmp-version=60 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK4-3287// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK4 -fopenmp -fopenmp-version=60 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s88// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-version=60 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK4-3289 90// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK4 -verify -Wno-vla -fopenmp-simd -fopenmp-version=60 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY4-64 %s91// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK4 -fopenmp-simd -fopenmp-version=60 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s92// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-version=60 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY4-64 %s93// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK4 -verify -Wno-vla -fopenmp-simd -fopenmp-version=60 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY4-32 %s94// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK4 -fopenmp-simd -fopenmp-version=60 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s95// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-version=60 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY4-32 %s96 97// Specified variable-category doesn't apply to referenced variable, so98// normal implicitly determined data-sharing applies.99#ifdef CK4100void foo4(){101 int p;102 103 #pragma omp target defaultmap(private : pointer)104 {105 p++;106 }107}108#endif109 110#endif // HEADER111// CK1-64-LABEL: define dso_local void @_Z4foo1i(112// CK1-64-SAME: i32 signext [[A:%.*]]) #[[ATTR0:[0-9]+]] {113// CK1-64-NEXT: [[ENTRY:.*:]]114// CK1-64-NEXT: [[A_ADDR:%.*]] = alloca i32, align 4115// CK1-64-NEXT: [[D:%.*]] = alloca double, align 8116// CK1-64-NEXT: [[D_CASTED:%.*]] = alloca i64, align 8117// CK1-64-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x ptr], align 8118// CK1-64-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x ptr], align 8119// CK1-64-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [1 x ptr], align 8120// CK1-64-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8121// CK1-64-NEXT: store i32 [[A]], ptr [[A_ADDR]], align 4122// CK1-64-NEXT: [[TMP0:%.*]] = load i32, ptr [[A_ADDR]], align 4123// CK1-64-NEXT: [[CONV:%.*]] = sitofp i32 [[TMP0]] to double124// CK1-64-NEXT: store double [[CONV]], ptr [[D]], align 8125// CK1-64-NEXT: [[TMP1:%.*]] = load double, ptr [[D]], align 8126// CK1-64-NEXT: store double [[TMP1]], ptr [[D_CASTED]], align 8127// CK1-64-NEXT: [[TMP2:%.*]] = load i64, ptr [[D_CASTED]], align 8128// CK1-64-NEXT: [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0129// CK1-64-NEXT: store i64 [[TMP2]], ptr [[TMP3]], align 8130// CK1-64-NEXT: [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0131// CK1-64-NEXT: store i64 [[TMP2]], ptr [[TMP4]], align 8132// CK1-64-NEXT: [[TMP5:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 0133// CK1-64-NEXT: store ptr null, ptr [[TMP5]], align 8134// CK1-64-NEXT: [[TMP6:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0135// CK1-64-NEXT: [[TMP7:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0136// CK1-64-NEXT: [[TMP8:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0137// CK1-64-NEXT: store i32 3, ptr [[TMP8]], align 4138// CK1-64-NEXT: [[TMP9:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1139// CK1-64-NEXT: store i32 1, ptr [[TMP9]], align 4140// CK1-64-NEXT: [[TMP10:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2141// CK1-64-NEXT: store ptr [[TMP6]], ptr [[TMP10]], align 8142// CK1-64-NEXT: [[TMP11:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3143// CK1-64-NEXT: store ptr [[TMP7]], ptr [[TMP11]], align 8144// CK1-64-NEXT: [[TMP12:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4145// CK1-64-NEXT: store ptr @.offload_sizes, ptr [[TMP12]], align 8146// CK1-64-NEXT: [[TMP13:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5147// CK1-64-NEXT: store ptr @.offload_maptypes, ptr [[TMP13]], align 8148// CK1-64-NEXT: [[TMP14:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6149// CK1-64-NEXT: store ptr null, ptr [[TMP14]], align 8150// CK1-64-NEXT: [[TMP15:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7151// CK1-64-NEXT: store ptr null, ptr [[TMP15]], align 8152// CK1-64-NEXT: [[TMP16:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8153// CK1-64-NEXT: store i64 0, ptr [[TMP16]], align 8154// CK1-64-NEXT: [[TMP17:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9155// CK1-64-NEXT: store i64 0, ptr [[TMP17]], align 8156// CK1-64-NEXT: [[TMP18:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10157// CK1-64-NEXT: store [3 x i32] [i32 -1, i32 0, i32 0], ptr [[TMP18]], align 4158// CK1-64-NEXT: [[TMP19:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11159// CK1-64-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP19]], align 4160// CK1-64-NEXT: [[TMP20:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12161// CK1-64-NEXT: store i32 0, ptr [[TMP20]], align 4162// CK1-64-NEXT: [[TMP21:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1:[0-9]+]], i64 -1, i32 -1, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z4foo1i_l24.region_id, ptr [[KERNEL_ARGS]])163// CK1-64-NEXT: [[TMP22:%.*]] = icmp ne i32 [[TMP21]], 0164// CK1-64-NEXT: br i1 [[TMP22]], label %[[OMP_OFFLOAD_FAILED:.*]], label %[[OMP_OFFLOAD_CONT:.*]]165// CK1-64: [[OMP_OFFLOAD_FAILED]]:166// CK1-64-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z4foo1i_l24(i64 [[TMP2]]) #[[ATTR2:[0-9]+]]167// CK1-64-NEXT: br label %[[OMP_OFFLOAD_CONT]]168// CK1-64: [[OMP_OFFLOAD_CONT]]:169// CK1-64-NEXT: ret void170//171//172// CK1-64-LABEL: define internal void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z4foo1i_l24(173// CK1-64-SAME: i64 [[D:%.*]]) #[[ATTR1:[0-9]+]] {174// CK1-64-NEXT: [[ENTRY:.*:]]175// CK1-64-NEXT: [[D_ADDR:%.*]] = alloca i64, align 8176// CK1-64-NEXT: [[D1:%.*]] = alloca double, align 8177// CK1-64-NEXT: store i64 [[D]], ptr [[D_ADDR]], align 8178// CK1-64-NEXT: [[TMP0:%.*]] = load double, ptr [[D1]], align 8179// CK1-64-NEXT: [[ADD:%.*]] = fadd double [[TMP0]], 1.000000e+00180// CK1-64-NEXT: store double [[ADD]], ptr [[D1]], align 8181// CK1-64-NEXT: ret void182//183//184// CK1-32-LABEL: define dso_local void @_Z4foo1i(185// CK1-32-SAME: i32 [[A:%.*]]) #[[ATTR0:[0-9]+]] {186// CK1-32-NEXT: [[ENTRY:.*:]]187// CK1-32-NEXT: [[A_ADDR:%.*]] = alloca i32, align 4188// CK1-32-NEXT: [[D:%.*]] = alloca double, align 8189// CK1-32-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x ptr], align 4190// CK1-32-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x ptr], align 4191// CK1-32-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [1 x ptr], align 4192// CK1-32-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8193// CK1-32-NEXT: store i32 [[A]], ptr [[A_ADDR]], align 4194// CK1-32-NEXT: [[TMP0:%.*]] = load i32, ptr [[A_ADDR]], align 4195// CK1-32-NEXT: [[CONV:%.*]] = sitofp i32 [[TMP0]] to double196// CK1-32-NEXT: store double [[CONV]], ptr [[D]], align 8197// CK1-32-NEXT: [[TMP1:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0198// CK1-32-NEXT: store ptr [[D]], ptr [[TMP1]], align 4199// CK1-32-NEXT: [[TMP2:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0200// CK1-32-NEXT: store ptr [[D]], ptr [[TMP2]], align 4201// CK1-32-NEXT: [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 0202// CK1-32-NEXT: store ptr null, ptr [[TMP3]], align 4203// CK1-32-NEXT: [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0204// CK1-32-NEXT: [[TMP5:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0205// CK1-32-NEXT: [[TMP6:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0206// CK1-32-NEXT: store i32 3, ptr [[TMP6]], align 4207// CK1-32-NEXT: [[TMP7:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1208// CK1-32-NEXT: store i32 1, ptr [[TMP7]], align 4209// CK1-32-NEXT: [[TMP8:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2210// CK1-32-NEXT: store ptr [[TMP4]], ptr [[TMP8]], align 4211// CK1-32-NEXT: [[TMP9:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3212// CK1-32-NEXT: store ptr [[TMP5]], ptr [[TMP9]], align 4213// CK1-32-NEXT: [[TMP10:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4214// CK1-32-NEXT: store ptr @.offload_sizes, ptr [[TMP10]], align 4215// CK1-32-NEXT: [[TMP11:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5216// CK1-32-NEXT: store ptr @.offload_maptypes, ptr [[TMP11]], align 4217// CK1-32-NEXT: [[TMP12:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6218// CK1-32-NEXT: store ptr null, ptr [[TMP12]], align 4219// CK1-32-NEXT: [[TMP13:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7220// CK1-32-NEXT: store ptr null, ptr [[TMP13]], align 4221// CK1-32-NEXT: [[TMP14:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8222// CK1-32-NEXT: store i64 0, ptr [[TMP14]], align 8223// CK1-32-NEXT: [[TMP15:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9224// CK1-32-NEXT: store i64 0, ptr [[TMP15]], align 8225// CK1-32-NEXT: [[TMP16:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10226// CK1-32-NEXT: store [3 x i32] [i32 -1, i32 0, i32 0], ptr [[TMP16]], align 4227// CK1-32-NEXT: [[TMP17:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11228// CK1-32-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP17]], align 4229// CK1-32-NEXT: [[TMP18:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12230// CK1-32-NEXT: store i32 0, ptr [[TMP18]], align 4231// CK1-32-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]+}}__Z4foo1i_l24.region_id, ptr [[KERNEL_ARGS]])232// CK1-32-NEXT: [[TMP20:%.*]] = icmp ne i32 [[TMP19]], 0233// CK1-32-NEXT: br i1 [[TMP20]], label %[[OMP_OFFLOAD_FAILED:.*]], label %[[OMP_OFFLOAD_CONT:.*]]234// CK1-32: [[OMP_OFFLOAD_FAILED]]:235// CK1-32-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z4foo1i_l24(ptr [[D]]) #[[ATTR2:[0-9]+]]236// CK1-32-NEXT: br label %[[OMP_OFFLOAD_CONT]]237// CK1-32: [[OMP_OFFLOAD_CONT]]:238// CK1-32-NEXT: ret void239//240//241// CK1-32-LABEL: define internal void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z4foo1i_l24(242// CK1-32-SAME: ptr nonnull align 4 dereferenceable(8) [[D:%.*]]) #[[ATTR1:[0-9]+]] {243// CK1-32-NEXT: [[ENTRY:.*:]]244// CK1-32-NEXT: [[D_ADDR:%.*]] = alloca ptr, align 4245// CK1-32-NEXT: [[D1:%.*]] = alloca double, align 8246// CK1-32-NEXT: store ptr [[D]], ptr [[D_ADDR]], align 4247// CK1-32-NEXT: [[TMP0:%.*]] = load ptr, ptr [[D_ADDR]], align 4, !nonnull [[META6:![0-9]+]], !align [[META7:![0-9]+]]248// CK1-32-NEXT: [[TMP1:%.*]] = load double, ptr [[D1]], align 8249// CK1-32-NEXT: [[ADD:%.*]] = fadd double [[TMP1]], 1.000000e+00250// CK1-32-NEXT: store double [[ADD]], ptr [[D1]], align 8251// CK1-32-NEXT: ret void252//253//254// SIMD-ONLY1-64-LABEL: define dso_local void @_Z4foo1i(255// SIMD-ONLY1-64-SAME: i32 signext [[A:%.*]]) #[[ATTR0:[0-9]+]] {256// SIMD-ONLY1-64-NEXT: [[ENTRY:.*:]]257// SIMD-ONLY1-64-NEXT: [[A_ADDR:%.*]] = alloca i32, align 4258// SIMD-ONLY1-64-NEXT: [[D:%.*]] = alloca double, align 8259// SIMD-ONLY1-64-NEXT: [[D1:%.*]] = alloca double, align 8260// SIMD-ONLY1-64-NEXT: store i32 [[A]], ptr [[A_ADDR]], align 4261// SIMD-ONLY1-64-NEXT: [[TMP0:%.*]] = load i32, ptr [[A_ADDR]], align 4262// SIMD-ONLY1-64-NEXT: [[CONV:%.*]] = sitofp i32 [[TMP0]] to double263// SIMD-ONLY1-64-NEXT: store double [[CONV]], ptr [[D]], align 8264// SIMD-ONLY1-64-NEXT: [[TMP1:%.*]] = load double, ptr [[D1]], align 8265// SIMD-ONLY1-64-NEXT: [[ADD:%.*]] = fadd double [[TMP1]], 1.000000e+00266// SIMD-ONLY1-64-NEXT: store double [[ADD]], ptr [[D1]], align 8267// SIMD-ONLY1-64-NEXT: ret void268//269//270// SIMD-ONLY1-32-LABEL: define dso_local void @_Z4foo1i(271// SIMD-ONLY1-32-SAME: i32 [[A:%.*]]) #[[ATTR0:[0-9]+]] {272// SIMD-ONLY1-32-NEXT: [[ENTRY:.*:]]273// SIMD-ONLY1-32-NEXT: [[A_ADDR:%.*]] = alloca i32, align 4274// SIMD-ONLY1-32-NEXT: [[D:%.*]] = alloca double, align 8275// SIMD-ONLY1-32-NEXT: [[D1:%.*]] = alloca double, align 8276// SIMD-ONLY1-32-NEXT: store i32 [[A]], ptr [[A_ADDR]], align 4277// SIMD-ONLY1-32-NEXT: [[TMP0:%.*]] = load i32, ptr [[A_ADDR]], align 4278// SIMD-ONLY1-32-NEXT: [[CONV:%.*]] = sitofp i32 [[TMP0]] to double279// SIMD-ONLY1-32-NEXT: store double [[CONV]], ptr [[D]], align 8280// SIMD-ONLY1-32-NEXT: [[TMP1:%.*]] = load double, ptr [[D1]], align 8281// SIMD-ONLY1-32-NEXT: [[ADD:%.*]] = fadd double [[TMP1]], 1.000000e+00282// SIMD-ONLY1-32-NEXT: store double [[ADD]], ptr [[D1]], align 8283// SIMD-ONLY1-32-NEXT: ret void284//285//286// CK2-64-LABEL: define dso_local void @_Z4foo2v(287// CK2-64-SAME: ) #[[ATTR0:[0-9]+]] {288// CK2-64-NEXT: [[ENTRY:.*:]]289// CK2-64-NEXT: [[PVTARR:%.*]] = alloca [10 x i32], align 4290// CK2-64-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x ptr], align 8291// CK2-64-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x ptr], align 8292// CK2-64-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [1 x ptr], align 8293// CK2-64-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8294// CK2-64-NEXT: [[TMP0:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0295// CK2-64-NEXT: store ptr [[PVTARR]], ptr [[TMP0]], align 8296// CK2-64-NEXT: [[TMP1:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0297// CK2-64-NEXT: store ptr [[PVTARR]], ptr [[TMP1]], align 8298// CK2-64-NEXT: [[TMP2:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 0299// CK2-64-NEXT: store ptr null, ptr [[TMP2]], align 8300// CK2-64-NEXT: [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0301// CK2-64-NEXT: [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0302// CK2-64-NEXT: [[TMP5:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0303// CK2-64-NEXT: store i32 3, ptr [[TMP5]], align 4304// CK2-64-NEXT: [[TMP6:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1305// CK2-64-NEXT: store i32 1, ptr [[TMP6]], align 4306// CK2-64-NEXT: [[TMP7:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2307// CK2-64-NEXT: store ptr [[TMP3]], ptr [[TMP7]], align 8308// CK2-64-NEXT: [[TMP8:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3309// CK2-64-NEXT: store ptr [[TMP4]], ptr [[TMP8]], align 8310// CK2-64-NEXT: [[TMP9:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4311// CK2-64-NEXT: store ptr @.offload_sizes, ptr [[TMP9]], align 8312// CK2-64-NEXT: [[TMP10:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5313// CK2-64-NEXT: store ptr @.offload_maptypes, ptr [[TMP10]], align 8314// CK2-64-NEXT: [[TMP11:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6315// CK2-64-NEXT: store ptr null, ptr [[TMP11]], align 8316// CK2-64-NEXT: [[TMP12:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7317// CK2-64-NEXT: store ptr null, ptr [[TMP12]], align 8318// CK2-64-NEXT: [[TMP13:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8319// CK2-64-NEXT: store i64 0, ptr [[TMP13]], align 8320// CK2-64-NEXT: [[TMP14:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9321// CK2-64-NEXT: store i64 0, ptr [[TMP14]], align 8322// CK2-64-NEXT: [[TMP15:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10323// CK2-64-NEXT: store [3 x i32] [i32 -1, i32 0, i32 0], ptr [[TMP15]], align 4324// CK2-64-NEXT: [[TMP16:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11325// CK2-64-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP16]], align 4326// CK2-64-NEXT: [[TMP17:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12327// CK2-64-NEXT: store i32 0, ptr [[TMP17]], align 4328// CK2-64-NEXT: [[TMP18:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1:[0-9]+]], i64 -1, i32 -1, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z4foo2v_l50.region_id, ptr [[KERNEL_ARGS]])329// CK2-64-NEXT: [[TMP19:%.*]] = icmp ne i32 [[TMP18]], 0330// CK2-64-NEXT: br i1 [[TMP19]], label %[[OMP_OFFLOAD_FAILED:.*]], label %[[OMP_OFFLOAD_CONT:.*]]331// CK2-64: [[OMP_OFFLOAD_FAILED]]:332// CK2-64-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z4foo2v_l50(ptr [[PVTARR]]) #[[ATTR2:[0-9]+]]333// CK2-64-NEXT: br label %[[OMP_OFFLOAD_CONT]]334// CK2-64: [[OMP_OFFLOAD_CONT]]:335// CK2-64-NEXT: ret void336//337//338// CK2-64-LABEL: define internal void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z4foo2v_l50(339// CK2-64-SAME: ptr nonnull align 4 dereferenceable(40) [[PVTARR:%.*]]) #[[ATTR1:[0-9]+]] {340// CK2-64-NEXT: [[ENTRY:.*:]]341// CK2-64-NEXT: [[PVTARR_ADDR:%.*]] = alloca ptr, align 8342// CK2-64-NEXT: [[PVTARR1:%.*]] = alloca [10 x i32], align 4343// CK2-64-NEXT: store ptr [[PVTARR]], ptr [[PVTARR_ADDR]], align 8344// CK2-64-NEXT: [[TMP0:%.*]] = load ptr, ptr [[PVTARR_ADDR]], align 8, !nonnull [[META5:![0-9]+]], !align [[META6:![0-9]+]]345// CK2-64-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x i32], ptr [[PVTARR1]], i64 0, i64 5346// CK2-64-NEXT: [[TMP1:%.*]] = load i32, ptr [[ARRAYIDX]], align 4347// CK2-64-NEXT: [[INC:%.*]] = add nsw i32 [[TMP1]], 1348// CK2-64-NEXT: store i32 [[INC]], ptr [[ARRAYIDX]], align 4349// CK2-64-NEXT: ret void350//351//352// CK2-32-LABEL: define dso_local void @_Z4foo2v(353// CK2-32-SAME: ) #[[ATTR0:[0-9]+]] {354// CK2-32-NEXT: [[ENTRY:.*:]]355// CK2-32-NEXT: [[PVTARR:%.*]] = alloca [10 x i32], align 4356// CK2-32-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x ptr], align 4357// CK2-32-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x ptr], align 4358// CK2-32-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [1 x ptr], align 4359// CK2-32-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8360// CK2-32-NEXT: [[TMP0:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0361// CK2-32-NEXT: store ptr [[PVTARR]], ptr [[TMP0]], align 4362// CK2-32-NEXT: [[TMP1:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0363// CK2-32-NEXT: store ptr [[PVTARR]], ptr [[TMP1]], align 4364// CK2-32-NEXT: [[TMP2:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 0365// CK2-32-NEXT: store ptr null, ptr [[TMP2]], align 4366// CK2-32-NEXT: [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0367// CK2-32-NEXT: [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0368// CK2-32-NEXT: [[TMP5:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0369// CK2-32-NEXT: store i32 3, ptr [[TMP5]], align 4370// CK2-32-NEXT: [[TMP6:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1371// CK2-32-NEXT: store i32 1, ptr [[TMP6]], align 4372// CK2-32-NEXT: [[TMP7:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2373// CK2-32-NEXT: store ptr [[TMP3]], ptr [[TMP7]], align 4374// CK2-32-NEXT: [[TMP8:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3375// CK2-32-NEXT: store ptr [[TMP4]], ptr [[TMP8]], align 4376// CK2-32-NEXT: [[TMP9:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4377// CK2-32-NEXT: store ptr @.offload_sizes, ptr [[TMP9]], align 4378// CK2-32-NEXT: [[TMP10:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5379// CK2-32-NEXT: store ptr @.offload_maptypes, ptr [[TMP10]], align 4380// CK2-32-NEXT: [[TMP11:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6381// CK2-32-NEXT: store ptr null, ptr [[TMP11]], align 4382// CK2-32-NEXT: [[TMP12:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7383// CK2-32-NEXT: store ptr null, ptr [[TMP12]], align 4384// CK2-32-NEXT: [[TMP13:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8385// CK2-32-NEXT: store i64 0, ptr [[TMP13]], align 8386// CK2-32-NEXT: [[TMP14:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9387// CK2-32-NEXT: store i64 0, ptr [[TMP14]], align 8388// CK2-32-NEXT: [[TMP15:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10389// CK2-32-NEXT: store [3 x i32] [i32 -1, i32 0, i32 0], ptr [[TMP15]], align 4390// CK2-32-NEXT: [[TMP16:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11391// CK2-32-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP16]], align 4392// CK2-32-NEXT: [[TMP17:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12393// CK2-32-NEXT: store i32 0, ptr [[TMP17]], align 4394// CK2-32-NEXT: [[TMP18:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1:[0-9]+]], i64 -1, i32 -1, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z4foo2v_l50.region_id, ptr [[KERNEL_ARGS]])395// CK2-32-NEXT: [[TMP19:%.*]] = icmp ne i32 [[TMP18]], 0396// CK2-32-NEXT: br i1 [[TMP19]], label %[[OMP_OFFLOAD_FAILED:.*]], label %[[OMP_OFFLOAD_CONT:.*]]397// CK2-32: [[OMP_OFFLOAD_FAILED]]:398// CK2-32-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z4foo2v_l50(ptr [[PVTARR]]) #[[ATTR2:[0-9]+]]399// CK2-32-NEXT: br label %[[OMP_OFFLOAD_CONT]]400// CK2-32: [[OMP_OFFLOAD_CONT]]:401// CK2-32-NEXT: ret void402//403//404// CK2-32-LABEL: define internal void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z4foo2v_l50(405// CK2-32-SAME: ptr nonnull align 4 dereferenceable(40) [[PVTARR:%.*]]) #[[ATTR1:[0-9]+]] {406// CK2-32-NEXT: [[ENTRY:.*:]]407// CK2-32-NEXT: [[PVTARR_ADDR:%.*]] = alloca ptr, align 4408// CK2-32-NEXT: [[PVTARR1:%.*]] = alloca [10 x i32], align 4409// CK2-32-NEXT: store ptr [[PVTARR]], ptr [[PVTARR_ADDR]], align 4410// CK2-32-NEXT: [[TMP0:%.*]] = load ptr, ptr [[PVTARR_ADDR]], align 4, !nonnull [[META6:![0-9]+]], !align [[META7:![0-9]+]]411// CK2-32-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x i32], ptr [[PVTARR1]], i32 0, i32 5412// CK2-32-NEXT: [[TMP1:%.*]] = load i32, ptr [[ARRAYIDX]], align 4413// CK2-32-NEXT: [[INC:%.*]] = add nsw i32 [[TMP1]], 1414// CK2-32-NEXT: store i32 [[INC]], ptr [[ARRAYIDX]], align 4415// CK2-32-NEXT: ret void416//417//418// SIMD-ONLY2-64-LABEL: define dso_local void @_Z4foo2v(419// SIMD-ONLY2-64-SAME: ) #[[ATTR0:[0-9]+]] {420// SIMD-ONLY2-64-NEXT: [[ENTRY:.*:]]421// SIMD-ONLY2-64-NEXT: [[PVTARR:%.*]] = alloca [10 x i32], align 4422// SIMD-ONLY2-64-NEXT: [[PVTARR1:%.*]] = alloca [10 x i32], align 4423// SIMD-ONLY2-64-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x i32], ptr [[PVTARR1]], i64 0, i64 5424// SIMD-ONLY2-64-NEXT: [[TMP0:%.*]] = load i32, ptr [[ARRAYIDX]], align 4425// SIMD-ONLY2-64-NEXT: [[INC:%.*]] = add nsw i32 [[TMP0]], 1426// SIMD-ONLY2-64-NEXT: store i32 [[INC]], ptr [[ARRAYIDX]], align 4427// SIMD-ONLY2-64-NEXT: ret void428//429//430// SIMD-ONLY2-32-LABEL: define dso_local void @_Z4foo2v(431// SIMD-ONLY2-32-SAME: ) #[[ATTR0:[0-9]+]] {432// SIMD-ONLY2-32-NEXT: [[ENTRY:.*:]]433// SIMD-ONLY2-32-NEXT: [[PVTARR:%.*]] = alloca [10 x i32], align 4434// SIMD-ONLY2-32-NEXT: [[PVTARR1:%.*]] = alloca [10 x i32], align 4435// SIMD-ONLY2-32-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x i32], ptr [[PVTARR1]], i32 0, i32 5436// SIMD-ONLY2-32-NEXT: [[TMP0:%.*]] = load i32, ptr [[ARRAYIDX]], align 4437// SIMD-ONLY2-32-NEXT: [[INC:%.*]] = add nsw i32 [[TMP0]], 1438// SIMD-ONLY2-32-NEXT: store i32 [[INC]], ptr [[ARRAYIDX]], align 4439// SIMD-ONLY2-32-NEXT: ret void440//441//442// CK3-64-LABEL: define dso_local void @_Z4foo3v(443// CK3-64-SAME: ) #[[ATTR0:[0-9]+]] {444// CK3-64-NEXT: [[ENTRY:.*:]]445// CK3-64-NEXT: [[PA:%.*]] = alloca ptr, align 8446// CK3-64-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x ptr], align 8447// CK3-64-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x ptr], align 8448// CK3-64-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [1 x ptr], align 8449// CK3-64-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8450// CK3-64-NEXT: [[TMP0:%.*]] = load ptr, ptr [[PA]], align 8451// CK3-64-NEXT: [[TMP1:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0452// CK3-64-NEXT: store ptr [[TMP0]], ptr [[TMP1]], align 8453// CK3-64-NEXT: [[TMP2:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0454// CK3-64-NEXT: store ptr [[TMP0]], ptr [[TMP2]], align 8455// CK3-64-NEXT: [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 0456// CK3-64-NEXT: store ptr null, ptr [[TMP3]], align 8457// CK3-64-NEXT: [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0458// CK3-64-NEXT: [[TMP5:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0459// CK3-64-NEXT: [[TMP6:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0460// CK3-64-NEXT: store i32 3, ptr [[TMP6]], align 4461// CK3-64-NEXT: [[TMP7:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1462// CK3-64-NEXT: store i32 1, ptr [[TMP7]], align 4463// CK3-64-NEXT: [[TMP8:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2464// CK3-64-NEXT: store ptr [[TMP4]], ptr [[TMP8]], align 8465// CK3-64-NEXT: [[TMP9:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3466// CK3-64-NEXT: store ptr [[TMP5]], ptr [[TMP9]], align 8467// CK3-64-NEXT: [[TMP10:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4468// CK3-64-NEXT: store ptr @.offload_sizes, ptr [[TMP10]], align 8469// CK3-64-NEXT: [[TMP11:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5470// CK3-64-NEXT: store ptr @.offload_maptypes, ptr [[TMP11]], align 8471// CK3-64-NEXT: [[TMP12:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6472// CK3-64-NEXT: store ptr null, ptr [[TMP12]], align 8473// CK3-64-NEXT: [[TMP13:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7474// CK3-64-NEXT: store ptr null, ptr [[TMP13]], align 8475// CK3-64-NEXT: [[TMP14:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8476// CK3-64-NEXT: store i64 0, ptr [[TMP14]], align 8477// CK3-64-NEXT: [[TMP15:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9478// CK3-64-NEXT: store i64 0, ptr [[TMP15]], align 8479// CK3-64-NEXT: [[TMP16:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10480// CK3-64-NEXT: store [3 x i32] [i32 -1, i32 0, i32 0], ptr [[TMP16]], align 4481// CK3-64-NEXT: [[TMP17:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11482// CK3-64-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP17]], align 4483// CK3-64-NEXT: [[TMP18:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12484// CK3-64-NEXT: store i32 0, ptr [[TMP18]], align 4485// CK3-64-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]+}}__Z4foo3v_l75.region_id, ptr [[KERNEL_ARGS]])486// CK3-64-NEXT: [[TMP20:%.*]] = icmp ne i32 [[TMP19]], 0487// CK3-64-NEXT: br i1 [[TMP20]], label %[[OMP_OFFLOAD_FAILED:.*]], label %[[OMP_OFFLOAD_CONT:.*]]488// CK3-64: [[OMP_OFFLOAD_FAILED]]:489// CK3-64-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z4foo3v_l75(ptr [[TMP0]]) #[[ATTR2:[0-9]+]]490// CK3-64-NEXT: br label %[[OMP_OFFLOAD_CONT]]491// CK3-64: [[OMP_OFFLOAD_CONT]]:492// CK3-64-NEXT: ret void493//494//495// CK3-64-LABEL: define internal void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z4foo3v_l75(496// CK3-64-SAME: ptr [[PA:%.*]]) #[[ATTR1:[0-9]+]] {497// CK3-64-NEXT: [[ENTRY:.*:]]498// CK3-64-NEXT: [[PA_ADDR:%.*]] = alloca ptr, align 8499// CK3-64-NEXT: [[PA1:%.*]] = alloca ptr, align 8500// CK3-64-NEXT: store ptr [[PA]], ptr [[PA_ADDR]], align 8501// CK3-64-NEXT: [[TMP0:%.*]] = load ptr, ptr [[PA1]], align 8502// CK3-64-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP0]], i64 50503// CK3-64-NEXT: [[TMP1:%.*]] = load i32, ptr [[ARRAYIDX]], align 4504// CK3-64-NEXT: [[INC:%.*]] = add nsw i32 [[TMP1]], 1505// CK3-64-NEXT: store i32 [[INC]], ptr [[ARRAYIDX]], align 4506// CK3-64-NEXT: ret void507//508//509// CK3-32-LABEL: define dso_local void @_Z4foo3v(510// CK3-32-SAME: ) #[[ATTR0:[0-9]+]] {511// CK3-32-NEXT: [[ENTRY:.*:]]512// CK3-32-NEXT: [[PA:%.*]] = alloca ptr, align 4513// CK3-32-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x ptr], align 4514// CK3-32-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x ptr], align 4515// CK3-32-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [1 x ptr], align 4516// CK3-32-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8517// CK3-32-NEXT: [[TMP0:%.*]] = load ptr, ptr [[PA]], align 4518// CK3-32-NEXT: [[TMP1:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0519// CK3-32-NEXT: store ptr [[TMP0]], ptr [[TMP1]], align 4520// CK3-32-NEXT: [[TMP2:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0521// CK3-32-NEXT: store ptr [[TMP0]], ptr [[TMP2]], align 4522// CK3-32-NEXT: [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 0523// CK3-32-NEXT: store ptr null, ptr [[TMP3]], align 4524// CK3-32-NEXT: [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0525// CK3-32-NEXT: [[TMP5:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0526// CK3-32-NEXT: [[TMP6:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0527// CK3-32-NEXT: store i32 3, ptr [[TMP6]], align 4528// CK3-32-NEXT: [[TMP7:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1529// CK3-32-NEXT: store i32 1, ptr [[TMP7]], align 4530// CK3-32-NEXT: [[TMP8:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2531// CK3-32-NEXT: store ptr [[TMP4]], ptr [[TMP8]], align 4532// CK3-32-NEXT: [[TMP9:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3533// CK3-32-NEXT: store ptr [[TMP5]], ptr [[TMP9]], align 4534// CK3-32-NEXT: [[TMP10:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4535// CK3-32-NEXT: store ptr @.offload_sizes, ptr [[TMP10]], align 4536// CK3-32-NEXT: [[TMP11:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5537// CK3-32-NEXT: store ptr @.offload_maptypes, ptr [[TMP11]], align 4538// CK3-32-NEXT: [[TMP12:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6539// CK3-32-NEXT: store ptr null, ptr [[TMP12]], align 4540// CK3-32-NEXT: [[TMP13:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7541// CK3-32-NEXT: store ptr null, ptr [[TMP13]], align 4542// CK3-32-NEXT: [[TMP14:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8543// CK3-32-NEXT: store i64 0, ptr [[TMP14]], align 8544// CK3-32-NEXT: [[TMP15:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9545// CK3-32-NEXT: store i64 0, ptr [[TMP15]], align 8546// CK3-32-NEXT: [[TMP16:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10547// CK3-32-NEXT: store [3 x i32] [i32 -1, i32 0, i32 0], ptr [[TMP16]], align 4548// CK3-32-NEXT: [[TMP17:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11549// CK3-32-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP17]], align 4550// CK3-32-NEXT: [[TMP18:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12551// CK3-32-NEXT: store i32 0, ptr [[TMP18]], align 4552// CK3-32-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]+}}__Z4foo3v_l75.region_id, ptr [[KERNEL_ARGS]])553// CK3-32-NEXT: [[TMP20:%.*]] = icmp ne i32 [[TMP19]], 0554// CK3-32-NEXT: br i1 [[TMP20]], label %[[OMP_OFFLOAD_FAILED:.*]], label %[[OMP_OFFLOAD_CONT:.*]]555// CK3-32: [[OMP_OFFLOAD_FAILED]]:556// CK3-32-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z4foo3v_l75(ptr [[TMP0]]) #[[ATTR2:[0-9]+]]557// CK3-32-NEXT: br label %[[OMP_OFFLOAD_CONT]]558// CK3-32: [[OMP_OFFLOAD_CONT]]:559// CK3-32-NEXT: ret void560//561//562// CK3-32-LABEL: define internal void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z4foo3v_l75(563// CK3-32-SAME: ptr [[PA:%.*]]) #[[ATTR1:[0-9]+]] {564// CK3-32-NEXT: [[ENTRY:.*:]]565// CK3-32-NEXT: [[PA_ADDR:%.*]] = alloca ptr, align 4566// CK3-32-NEXT: [[PA1:%.*]] = alloca ptr, align 4567// CK3-32-NEXT: store ptr [[PA]], ptr [[PA_ADDR]], align 4568// CK3-32-NEXT: [[TMP0:%.*]] = load ptr, ptr [[PA1]], align 4569// CK3-32-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP0]], i32 50570// CK3-32-NEXT: [[TMP1:%.*]] = load i32, ptr [[ARRAYIDX]], align 4571// CK3-32-NEXT: [[INC:%.*]] = add nsw i32 [[TMP1]], 1572// CK3-32-NEXT: store i32 [[INC]], ptr [[ARRAYIDX]], align 4573// CK3-32-NEXT: ret void574//575//576// SIMD-ONLY3-64-LABEL: define dso_local void @_Z4foo3v(577// SIMD-ONLY3-64-SAME: ) #[[ATTR0:[0-9]+]] {578// SIMD-ONLY3-64-NEXT: [[ENTRY:.*:]]579// SIMD-ONLY3-64-NEXT: [[PA:%.*]] = alloca ptr, align 8580// SIMD-ONLY3-64-NEXT: [[PA1:%.*]] = alloca ptr, align 8581// SIMD-ONLY3-64-NEXT: [[TMP0:%.*]] = load ptr, ptr [[PA1]], align 8582// SIMD-ONLY3-64-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP0]], i64 50583// SIMD-ONLY3-64-NEXT: [[TMP1:%.*]] = load i32, ptr [[ARRAYIDX]], align 4584// SIMD-ONLY3-64-NEXT: [[INC:%.*]] = add nsw i32 [[TMP1]], 1585// SIMD-ONLY3-64-NEXT: store i32 [[INC]], ptr [[ARRAYIDX]], align 4586// SIMD-ONLY3-64-NEXT: ret void587//588//589// SIMD-ONLY3-32-LABEL: define dso_local void @_Z4foo3v(590// SIMD-ONLY3-32-SAME: ) #[[ATTR0:[0-9]+]] {591// SIMD-ONLY3-32-NEXT: [[ENTRY:.*:]]592// SIMD-ONLY3-32-NEXT: [[PA:%.*]] = alloca ptr, align 4593// SIMD-ONLY3-32-NEXT: [[PA1:%.*]] = alloca ptr, align 4594// SIMD-ONLY3-32-NEXT: [[TMP0:%.*]] = load ptr, ptr [[PA1]], align 4595// SIMD-ONLY3-32-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP0]], i32 50596// SIMD-ONLY3-32-NEXT: [[TMP1:%.*]] = load i32, ptr [[ARRAYIDX]], align 4597// SIMD-ONLY3-32-NEXT: [[INC:%.*]] = add nsw i32 [[TMP1]], 1598// SIMD-ONLY3-32-NEXT: store i32 [[INC]], ptr [[ARRAYIDX]], align 4599// SIMD-ONLY3-32-NEXT: ret void600//601//602// CK4-64-LABEL: define dso_local void @_Z4foo4v(603// CK4-64-SAME: ) #[[ATTR0:[0-9]+]] {604// CK4-64-NEXT: [[ENTRY:.*:]]605// CK4-64-NEXT: [[P:%.*]] = alloca i32, align 4606// CK4-64-NEXT: [[P_CASTED:%.*]] = alloca i64, align 8607// CK4-64-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x ptr], align 8608// CK4-64-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x ptr], align 8609// CK4-64-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [1 x ptr], align 8610// CK4-64-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8611// CK4-64-NEXT: [[TMP0:%.*]] = load i32, ptr [[P]], align 4612// CK4-64-NEXT: store i32 [[TMP0]], ptr [[P_CASTED]], align 4613// CK4-64-NEXT: [[TMP1:%.*]] = load i64, ptr [[P_CASTED]], align 8614// CK4-64-NEXT: [[TMP2:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0615// CK4-64-NEXT: store i64 [[TMP1]], ptr [[TMP2]], align 8616// CK4-64-NEXT: [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0617// CK4-64-NEXT: store i64 [[TMP1]], ptr [[TMP3]], align 8618// CK4-64-NEXT: [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 0619// CK4-64-NEXT: store ptr null, ptr [[TMP4]], align 8620// CK4-64-NEXT: [[TMP5:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0621// CK4-64-NEXT: [[TMP6:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0622// CK4-64-NEXT: [[TMP7:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0623// CK4-64-NEXT: store i32 3, ptr [[TMP7]], align 4624// CK4-64-NEXT: [[TMP8:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1625// CK4-64-NEXT: store i32 1, ptr [[TMP8]], align 4626// CK4-64-NEXT: [[TMP9:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2627// CK4-64-NEXT: store ptr [[TMP5]], ptr [[TMP9]], align 8628// CK4-64-NEXT: [[TMP10:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3629// CK4-64-NEXT: store ptr [[TMP6]], ptr [[TMP10]], align 8630// CK4-64-NEXT: [[TMP11:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4631// CK4-64-NEXT: store ptr @.offload_sizes, ptr [[TMP11]], align 8632// CK4-64-NEXT: [[TMP12:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5633// CK4-64-NEXT: store ptr @.offload_maptypes, ptr [[TMP12]], align 8634// CK4-64-NEXT: [[TMP13:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6635// CK4-64-NEXT: store ptr null, ptr [[TMP13]], align 8636// CK4-64-NEXT: [[TMP14:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7637// CK4-64-NEXT: store ptr null, ptr [[TMP14]], align 8638// CK4-64-NEXT: [[TMP15:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8639// CK4-64-NEXT: store i64 0, ptr [[TMP15]], align 8640// CK4-64-NEXT: [[TMP16:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9641// CK4-64-NEXT: store i64 0, ptr [[TMP16]], align 8642// CK4-64-NEXT: [[TMP17:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10643// CK4-64-NEXT: store [3 x i32] [i32 -1, i32 0, i32 0], ptr [[TMP17]], align 4644// CK4-64-NEXT: [[TMP18:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11645// CK4-64-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP18]], align 4646// CK4-64-NEXT: [[TMP19:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12647// CK4-64-NEXT: store i32 0, ptr [[TMP19]], align 4648// CK4-64-NEXT: [[TMP20:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1:[0-9]+]], i64 -1, i32 -1, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z4foo4v_l103.region_id, ptr [[KERNEL_ARGS]])649// CK4-64-NEXT: [[TMP21:%.*]] = icmp ne i32 [[TMP20]], 0650// CK4-64-NEXT: br i1 [[TMP21]], label %[[OMP_OFFLOAD_FAILED:.*]], label %[[OMP_OFFLOAD_CONT:.*]]651// CK4-64: [[OMP_OFFLOAD_FAILED]]:652// CK4-64-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z4foo4v_l103(i64 [[TMP1]]) #[[ATTR2:[0-9]+]]653// CK4-64-NEXT: br label %[[OMP_OFFLOAD_CONT]]654// CK4-64: [[OMP_OFFLOAD_CONT]]:655// CK4-64-NEXT: ret void656//657//658// CK4-64-LABEL: define internal void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z4foo4v_l103(659// CK4-64-SAME: i64 [[P:%.*]]) #[[ATTR1:[0-9]+]] {660// CK4-64-NEXT: [[ENTRY:.*:]]661// CK4-64-NEXT: [[P_ADDR:%.*]] = alloca i64, align 8662// CK4-64-NEXT: store i64 [[P]], ptr [[P_ADDR]], align 8663// CK4-64-NEXT: [[TMP0:%.*]] = load i32, ptr [[P_ADDR]], align 4664// CK4-64-NEXT: [[INC:%.*]] = add nsw i32 [[TMP0]], 1665// CK4-64-NEXT: store i32 [[INC]], ptr [[P_ADDR]], align 4666// CK4-64-NEXT: ret void667//668//669// CK4-32-LABEL: define dso_local void @_Z4foo4v(670// CK4-32-SAME: ) #[[ATTR0:[0-9]+]] {671// CK4-32-NEXT: [[ENTRY:.*:]]672// CK4-32-NEXT: [[P:%.*]] = alloca i32, align 4673// CK4-32-NEXT: [[P_CASTED:%.*]] = alloca i32, align 4674// CK4-32-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x ptr], align 4675// CK4-32-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x ptr], align 4676// CK4-32-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [1 x ptr], align 4677// CK4-32-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8678// CK4-32-NEXT: [[TMP0:%.*]] = load i32, ptr [[P]], align 4679// CK4-32-NEXT: store i32 [[TMP0]], ptr [[P_CASTED]], align 4680// CK4-32-NEXT: [[TMP1:%.*]] = load i32, ptr [[P_CASTED]], align 4681// CK4-32-NEXT: [[TMP2:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0682// CK4-32-NEXT: store i32 [[TMP1]], ptr [[TMP2]], align 4683// CK4-32-NEXT: [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0684// CK4-32-NEXT: store i32 [[TMP1]], ptr [[TMP3]], align 4685// CK4-32-NEXT: [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 0686// CK4-32-NEXT: store ptr null, ptr [[TMP4]], align 4687// CK4-32-NEXT: [[TMP5:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0688// CK4-32-NEXT: [[TMP6:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0689// CK4-32-NEXT: [[TMP7:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0690// CK4-32-NEXT: store i32 3, ptr [[TMP7]], align 4691// CK4-32-NEXT: [[TMP8:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1692// CK4-32-NEXT: store i32 1, ptr [[TMP8]], align 4693// CK4-32-NEXT: [[TMP9:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2694// CK4-32-NEXT: store ptr [[TMP5]], ptr [[TMP9]], align 4695// CK4-32-NEXT: [[TMP10:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3696// CK4-32-NEXT: store ptr [[TMP6]], ptr [[TMP10]], align 4697// CK4-32-NEXT: [[TMP11:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4698// CK4-32-NEXT: store ptr @.offload_sizes, ptr [[TMP11]], align 4699// CK4-32-NEXT: [[TMP12:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5700// CK4-32-NEXT: store ptr @.offload_maptypes, ptr [[TMP12]], align 4701// CK4-32-NEXT: [[TMP13:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6702// CK4-32-NEXT: store ptr null, ptr [[TMP13]], align 4703// CK4-32-NEXT: [[TMP14:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7704// CK4-32-NEXT: store ptr null, ptr [[TMP14]], align 4705// CK4-32-NEXT: [[TMP15:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8706// CK4-32-NEXT: store i64 0, ptr [[TMP15]], align 8707// CK4-32-NEXT: [[TMP16:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9708// CK4-32-NEXT: store i64 0, ptr [[TMP16]], align 8709// CK4-32-NEXT: [[TMP17:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10710// CK4-32-NEXT: store [3 x i32] [i32 -1, i32 0, i32 0], ptr [[TMP17]], align 4711// CK4-32-NEXT: [[TMP18:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11712// CK4-32-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP18]], align 4713// CK4-32-NEXT: [[TMP19:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12714// CK4-32-NEXT: store i32 0, ptr [[TMP19]], align 4715// CK4-32-NEXT: [[TMP20:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1:[0-9]+]], i64 -1, i32 -1, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z4foo4v_l103.region_id, ptr [[KERNEL_ARGS]])716// CK4-32-NEXT: [[TMP21:%.*]] = icmp ne i32 [[TMP20]], 0717// CK4-32-NEXT: br i1 [[TMP21]], label %[[OMP_OFFLOAD_FAILED:.*]], label %[[OMP_OFFLOAD_CONT:.*]]718// CK4-32: [[OMP_OFFLOAD_FAILED]]:719// CK4-32-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z4foo4v_l103(i32 [[TMP1]]) #[[ATTR2:[0-9]+]]720// CK4-32-NEXT: br label %[[OMP_OFFLOAD_CONT]]721// CK4-32: [[OMP_OFFLOAD_CONT]]:722// CK4-32-NEXT: ret void723//724//725// CK4-32-LABEL: define internal void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z4foo4v_l103(726// CK4-32-SAME: i32 [[P:%.*]]) #[[ATTR1:[0-9]+]] {727// CK4-32-NEXT: [[ENTRY:.*:]]728// CK4-32-NEXT: [[P_ADDR:%.*]] = alloca i32, align 4729// CK4-32-NEXT: store i32 [[P]], ptr [[P_ADDR]], align 4730// CK4-32-NEXT: [[TMP0:%.*]] = load i32, ptr [[P_ADDR]], align 4731// CK4-32-NEXT: [[INC:%.*]] = add nsw i32 [[TMP0]], 1732// CK4-32-NEXT: store i32 [[INC]], ptr [[P_ADDR]], align 4733// CK4-32-NEXT: ret void734//735//736// SIMD-ONLY4-64-LABEL: define dso_local void @_Z4foo4v(737// SIMD-ONLY4-64-SAME: ) #[[ATTR0:[0-9]+]] {738// SIMD-ONLY4-64-NEXT: [[ENTRY:.*:]]739// SIMD-ONLY4-64-NEXT: [[P:%.*]] = alloca i32, align 4740// SIMD-ONLY4-64-NEXT: [[TMP0:%.*]] = load i32, ptr [[P]], align 4741// SIMD-ONLY4-64-NEXT: [[INC:%.*]] = add nsw i32 [[TMP0]], 1742// SIMD-ONLY4-64-NEXT: store i32 [[INC]], ptr [[P]], align 4743// SIMD-ONLY4-64-NEXT: ret void744//745//746// SIMD-ONLY4-32-LABEL: define dso_local void @_Z4foo4v(747// SIMD-ONLY4-32-SAME: ) #[[ATTR0:[0-9]+]] {748// SIMD-ONLY4-32-NEXT: [[ENTRY:.*:]]749// SIMD-ONLY4-32-NEXT: [[P:%.*]] = alloca i32, align 4750// SIMD-ONLY4-32-NEXT: [[TMP0:%.*]] = load i32, ptr [[P]], align 4751// SIMD-ONLY4-32-NEXT: [[INC:%.*]] = add nsw i32 [[TMP0]], 1752// SIMD-ONLY4-32-NEXT: store i32 [[INC]], ptr [[P]], align 4753// SIMD-ONLY4-32-NEXT: ret void754//755//.756// CK1-32: [[META6]] = !{}757// CK1-32: [[META7]] = !{i64 4}758//.759// CK2-64: [[META5]] = !{}760// CK2-64: [[META6]] = !{i64 4}761//.762// CK2-32: [[META6]] = !{}763// CK2-32: [[META7]] = !{i64 4}764//.765