895 lines · cpp
1// Test host codegen.2// RUN: %clang_cc1 -verify -Wno-vla -fopenmp -fopenmp-version=45 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-64 --check-prefix OMP453// RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s4// RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-64 --check-prefix OMP455// RUN: %clang_cc1 -verify -Wno-vla -fopenmp -fopenmp-version=45 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-32 --check-prefix OMP456// RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s7// RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-32 --check-prefix OMP458 9// RUN: %clang_cc1 -verify -Wno-vla -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-64 --check-prefix OMP5010// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s11// RUN: %clang_cc1 -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-64 --check-prefix OMP5012// RUN: %clang_cc1 -verify -Wno-vla -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-32 --check-prefix OMP5013// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s14// RUN: %clang_cc1 -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-32 --check-prefix OMP5015 16// RUN: %clang_cc1 -verify -Wno-vla -fopenmp -fopenmp-version=51 -D_DOMP51 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-64 --check-prefix OMP5117// RUN: %clang_cc1 -fopenmp -x c++ -fopenmp-version=51 -D_DOMP51 -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s18// RUN: %clang_cc1 -fopenmp -x c++ -fopenmp-version=51 -D_DOMP51 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-64 --check-prefix OMP5119// RUN: %clang_cc1 -verify -Wno-vla -fopenmp -fopenmp-version=51 -D_DOMP51 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-32 --check-prefix OMP5120// RUN: %clang_cc1 -fopenmp -fopenmp-version=51 -D_DOMP51 -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s21// RUN: %clang_cc1 -fopenmp -fopenmp-version=51 -D_DOMP51 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-32 --check-prefix OMP5122 23// RUN: %clang_cc1 -verify -Wno-vla -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s24// RUN: %clang_cc1 -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s25// RUN: %clang_cc1 -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s26// RUN: %clang_cc1 -verify -Wno-vla -fopenmp-simd -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s27// RUN: %clang_cc1 -fopenmp-simd -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s28// RUN: %clang_cc1 -fopenmp-simd -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s29// SIMD-ONLY0-NOT: {{__kmpc|__tgt}}30 31// Test target codegen - host bc file has to be created first.32// RUN: %clang_cc1 -verify -Wno-vla -fopenmp -fopenmp-version=45 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm-bc %s -o %t-ppc-host.bc33// RUN: %clang_cc1 -verify -Wno-vla -fopenmp -fopenmp-version=45 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --check-prefix TCHECK --check-prefix TCHECK-6434// RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o %t %s35// RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck %s --check-prefix TCHECK --check-prefix TCHECK-6436// RUN: %clang_cc1 -verify -Wno-vla -fopenmp -fopenmp-version=45 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm-bc %s -o %t-x86-host.bc37// RUN: %clang_cc1 -verify -Wno-vla -fopenmp -fopenmp-version=45 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s --check-prefix TCHECK --check-prefix TCHECK-3238// RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-x86-host.bc -o %t %s39// RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-x86-host.bc -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck %s --check-prefix TCHECK --check-prefix TCHECK-3240 41// RUN: %clang_cc1 -verify -Wno-vla -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm-bc %s -o %t-ppc-host.bc42// RUN: %clang_cc1 -verify -Wno-vla -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --check-prefix TCHECK --check-prefix TCHECK-6443// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o %t %s44// RUN: %clang_cc1 -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck %s --check-prefix TCHECK --check-prefix TCHECK-6445// RUN: %clang_cc1 -verify -Wno-vla -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm-bc %s -o %t-x86-host.bc46// RUN: %clang_cc1 -verify -Wno-vla -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s --check-prefix TCHECK --check-prefix TCHECK-3247// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-x86-host.bc -o %t %s48// RUN: %clang_cc1 -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-x86-host.bc -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck %s --check-prefix TCHECK --check-prefix TCHECK-3249 50// RUN: %clang_cc1 -verify -Wno-vla -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm-bc %s -o %t-ppc-host.bc51// RUN: %clang_cc1 -verify -Wno-vla -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck --check-prefix SIMD-ONLY1 %s52// RUN: %clang_cc1 -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o %t %s53// RUN: %clang_cc1 -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY1 %s54// RUN: %clang_cc1 -verify -Wno-vla -fopenmp-simd -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm-bc %s -o %t-x86-host.bc55// RUN: %clang_cc1 -verify -Wno-vla -fopenmp-simd -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck --check-prefix SIMD-ONLY1 %s56// RUN: %clang_cc1 -fopenmp-simd -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-x86-host.bc -o %t %s57// RUN: %clang_cc1 -fopenmp-simd -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-x86-host.bc -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY1 %s58// SIMD-ONLY1-NOT: {{__kmpc|__tgt}}59 60// expected-no-diagnostics61#ifndef HEADER62#define HEADER63 64// CHECK-DAG: [[IDENT_T:%.+]] = type { i32, i32, i32, i32, ptr }65// CHECK-DAG: [[KMP_TASK_T_WITH_PRIVATES:%.+]] = type { [[KMP_TASK_T:%.+]], [[KMP_PRIVATES_T:%.+]] }66// CHECK-DAG: [[KMP_TASK_T]] = type { ptr, ptr, i32, {{%.+}}, {{%.+}} }67// CHECK-DAG: [[TT:%.+]] = type { i64, i8 }68// CHECK-DAG: [[S1:%.+]] = type { double }69// CHECK-DAG: [[S2:%.+]] = type { i32, i32, i32 }70// CHECK-DAG: [[ENTTY:%.+]] = type { i64, i16, i16, i32, ptr, ptr, i64, i64, ptr }71// CHECK-DAG: [[ANON_T:%.+]] = type { ptr, i32, i32 }72// CHECK-32-DAG: [[KMP_PRIVATES_T]] = type { [2 x i64], ptr, i32, [2 x ptr], [2 x ptr] }73// CHECK-64-DAG: [[KMP_PRIVATES_T]] = type { ptr, [2 x ptr], [2 x ptr], [2 x i64], i32 }74 75// TCHECK: [[ENTTY:%.+]] = type { i64, i16, i16, i32, ptr, ptr, i64, i64, ptr }76 77// We have 9 target regions, but only 8 that actually will generate offloading78// code and have mapped arguments, and only 6 have all-constant map sizes.79 80// CHECK-DAG: [[SIZET:@.+]] = private unnamed_addr constant [2 x i64] [i64 0, i64 4]81// CHECK-DAG: [[MAPT:@.+]] = private unnamed_addr constant [2 x i64] [i64 544, i64 800]82// CHECK-DAG: [[SIZET2:@.+]] = private unnamed_addr constant [1 x i{{32|64}}] [i64 2]83// CHECK-DAG: [[MAPT2:@.+]] = private unnamed_addr constant [1 x i64] [i64 800]84// CHECK-DAG: [[SIZET3:@.+]] = private unnamed_addr constant [2 x i64] [i64 4, i64 2]85// CHECK-DAG: [[MAPT3:@.+]] = private unnamed_addr constant [2 x i64] [i64 800, i64 800]86// CHECK-DAG: [[SIZET4:@.+]] = private unnamed_addr constant [9 x i64] [i64 4, i64 40, i64 {{4|8}}, i64 0, i64 400, i64 {{4|8}}, i64 {{4|8}}, i64 0, i64 {{12|16}}]87// CHECK-DAG: [[MAPT4:@.+]] = private unnamed_addr constant [9 x i64] [i64 800, i64 547, i64 800, i64 547, i64 547, i64 800, i64 800, i64 547, i64 547]88// CHECK-DAG: [[SIZET5:@.+]] = private unnamed_addr constant [3 x i64] [i64 4, i64 2, i64 40]89// CHECK-DAG: [[MAPT5:@.+]] = private unnamed_addr constant [3 x i64] [i64 800, i64 800, i64 547]90// CHECK-DAG: [[SIZET6:@.+]] = private unnamed_addr constant [4 x i64] [i64 4, i64 2, i64 1, i64 40]91// CHECK-DAG: [[MAPT6:@.+]] = private unnamed_addr constant [4 x i64] [i64 800, i64 800, i64 800, i64 547]92// CHECK-DAG: [[SIZET7:@.+]] = private unnamed_addr constant [5 x i64] [i64 8, i64 4, i64 {{4|8}}, i64 {{4|8}}, i64 0]93// CHECK-DAG: [[MAPT7:@.+]] = private unnamed_addr constant [5 x i64] [i64 547, i64 800, i64 800, i64 800, i64 547]94// CHECK-DAG: [[SIZET9:@.+]] = private unnamed_addr constant [1 x i64] [i64 12]95// CHECK-DAG: [[MAPT10:@.+]] = private unnamed_addr constant [1 x i64] [i64 35]96// CHECK-DAG: @{{.*}} = weak constant i8 097// CHECK-DAG: @{{.*}} = weak constant i8 098// CHECK-DAG: @{{.*}} = weak constant i8 099// CHECK-DAG: @{{.*}} = weak constant i8 0100// CHECK-DAG: @{{.*}} = weak constant i8 0101// CHECK-DAG: @{{.*}} = weak constant i8 0102// CHECK-DAG: @{{.*}} = weak constant i8 0103// CHECK-DAG: @{{.*}} = weak constant i8 0104 105// TCHECK: @{{.+}} = weak constant [[ENTTY]]106// TCHECK: @{{.+}} = weak constant [[ENTTY]]107// TCHECK: @{{.+}} = weak constant [[ENTTY]]108// TCHECK: @{{.+}} = weak constant [[ENTTY]]109// TCHECK: @{{.+}} = weak constant [[ENTTY]]110// TCHECK: @{{.+}} = weak constant [[ENTTY]]111// TCHECK: @{{.+}} = weak constant [[ENTTY]]112// TCHECK: @{{.+}} = weak constant [[ENTTY]]113// TCHECK: @{{.+}} = weak constant [[ENTTY]]114// TCHECK: @{{.+}} = weak constant [[ENTTY]]115// TCHECK-NOT: @{{.+}} = weak constant [[ENTTY]]116 117template<typename tx, typename ty>118struct TT{119 tx X;120 ty Y;121};122 123int global;124extern int global;125 126// CHECK: define {{.*}}[[FOO:@.+]](127int foo(int n) {128 // CHECK: [[OFFLOADBPTR:%.+]] = alloca [2 x ptr], align129 // CHECK: [[OFFLOADPTR:%.+]] = alloca [2 x ptr], align130 // CHECK: [[OFFLOADMAPPER:%.+]] = alloca [2 x ptr], align131 int a = 0;132 short aa = 0;133 float b[10];134 float bn[n];135 double c[5][10];136 double cn[5][n];137 TT<long long, char> d;138 static long *plocal;139 140// CHECK: [[ADD:%.+]] = add nsw i32141// CHECK: store i32 [[ADD]], ptr [[DEVICE_CAP:%.+]],142// CHECK: [[DEV:%.+]] = load i32, ptr [[DEVICE_CAP]],143// CHECK: [[DEVICE:%.+]] = sext i32 [[DEV]] to i64144// CHECK: [[RET:%.+]] = call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 [[DEVICE]], i32 {{.+}}, i32 {{.+}}, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])145// CHECK-NEXT: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0146// CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:[^,]+]], label %[[END:[^,]+]]147// CHECK: [[FAIL]]148// CHECK: call void [[HVT0:@.+]]()149// CHECK-NEXT: br label %[[END]]150// CHECK: [[END]]151#pragma omp target device(global + a)152 {153 }154 155 // CHECK: [[BPRGEP:%.+]] = getelementptr inbounds [2 x ptr], ptr [[OFFLOADBPTR]], i32 0, i32 0156 // CHECK: [[PRGEP:%.+]] = getelementptr inbounds [2 x ptr], ptr [[OFFLOADPTR]], i32 0, i32 0157 // CHECK: [[BPRGEP:%.+]] = getelementptr inbounds [2 x ptr], ptr [[OFFLOADBPTR]], i32 0, i32 0158 // CHECK: [[PRGEP:%.+]] = getelementptr inbounds [2 x ptr], ptr [[OFFLOADPTR]], i32 0, i32 0159 // CHECK: [[DEVICE:%.+]] = sext i32 {{%.+}} to i64160 // CHECK-32: [[TASK:%.+]] = call ptr @__kmpc_omp_target_task_alloc(ptr {{.+}}, i32 %0, i32 1, i32 60, i32 12, ptr [[OMP_TASK_ENTRY:@.+]], i64 [[DEVICE]])161 // CHECK-64: [[TASK:%.+]] = call ptr @__kmpc_omp_target_task_alloc(ptr {{.+}}, i32 %0, i32 1, i64 104, i64 16, ptr [[OMP_TASK_ENTRY:@.+]], i64 [[DEVICE]])162 // CHECK: [[TASK_WITH_PRIVATES_GEP:%.+]] = getelementptr inbounds nuw [[KMP_TASK_T_WITH_PRIVATES]], ptr [[TASK]], i32 0, i32 1163 // CHECK-32: [[SIZEGEP:%.+]] = getelementptr inbounds nuw [[KMP_PRIVATES_T]], ptr [[TASK_WITH_PRIVATES_GEP]], i32 0, i32 0164 // CHECK-32: call void @llvm.memcpy.p0.p0.i32(ptr align 4 [[SIZEGEP]], ptr align 4 [[SIZET]], i32 16, i1 false)165 // CHECK-32: [[FPBPRGEP:%.+]] = getelementptr inbounds nuw [[KMP_PRIVATES_T]], ptr [[TASK_WITH_PRIVATES_GEP]], i32 0, i32 3166 // CHECK-32: call void @llvm.memcpy.p0.p0.i32(ptr align 4 [[FPBPRGEP]], ptr align 4 [[BPRGEP]], i32 8, i1 false)167 // CHECK-32: [[FPPRGEP:%.+]] = getelementptr inbounds nuw [[KMP_PRIVATES_T]], ptr [[TASK_WITH_PRIVATES_GEP]], i32 0, i32 4168 // CHECK-32: call void @llvm.memcpy.p0.p0.i32(ptr align 4 [[FPPRGEP]], ptr align 4 [[PRGEP]], i32 8, i1 false)169 // CHECK-64: [[FPBPRGEP:%.+]] = getelementptr inbounds nuw [[KMP_PRIVATES_T]], ptr [[TASK_WITH_PRIVATES_GEP]], i32 0, i32 1170 // CHECK-64: call void @llvm.memcpy.p0.p0.i64(ptr align 8 [[FPBPRGEP]], ptr align 8 [[BPRGEP]], i64 16, i1 false)171 // CHECK-64: [[FPPRGEP:%.+]] = getelementptr inbounds nuw [[KMP_PRIVATES_T]], ptr [[TASK_WITH_PRIVATES_GEP]], i32 0, i32 2172 // CHECK-64: call void @llvm.memcpy.p0.p0.i64(ptr align 8 [[FPPRGEP]], ptr align 8 [[PRGEP]], i64 16, i1 false)173 // CHECK-64: [[SIZEGEP:%.+]] = getelementptr inbounds nuw [[KMP_PRIVATES_T]], ptr [[TASK_WITH_PRIVATES_GEP]], i32 0, i32 3174 // CHECK-64: call void @llvm.memcpy.p0.p0.i64(ptr align 8 [[SIZEGEP]], ptr align 8 [[SIZET]], i64 16, i1 false)175 // CHECK: call i32 @__kmpc_omp_task(ptr {{.+}}, i32 {{.+}}, ptr [[TASK]])176 #pragma omp target device(global + a) nowait177 {178 static int local1;179 *plocal = global;180 local1 = global;181 }182 183 // CHECK: call void [[HVT1:@.+]](i[[SZ:32|64]] {{[^,]+}})184 #pragma omp target if(0) firstprivate(global)185 {186 global += 1;187 }188 189// CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 {{.+}}, i32 {{.+}}, i32 {{.+}}, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])190// CHECK-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2191// CHECK-DAG: store ptr [[BP:%.+]], ptr [[BPARG]]192// CHECK-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3193// CHECK-DAG: store ptr [[P:%.+]], ptr [[PARG]]194// CHECK-DAG: [[BP]] = getelementptr inbounds [1 x ptr], ptr [[BPR:%[^,]+]], i32 0, i32 0195// CHECK-DAG: [[P]] = getelementptr inbounds [1 x ptr], ptr [[PR:%[^,]+]], i32 0, i32 0196// CHECK-DAG: [[BPADDR0:%.+]] = getelementptr inbounds [1 x ptr], ptr [[BPR]], i32 0, i32 [[IDX0:[0-9]+]]197// CHECK-DAG: [[PADDR0:%.+]] = getelementptr inbounds [1 x ptr], ptr [[PR]], i32 0, i32 [[IDX0]]198// CHECK-DAG: store i[[SZ]] [[BP0:%[^,]+]], ptr [[BPADDR0]]199// CHECK-DAG: store i[[SZ]] [[P0:%[^,]+]], ptr [[PADDR0]]200 201// CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0202// CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:[^,]+]], label %[[END:[^,]+]]203// CHECK: [[FAIL]]204// CHECK: call void [[HVT2:@.+]](i[[SZ]] {{[^,]+}})205// CHECK-NEXT: br label %[[END]]206// CHECK: [[END]]207#pragma omp target if (1)208 {209 aa += 1;210 }211 212// CHECK: [[IF:%.+]] = icmp sgt i32 {{[^,]+}}, 10213// CHECK: br i1 [[IF]], label %[[IFTHEN:[^,]+]], label %[[IFELSE:[^,]+]]214// CHECK: [[IFTHEN]]215// CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 {{.+}}, i32 {{.+}}, i32 {{.+}}, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])216// CHECK-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2217// CHECK-DAG: store ptr [[BPR:%.+]], ptr [[BPARG]]218// CHECK-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3219// CHECK-DAG: store ptr [[PR:%.+]], ptr [[PARG]]220// CHECK-DAG: [[BPR]] = getelementptr inbounds [2 x ptr], ptr [[BP:%[^,]+]], i32 0, i32 0221// CHECK-DAG: [[PR]] = getelementptr inbounds [2 x ptr], ptr [[P:%[^,]+]], i32 0, i32 0222 223// CHECK-DAG: [[BPADDR0:%.+]] = getelementptr inbounds [2 x ptr], ptr [[BP]], i32 0, i32 0224// CHECK-DAG: [[PADDR0:%.+]] = getelementptr inbounds [2 x ptr], ptr [[P]], i32 0, i32 0225// CHECK-DAG: store i[[SZ]] [[BP0:%[^,]+]], ptr [[BPADDR0]]226// CHECK-DAG: store i[[SZ]] [[P0:%[^,]+]], ptr [[PADDR0]]227 228// CHECK-DAG: [[BPADDR1:%.+]] = getelementptr inbounds [2 x ptr], ptr [[BP]], i32 0, i32 1229// CHECK-DAG: [[PADDR1:%.+]] = getelementptr inbounds [2 x ptr], ptr [[P]], i32 0, i32 1230// CHECK-DAG: store i[[SZ]] [[BP1:%[^,]+]], ptr [[BPADDR1]]231// CHECK-DAG: store i[[SZ]] [[P1:%[^,]+]], ptr [[PADDR1]]232// CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0233// CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]]234// CHECK: [[FAIL]]235// CHECK: call void [[HVT3:@.+]]({{[^,]+}}, {{[^,]+}})236// CHECK-NEXT: br label %[[END]]237// CHECK: [[END]]238// CHECK-NEXT: br label %[[IFEND:.+]]239// CHECK: [[IFELSE]]240// CHECK: call void [[HVT3]]({{[^,]+}}, {{[^,]+}})241// CHECK-NEXT: br label %[[IFEND]]242 243// CHECK: [[IFEND]]244#pragma omp target if (n > 10)245 {246 a += 1;247 aa += 1;248 }249 250 // We capture 3 VLA sizes in this target region251 // CHECK-64: [[A_VAL:%.+]] = load i32, ptr %{{.+}},252 // CHECK-64: store i32 [[A_VAL]], ptr [[A_CADDR:%.+]],253 // CHECK-64: [[A_CVAL:%.+]] = load i[[SZ]], ptr [[A_CADDR]],254 255 // CHECK-32: [[A_VAL:%.+]] = load i32, ptr %{{.+}},256 // CHECK-32: store i32 [[A_VAL]], ptr [[A_CADDR:%.+]],257 // CHECK-32: [[A_CVAL:%.+]] = load i[[SZ]], ptr [[A_CADDR]],258 259 // CHECK: [[IF:%.+]] = icmp sgt i32 {{[^,]+}}, 20260 // CHECK: br i1 [[IF]], label %[[TRY:[^,]+]], label %[[IFELSE:[^,]+]]261 // CHECK: [[TRY]]262 // CHECK-64: [[BNSIZE:%.+]] = mul nuw i64 [[VLA0:%.+]], 4263 // CHECK-32: [[BNSZSIZE:%.+]] = mul nuw i32 [[VLA0:%.+]], 4264 // CHECK-32: [[BNSIZE:%.+]] = sext i32 [[BNSZSIZE]] to i64265 // CHECK: [[CNELEMSIZE2:%.+]] = mul nuw i[[SZ]] 5, [[VLA1:%.+]]266 // CHECK-64: [[CNSIZE:%.+]] = mul nuw i64 [[CNELEMSIZE2]], 8267 // CHECK-32: [[CNSZSIZE:%.+]] = mul nuw i32 [[CNELEMSIZE2]], 8268 // CHECK-32: [[CNSIZE:%.+]] = sext i32 [[CNSZSIZE]] to i64269 270// CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 {{.+}}, i32 {{.+}}, i32 {{.+}}, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])271// CHECK-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2272// CHECK-DAG: store ptr [[BPR:%.+]], ptr [[BPARG]]273// CHECK-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3274// CHECK-DAG: store ptr [[PR:%.+]], ptr [[PARG]]275// CHECK-DAG: [[SARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 4276// CHECK-DAG: store ptr [[SZ4:%.+]], ptr [[SARG]]277// CHECK-DAG: [[BPR]] = getelementptr inbounds [9 x ptr], ptr [[BP:%[^,]+]], i32 0, i32 0278// CHECK-DAG: [[PR]] = getelementptr inbounds [9 x ptr], ptr [[P:%[^,]+]], i32 0, i32 0279// CHECK-DAG: [[SZ4]] = getelementptr inbounds [9 x i64], ptr [[PSZ:%[^,]+]], i32 0, i32 0280 281// CHECK-DAG: [[BPADDR0:%.+]] = getelementptr inbounds [9 x ptr], ptr [[BP]], i32 0, i32 [[IDX0:0]]282// CHECK-DAG: [[PADDR0:%.+]] = getelementptr inbounds [9 x ptr], ptr [[P]], i32 0, i32 [[IDX0]]283// CHECK-DAG: [[BPADDR1:%.+]] = getelementptr inbounds [9 x ptr], ptr [[BP]], i32 0, i32 [[IDX1:1]]284// CHECK-DAG: [[PADDR1:%.+]] = getelementptr inbounds [9 x ptr], ptr [[P]], i32 0, i32 [[IDX1]]285// CHECK-DAG: [[BPADDR2:%.+]] = getelementptr inbounds [9 x ptr], ptr [[BP]], i32 0, i32 [[IDX2:2]]286// CHECK-DAG: [[PADDR2:%.+]] = getelementptr inbounds [9 x ptr], ptr [[P]], i32 0, i32 [[IDX2]]287// CHECK-DAG: [[BPADDR3:%.+]] = getelementptr inbounds [9 x ptr], ptr [[BP]], i32 0, i32 [[IDX3:3]]288// CHECK-DAG: [[PADDR3:%.+]] = getelementptr inbounds [9 x ptr], ptr [[P]], i32 0, i32 [[IDX3]]289// CHECK-DAG: [[PSZ3:%.+]] = getelementptr inbounds [9 x i64], ptr [[PSZ]], i32 0, i32 [[IDX3]]290// CHECK-DAG: [[BPADDR4:%.+]] = getelementptr inbounds [9 x ptr], ptr [[BP]], i32 0, i32 [[IDX4:4]]291// CHECK-DAG: [[PADDR4:%.+]] = getelementptr inbounds [9 x ptr], ptr [[P]], i32 0, i32 [[IDX4]]292// CHECK-DAG: [[BPADDR5:%.+]] = getelementptr inbounds [9 x ptr], ptr [[BP]], i32 0, i32 [[IDX5:5]]293// CHECK-DAG: [[PADDR5:%.+]] = getelementptr inbounds [9 x ptr], ptr [[P]], i32 0, i32 [[IDX5]]294// CHECK-DAG: [[BPADDR6:%.+]] = getelementptr inbounds [9 x ptr], ptr [[BP]], i32 0, i32 [[IDX6:6]]295// CHECK-DAG: [[PADDR6:%.+]] = getelementptr inbounds [9 x ptr], ptr [[P]], i32 0, i32 [[IDX6]]296// CHECK-DAG: [[BPADDR7:%.+]] = getelementptr inbounds [9 x ptr], ptr [[BP]], i32 0, i32 [[IDX7:7]]297// CHECK-DAG: [[PADDR7:%.+]] = getelementptr inbounds [9 x ptr], ptr [[P]], i32 0, i32 [[IDX7]]298// CHECK-DAG: [[PSZ7:%.+]] = getelementptr inbounds [9 x i64], ptr [[PSZ]], i32 0, i32 [[IDX7]]299// CHECK-DAG: [[BPADDR8:%.+]] = getelementptr inbounds [9 x ptr], ptr [[BP]], i32 0, i32 [[IDX8:8]]300// CHECK-DAG: [[PADDR8:%.+]] = getelementptr inbounds [9 x ptr], ptr [[P]], i32 0, i32 [[IDX8]]301 302// The names below are not necessarily consistent with the names used for the303// addresses above as some are repeated.304// CHECK-DAG: store i[[SZ]] [[VLA0]], ptr [[BPADDR2]]305// CHECK-DAG: store i[[SZ]] [[VLA0]], ptr [[PADDR2]]306 307// CHECK-DAG: store i[[SZ]] [[VLA1]], ptr [[BPADDR6]]308// CHECK-DAG: store i[[SZ]] [[VLA1]], ptr [[PADDR6]]309 310// CHECK-DAG: store i[[SZ]] 5, ptr [[BPADDR5]]311// CHECK-DAG: store i[[SZ]] 5, ptr [[PADDR5]]312 313// CHECK-DAG: store i[[SZ]] [[A_CVAL]], ptr [[BPADDR0]]314// CHECK-DAG: store i[[SZ]] [[A_CVAL]], ptr [[PADDR0]]315 316// CHECK-DAG: store ptr %{{.+}}, ptr [[BPADDR1]]317// CHECK-DAG: store ptr %{{.+}}, ptr [[PADDR1]]318 319// CHECK-DAG: store ptr %{{.+}}, ptr [[BPADDR3]]320// CHECK-DAG: store ptr %{{.+}}, ptr [[PADDR3]]321// CHECK-DAG: store i64 [[BNSIZE]], ptr [[PSZ3]]322 323// CHECK-DAG: store ptr %{{.+}}, ptr [[BPADDR4]]324// CHECK-DAG: store ptr %{{.+}}, ptr [[PADDR4]]325 326// CHECK-DAG: store ptr %{{.+}}, ptr [[BPADDR7]]327// CHECK-DAG: store ptr %{{.+}}, ptr [[PADDR7]]328// CHECK-DAG: store i64 [[CNSIZE]], ptr [[PSZ7]]329 330// CHECK-DAG: store ptr %{{.+}}, ptr [[BPADDR8]]331// CHECK-DAG: store ptr %{{.+}}, ptr [[PADDR8]]332 333// CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0334// CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]]335// CHECK: [[FAIL]]336// CHECK: call void [[HVT4:@.+]]({{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}})337// CHECK-NEXT: br label %[[END]]338// CHECK: [[END]]339// CHECK-NEXT: br label %[[IFEND:.+]]340// CHECK: [[IFELSE]]341// CHECK: call void [[HVT4]]({{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}})342// CHECK-NEXT: br label %[[IFEND]]343 344// CHECK: [[IFEND]]345#pragma omp target if (n > 20)346 {347 a += 1;348 b[2] += 1.0;349 bn[3] += 1.0;350 c[1][2] += 1.0;351 cn[1][3] += 1.0;352 d.X += 1;353 d.Y += 1;354 }355 356 return a;357}358 359// Check that the offloading functions are emitted and that the arguments are360// correct and loaded correctly for the target regions in foo().361 362// CHECK: define internal void [[HVT0]]()363 364// CHECK: define internal void [[HVT0_:@.+]](ptr noundef {{%[^,]+}}, i[[SZ]] noundef {{%[^,]+}})365// CHECK: define internal {{.*}}i32 [[OMP_TASK_ENTRY]](i32 {{.*}}%0, ptr noalias noundef %1)366// CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 [[DEVICE:%.+]], i32 {{.+}}, i32 {{.+}}, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])367// CHECK-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2368// CHECK-DAG: store ptr [[BPR:%.+]], ptr [[BPARG]]369// CHECK-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3370// CHECK-DAG: store ptr [[PR:%.+]], ptr [[PARG]]371// CHECK-DAG: [[SARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 4372// CHECK-DAG: store ptr [[SIZE:%.+]], ptr [[SARG]]373// CHECK-DAG: [[DEVICE]] = sext i32 [[DEV:%.+]] to i64374// CHECK-DAG: [[DEV]] = load i32, ptr [[DEVADDR:%.+]], align375// CHECK-DAG: [[DEVADDR]] = getelementptr inbounds nuw [[ANON_T]], ptr {{%.+}}, i32 0, i32 2376// CHECK-DAG: [[BPR]] = load ptr, ptr [[FPPTR_BPR:%.+]], align377// CHECK-DAG: [[PR]] = load ptr, ptr [[FPPTR_PR:%.+]], align378// CHECK-DAG: [[SIZE]] = load ptr, ptr [[FPPTR_SIZE:%.+]], align379// CHECK-DAG: call void {{%[0-9]+}}(ptr {{%[^,]+}}, ptr [[FPPTR_PLOCAL:%.+]], ptr [[FPPTR_GLOBAL:%.+]], ptr [[FPPTR_BPR]], ptr [[FPPTR_PR]], ptr [[FPPTR_SIZE]])380// CHECK-DAG: [[PLOCALADDR:%.+]] = load ptr, ptr [[FPPTR_PLOCAL]], align381// CHECK-DAG: {{%.+}} = load ptr, ptr [[FPPTR_GLOBAL:%.+]], align382// CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0383// CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:[^,]+]], label %[[END:[^,]+]]384// CHECK: [[FAIL]]385// CHECK: [[PLOCAL:%.+]] = load ptr, ptr [[PLOCALADDR]], align386// CHECK: [[GLOBAL:%.+]] = load i32, ptr {{@.+}}, align387// CHECK-32: store i32 [[GLOBAL]], ptr [[GLOBALCAST:%.+]], align388// CHECK-64: store i32 [[GLOBAL]], ptr [[GLOBALCAST:%.+]], align389// CHECK: [[GLOBAL:%.+]] = load i[[SZ]], ptr [[GLOBALCAST]], align390// CHECK: call void [[HVT0_]](ptr [[PLOCAL]], i[[SZ]] [[GLOBAL]])391// CHECK-NEXT: br label %[[END]]392// CHECK: [[END]]393 394// CHECK: define internal void [[HVT1]](i[[SZ]] noundef %{{.+}})395// Create stack storage and store argument in there.396// CHECK: [[AA_ADDR:%.+]] = alloca i[[SZ]], align397// CHECK: store i[[SZ]] %{{.+}}, ptr [[AA_ADDR]], align398// CHECK-64: load i32, ptr [[AA_ADDR]], align399// CHECK-32: load i32, ptr [[AA_ADDR]], align400 401// CHECK: define internal void [[HVT2]](i[[SZ]] noundef %{{.+}})402// Create stack storage and store argument in there.403// CHECK: [[AA_ADDR:%.+]] = alloca i[[SZ]], align404// CHECK: store i[[SZ]] %{{.+}}, ptr [[AA_ADDR]], align405// CHECK: load i16, ptr [[AA_ADDR]], align406 407// CHECK: define internal void [[HVT3]]408// Create stack storage and store argument in there.409// CHECK: [[A_ADDR:%.+]] = alloca i[[SZ]], align410// CHECK: [[AA_ADDR:%.+]] = alloca i[[SZ]], align411// CHECK-DAG: store i[[SZ]] %{{.+}}, ptr [[A_ADDR]], align412// CHECK-DAG: store i[[SZ]] %{{.+}}, ptr [[AA_ADDR]], align413// CHECK-64-DAG:load i32, ptr [[A_ADDR]], align414// CHECK-32-DAG:load i32, ptr [[A_ADDR]], align415// CHECK-DAG: load i16, ptr [[AA_ADDR]], align416 417// CHECK: define internal void [[HVT4]]418// Create local storage for each capture.419// CHECK: [[LOCAL_A:%.+]] = alloca i[[SZ]]420// CHECK: [[LOCAL_B:%.+]] = alloca ptr421// CHECK: [[LOCAL_VLA1:%.+]] = alloca i[[SZ]]422// CHECK: [[LOCAL_BN:%.+]] = alloca ptr423// CHECK: [[LOCAL_C:%.+]] = alloca ptr424// CHECK: [[LOCAL_VLA2:%.+]] = alloca i[[SZ]]425// CHECK: [[LOCAL_VLA3:%.+]] = alloca i[[SZ]]426// CHECK: [[LOCAL_CN:%.+]] = alloca ptr427// CHECK: [[LOCAL_D:%.+]] = alloca ptr428// CHECK-DAG: store i[[SZ]] [[ARG_A:%.+]], ptr [[LOCAL_A]]429// CHECK-DAG: store ptr [[ARG_B:%.+]], ptr [[LOCAL_B]]430// CHECK-DAG: store i[[SZ]] [[ARG_VLA1:%.+]], ptr [[LOCAL_VLA1]]431// CHECK-DAG: store ptr [[ARG_BN:%.+]], ptr [[LOCAL_BN]]432// CHECK-DAG: store ptr [[ARG_C:%.+]], ptr [[LOCAL_C]]433// CHECK-DAG: store i[[SZ]] [[ARG_VLA2:%.+]], ptr [[LOCAL_VLA2]]434// CHECK-DAG: store i[[SZ]] [[ARG_VLA3:%.+]], ptr [[LOCAL_VLA3]]435// CHECK-DAG: store ptr [[ARG_CN:%.+]], ptr [[LOCAL_CN]]436// CHECK-DAG: store ptr [[ARG_D:%.+]], ptr [[LOCAL_D]]437 438// CHECK-DAG: [[REF_B:%.+]] = load ptr, ptr [[LOCAL_B]],439// CHECK-DAG: [[VAL_VLA1:%.+]] = load i[[SZ]], ptr [[LOCAL_VLA1]],440// CHECK-DAG: [[REF_BN:%.+]] = load ptr, ptr [[LOCAL_BN]],441// CHECK-DAG: [[REF_C:%.+]] = load ptr, ptr [[LOCAL_C]],442// CHECK-DAG: [[VAL_VLA2:%.+]] = load i[[SZ]], ptr [[LOCAL_VLA2]],443// CHECK-DAG: [[VAL_VLA3:%.+]] = load i[[SZ]], ptr [[LOCAL_VLA3]],444// CHECK-DAG: [[REF_CN:%.+]] = load ptr, ptr [[LOCAL_CN]],445// CHECK-DAG: [[REF_D:%.+]] = load ptr, ptr [[LOCAL_D]],446 447// Use captures.448// CHECK-64-DAG: load i32, ptr [[LOCAL_A]]449// CHECK-32-DAG: load i32, ptr [[LOCAL_A]]450// CHECK-DAG: getelementptr inbounds [10 x float], ptr [[REF_B]], i[[SZ]] 0, i[[SZ]] 2451// CHECK-DAG: getelementptr inbounds float, ptr [[REF_BN]], i[[SZ]] 3452// CHECK-DAG: getelementptr inbounds [5 x [10 x double]], ptr [[REF_C]], i[[SZ]] 0, i[[SZ]] 1453// CHECK-DAG: getelementptr inbounds double, ptr [[REF_CN]], i[[SZ]] %{{.+}}454// CHECK-DAG: getelementptr inbounds nuw [[TT]], ptr [[REF_D]], i32 0, i32 0455 456template<typename tx>457tx ftemplate(int n) {458 tx a = 0;459 short aa = 0;460 tx b[10];461 462 #pragma omp target if(n>40)463 {464 a += 1;465 aa += 1;466 b[2] += 1;467 }468 469 return a;470}471 472static473int fstatic(int n) {474 int a = 0;475 short aa = 0;476 char aaa = 0;477 int b[10];478 479 #pragma omp target if(n>50)480 {481 a += 1;482 aa += 1;483 aaa += 1;484 b[2] += 1;485 }486 487 return a;488}489 490struct S1 {491 double a;492 493 int r1(int n){494 int b = n+1;495 short int c[2][n];496 497 #pragma omp target if(n>60)498 {499 this->a = (double)b + 1.5;500 c[1][1] = ++a;501 }502 503 return c[1][1] + (int)b;504 }505};506 507// CHECK: define {{.*}}@{{.*}}bar{{.*}}508int bar(int n){509 int a = 0;510 511 // CHECK: call {{.*}}i32 [[FOO]](i32 {{.*}})512 a += foo(n);513 514 S1 S;515 // CHECK: call {{.*}}i32 [[FS1:@.+]](ptr {{.*}}, i32 {{.*}})516 a += S.r1(n);517 518 // CHECK: call {{.*}}i32 [[FSTATIC:@.+]](i32 {{.*}})519 a += fstatic(n);520 521 // CHECK: call {{.*}}i32 [[FTEMPLATE:@.+]](i32 {{.*}})522 a += ftemplate<int>(n);523 524 return a;525}526 527//528// CHECK: define {{.*}}[[FS1]]529//530// CHECK: ptr @llvm.stacksave.p0()531// CHECK-64: store i32 %{{.+}}, ptr [[B_CADDR:%.+]],532// CHECK-64: [[B_CVAL:%.+]] = load i[[SZ]], ptr [[B_CADDR]],533 534// CHECK-32: store i32 %{{.+}}, ptr %__vla_expr535// CHECK-32: store i32 %{{.+}}, ptr [[B_CADDR:%.+]],536// CHECK-32: [[B_CVAL:%.+]] = load i[[SZ]], ptr [[B_CADDR]],537 538// CHECK: [[IF:%.+]] = icmp sgt i32 {{[^,]+}}, 60539// CHECK: br i1 [[IF]], label %[[TRY:[^,]+]], label %[[IFELSE:[^,]+]]540// CHECK: [[TRY]]541// We capture 2 VLA sizes in this target region542// CHECK: [[CELEMSIZE2:%.+]] = mul nuw i[[SZ]] 2, [[VLA0:%.+]]543// CHECK-64: [[CSIZE:%.+]] = mul nuw i64 [[CELEMSIZE2]], 2544// CHECK-32: [[CSZSIZE:%.+]] = mul nuw i32 [[CELEMSIZE2]], 2545// CHECK-32: [[CSIZE:%.+]] = sext i32 [[CSZSIZE]] to i64546 547// CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 {{.+}}, i32 {{.+}}, i32 {{.+}}, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])548// CHECK-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2549// CHECK-DAG: store ptr [[BPR:%.+]], ptr [[BPARG]]550// CHECK-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3551// CHECK-DAG: store ptr [[PR:%.+]], ptr [[PARG]]552// CHECK-DAG: [[SARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 4553// CHECK-DAG: store ptr [[SZ7:%.+]], ptr [[SARG]]554// CHECK-DAG: [[BPR]] = getelementptr inbounds [5 x ptr], ptr [[BP:%.+]], i32 0, i32 0555// CHECK-DAG: [[PR]] = getelementptr inbounds [5 x ptr], ptr [[P:%.+]], i32 0, i32 0556// CHECK-DAG: [[SZ7]] = getelementptr inbounds [5 x i64], ptr [[PSZ:%.+]], i32 0, i32 0557// CHECK-DAG: [[BPADDR0:%.+]] = getelementptr inbounds [5 x ptr], ptr [[BP]], i32 0, i32 [[IDX0:0]]558// CHECK-DAG: [[PADDR0:%.+]] = getelementptr inbounds [5 x ptr], ptr [[P]], i32 0, i32 [[IDX0]]559// CHECK-DAG: [[BPADDR1:%.+]] = getelementptr inbounds [5 x ptr], ptr [[BP]], i32 0, i32 [[IDX1:1]]560// CHECK-DAG: [[PADDR1:%.+]] = getelementptr inbounds [5 x ptr], ptr [[P]], i32 0, i32 [[IDX1]]561// CHECK-DAG: [[BPADDR2:%.+]] = getelementptr inbounds [5 x ptr], ptr [[BP]], i32 0, i32 [[IDX2:2]]562// CHECK-DAG: [[PADDR2:%.+]] = getelementptr inbounds [5 x ptr], ptr [[P]], i32 0, i32 [[IDX2]]563// CHECK-DAG: [[BPADDR3:%.+]] = getelementptr inbounds [5 x ptr], ptr [[BP]], i32 0, i32 [[IDX3:3]]564// CHECK-DAG: [[PADDR3:%.+]] = getelementptr inbounds [5 x ptr], ptr [[P]], i32 0, i32 [[IDX3]]565// CHECK-DAG: [[BPADDR4:%.+]] = getelementptr inbounds [5 x ptr], ptr [[BP]], i32 0, i32 [[IDX4:4]]566// CHECK-DAG: [[PADDR4:%.+]] = getelementptr inbounds [5 x ptr], ptr [[P]], i32 0, i32 [[IDX4]]567// CHECK-DAG: [[PSZ4:%.+]] = getelementptr inbounds [5 x i64], ptr [[PSZ:%.+]], i32 0, i32 [[IDX4]]568 569// The names below are not necessarily consistent with the names used for the570// addresses above as some are repeated.571// CHECK-DAG: store ptr %{{.+}}, ptr [[BPADDR4]]572// CHECK-DAG: store ptr %{{.+}}, ptr [[PADDR4]]573// CHECK-DAG: store i64 [[CSIZE]], ptr [[PSZ4]]574 575// CHECK-DAG: store i[[SZ]] [[VLA0]], ptr [[BPADDR3]]576// CHECK-DAG: store i[[SZ]] [[VLA0]], ptr [[PADDR3]]577 578// CHECK-DAG: store i[[SZ]] 2, ptr [[BPADDR2]]579// CHECK-DAG: store i[[SZ]] 2, ptr [[PADDR2]]580 581// CHECK-DAG: store i[[SZ]] [[B_CVAL]], ptr [[BPADDR1]]582// CHECK-DAG: store i[[SZ]] [[B_CVAL]], ptr [[PADDR1]]583 584// CHECK-DAG: store ptr [[THIS:%.+]], ptr [[BPADDR0]]585// CHECK-DAG: store ptr [[A:%.+]], ptr [[PADDR0]]586 587// CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0588// CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]]589// CHECK: [[FAIL]]590// CHECK: call void [[HVT7:@.+]]({{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}})591// CHECK-NEXT: br label %[[END]]592// CHECK: [[END]]593// CHECK-NEXT: br label %[[IFEND:.+]]594// CHECK: [[IFELSE]]595// CHECK: call void [[HVT7]]({{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}})596// CHECK-NEXT: br label %[[IFEND]]597 598// CHECK: [[IFEND]]599 600//601// CHECK: define {{.*}}[[FSTATIC]]602//603// CHECK: [[IF:%.+]] = icmp sgt i32 {{[^,]+}}, 50604// CHECK: br i1 [[IF]], label %[[IFTHEN:[^,]+]], label %[[IFELSE:[^,]+]]605// CHECK: [[IFTHEN]]606// CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 {{.+}}, i32 {{.+}}, i32 {{.+}}, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])607// CHECK-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2608// CHECK-DAG: store ptr [[BPR:%.+]], ptr [[BPARG]]609// CHECK-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3610// CHECK-DAG: store ptr [[PR:%.+]], ptr [[PARG]]611// CHECK-DAG: [[BPR]] = getelementptr inbounds [4 x ptr], ptr [[BP:%.+]], i32 0, i32 0612// CHECK-DAG: [[PR]] = getelementptr inbounds [4 x ptr], ptr [[P:%.+]], i32 0, i32 0613 614// CHECK-DAG: [[BPADDR0:%.+]] = getelementptr inbounds [4 x ptr], ptr [[BP]], i32 0, i32 0615// CHECK-DAG: [[PADDR0:%.+]] = getelementptr inbounds [4 x ptr], ptr [[P]], i32 0, i32 0616// CHECK-DAG: store i[[SZ]] [[VAL0:%[^,]+]], ptr [[BPADDR0]]617// CHECK-DAG: store i[[SZ]] [[VAL0]], ptr [[PADDR0]]618 619// CHECK-DAG: [[BPADDR1:%.+]] = getelementptr inbounds [4 x ptr], ptr [[BP]], i32 0, i32 1620// CHECK-DAG: [[PADDR1:%.+]] = getelementptr inbounds [4 x ptr], ptr [[P]], i32 0, i32 1621// CHECK-DAG: store i[[SZ]] [[VAL1:%[^,]+]], ptr [[BPADDR1]]622// CHECK-DAG: store i[[SZ]] [[VAL1]], ptr [[PADDR1]]623 624// CHECK-DAG: [[BPADDR2:%.+]] = getelementptr inbounds [4 x ptr], ptr [[BP]], i32 0, i32 2625// CHECK-DAG: [[PADDR2:%.+]] = getelementptr inbounds [4 x ptr], ptr [[P]], i32 0, i32 2626// CHECK-DAG: store i[[SZ]] [[VAL2:%[^,]+]], ptr [[BPADDR2]]627// CHECK-DAG: store i[[SZ]] [[VAL2]], ptr [[PADDR2]]628 629// CHECK-DAG: [[BPADDR3:%.+]] = getelementptr inbounds [4 x ptr], ptr [[BP]], i32 0, i32 3630// CHECK-DAG: [[PADDR3:%.+]] = getelementptr inbounds [4 x ptr], ptr [[P]], i32 0, i32 3631// CHECK-DAG: store ptr [[VAL3:%[^,]+]], ptr [[BPADDR3]]632// CHECK-DAG: store ptr [[VAL3]], ptr [[PADDR3]]633 634// CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0635// CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]]636// CHECK: [[FAIL]]637// CHECK: call void [[HVT6:@.+]]({{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}})638// CHECK-NEXT: br label %[[END]]639// CHECK: [[END]]640// CHECK-NEXT: br label %[[IFEND:.+]]641// CHECK: [[IFELSE]]642// CHECK: call void [[HVT6]]({{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}})643// CHECK-NEXT: br label %[[IFEND]]644 645// CHECK: [[IFEND]]646 647//648// CHECK: define {{.*}}[[FTEMPLATE]]649//650// CHECK: [[IF:%.+]] = icmp sgt i32 {{[^,]+}}, 40651// CHECK: br i1 [[IF]], label %[[IFTHEN:[^,]+]], label %[[IFELSE:[^,]+]]652// CHECK: [[IFTHEN]]653// CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 {{.+}}, i32 {{.+}}, i32 {{.+}}, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])654// CHECK-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2655// CHECK-DAG: store ptr [[BPR:%.+]], ptr [[BPARG]]656// CHECK-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3657// CHECK-DAG: store ptr [[PR:%.+]], ptr [[PARG]]658// CHECK-DAG: [[BPR]] = getelementptr inbounds [3 x ptr], ptr [[BP:%.+]], i32 0, i32 0659// CHECK-DAG: [[PR]] = getelementptr inbounds [3 x ptr], ptr [[P:%.+]], i32 0, i32 0660 661// CHECK-DAG: [[BPADDR0:%.+]] = getelementptr inbounds [3 x ptr], ptr [[BP]], i32 0, i32 0662// CHECK-DAG: [[PADDR0:%.+]] = getelementptr inbounds [3 x ptr], ptr [[P]], i32 0, i32 0663// CHECK-DAG: store i[[SZ]] [[VAL0:%[^,]+]], ptr [[BPADDR0]]664// CHECK-DAG: store i[[SZ]] [[VAL0]], ptr [[PADDR0]]665 666// CHECK-DAG: [[BPADDR1:%.+]] = getelementptr inbounds [3 x ptr], ptr [[BP]], i32 0, i32 1667// CHECK-DAG: [[PADDR1:%.+]] = getelementptr inbounds [3 x ptr], ptr [[P]], i32 0, i32 1668// CHECK-DAG: store i[[SZ]] [[VAL1:%[^,]+]], ptr [[BPADDR1]]669// CHECK-DAG: store i[[SZ]] [[VAL1]], ptr [[PADDR1]]670 671// CHECK-DAG: [[BPADDR2:%.+]] = getelementptr inbounds [3 x ptr], ptr [[BP]], i32 0, i32 2672// CHECK-DAG: [[PADDR2:%.+]] = getelementptr inbounds [3 x ptr], ptr [[P]], i32 0, i32 2673// CHECK-DAG: store ptr [[VAL2:%[^,]+]], ptr [[BPADDR2]]674// CHECK-DAG: store ptr [[VAL2]], ptr [[PADDR2]]675 676// CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0677// CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]]678// CHECK: [[FAIL]]679// CHECK: call void [[HVT5:@.+]]({{[^,]+}}, {{[^,]+}}, {{[^,]+}})680// CHECK-NEXT: br label %[[END]]681// CHECK: [[END]]682// CHECK-NEXT: br label %[[IFEND:.+]]683// CHECK: [[IFELSE]]684// CHECK: call void [[HVT5]]({{[^,]+}}, {{[^,]+}}, {{[^,]+}})685// CHECK-NEXT: br label %[[IFEND]]686 687// CHECK: [[IFEND]]688 689// OMP45: define internal void @__omp_offloading_{{.+}}_{{.+}}bar{{.+}}_l{{[0-9]+}}(i[[SZ]] noundef %{{.+}})690 691// OMP45: define {{.*}}@{{.*}}zee{{.*}}692 693// OMP45: [[LOCAL_THIS:%.+]] = alloca ptr694// OMP45: [[BP:%.+]] = alloca [1 x ptr]695// OMP45: [[P:%.+]] = alloca [1 x ptr]696// OMP45: [[LOCAL_THIS1:%.+]] = load ptr, ptr [[LOCAL_THIS]]697 698// OMP45: call void @__kmpc_critical(699// OMP45: [[ARR_IDX:%.+]] = getelementptr inbounds [[S2]], ptr [[LOCAL_THIS1]], i[[SZ]] 0700// OMP45: [[ARR_IDX2:%.+]] = getelementptr inbounds [[S2]], ptr [[LOCAL_THIS1]], i[[SZ]] 0701 702// OMP45-DAG: [[BPADDR0:%.+]] = getelementptr inbounds [1 x ptr], ptr [[BP]], i32 0, i32 0703// OMP45-DAG: [[PADDR0:%.+]] = getelementptr inbounds [1 x ptr], ptr [[P]], i32 0, i32 0704// OMP45-DAG: store ptr [[ARR_IDX]], ptr [[BPADDR0]]705// OMP45-DAG: store ptr [[ARR_IDX2]], ptr [[PADDR0]]706 707// OMP45: [[BPR:%.+]] = getelementptr inbounds [1 x ptr], ptr [[BP]], i32 0, i32 0708// OMP45: [[PR:%.+]] = getelementptr inbounds [1 x ptr], ptr [[P]], i32 0, i32 0709// OMP45: [[RET:%.+]] = call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 {{.+}}, i32 {{.+}}, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])710// OMP45-NEXT: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0711// OMP45-NEXT: br i1 [[ERROR]], label %[[FAIL:[^,]+]], label %[[END:[^,]+]]712// OMP45: [[FAIL]]713// OMP45: call void [[HVT0:@.+]](ptr [[LOCAL_THIS1]])714// OMP45-NEXT: br label %[[END]]715// OMP45: [[END]]716// OMP45: call void @__kmpc_end_critical(717 718// Check that the offloading functions are emitted and that the arguments are719// correct and loaded correctly for the target regions of the callees of bar().720 721// CHECK: define internal void [[HVT7]]722// Create local storage for each capture.723// CHECK: [[LOCAL_THIS:%.+]] = alloca ptr724// CHECK: [[LOCAL_B:%.+]] = alloca i[[SZ]]725// CHECK: [[LOCAL_VLA1:%.+]] = alloca i[[SZ]]726// CHECK: [[LOCAL_VLA2:%.+]] = alloca i[[SZ]]727// CHECK: [[LOCAL_C:%.+]] = alloca ptr728// CHECK-DAG: store ptr [[ARG_THIS:%.+]], ptr [[LOCAL_THIS]]729// CHECK-DAG: store i[[SZ]] [[ARG_B:%.+]], ptr [[LOCAL_B]]730// CHECK-DAG: store i[[SZ]] [[ARG_VLA1:%.+]], ptr [[LOCAL_VLA1]]731// CHECK-DAG: store i[[SZ]] [[ARG_VLA2:%.+]], ptr [[LOCAL_VLA2]]732// CHECK-DAG: store ptr [[ARG_C:%.+]], ptr [[LOCAL_C]]733// Store captures in the context.734// CHECK-DAG: [[REF_THIS:%.+]] = load ptr, ptr [[LOCAL_THIS]],735// CHECK-DAG: [[VAL_VLA1:%.+]] = load i[[SZ]], ptr [[LOCAL_VLA1]],736// CHECK-DAG: [[VAL_VLA2:%.+]] = load i[[SZ]], ptr [[LOCAL_VLA2]],737// CHECK-DAG: [[REF_C:%.+]] = load ptr, ptr [[LOCAL_C]],738// Use captures.739// CHECK-DAG: getelementptr inbounds nuw [[S1]], ptr [[REF_THIS]], i32 0, i32 0740// CHECK-64-DAG:load i32, ptr [[LOCAL_B]]741// CHECK-32-DAG:load i32, ptr [[LOCAL_B]]742// CHECK-DAG: getelementptr inbounds i16, ptr [[REF_C]], i[[SZ]] %{{.+}}743 744 745// CHECK: define internal void [[HVT6]]746// Create local storage for each capture.747// CHECK: [[LOCAL_A:%.+]] = alloca i[[SZ]]748// CHECK: [[LOCAL_AA:%.+]] = alloca i[[SZ]]749// CHECK: [[LOCAL_AAA:%.+]] = alloca i[[SZ]]750// CHECK: [[LOCAL_B:%.+]] = alloca ptr751// CHECK-DAG: store i[[SZ]] [[ARG_A:%.+]], ptr [[LOCAL_A]]752// CHECK-DAG: store i[[SZ]] [[ARG_AA:%.+]], ptr [[LOCAL_AA]]753// CHECK-DAG: store i[[SZ]] [[ARG_AAA:%.+]], ptr [[LOCAL_AAA]]754// CHECK-DAG: store ptr [[ARG_B:%.+]], ptr [[LOCAL_B]]755// Store captures in the context.756// CHECK-DAG: [[REF_B:%.+]] = load ptr, ptr [[LOCAL_B]],757// Use captures.758// CHECK-64-DAG: load i32, ptr [[LOCAL_A]]759// CHECK-DAG: load i16, ptr [[LOCAL_AA]]760// CHECK-DAG: load i8, ptr [[LOCAL_AAA]]761// CHECK-32-DAG: load i32, ptr [[LOCAL_A]]762// CHECK-DAG: getelementptr inbounds [10 x i32], ptr [[REF_B]], i[[SZ]] 0, i[[SZ]] 2763 764// CHECK: define internal void [[HVT5]]765// Create local storage for each capture.766// CHECK: [[LOCAL_A:%.+]] = alloca i[[SZ]]767// CHECK: [[LOCAL_AA:%.+]] = alloca i[[SZ]]768// CHECK: [[LOCAL_B:%.+]] = alloca ptr769// CHECK-DAG: store i[[SZ]] [[ARG_A:%.+]], ptr [[LOCAL_A]]770// CHECK-DAG: store i[[SZ]] [[ARG_AA:%.+]], ptr [[LOCAL_AA]]771// CHECK-DAG: store ptr [[ARG_B:%.+]], ptr [[LOCAL_B]]772// Store captures in the context.773// CHECK-DAG: [[REF_B:%.+]] = load ptr, ptr [[LOCAL_B]],774// Use captures.775// CHECK-64-DAG: load i32, ptr [[LOCAL_A]]776// CHECK-32-DAG: load i32, ptr [[LOCAL_A]]777// CHECK-DAG: load i16, ptr [[LOCAL_AA]]778// CHECK-DAG: getelementptr inbounds [10 x i32], ptr [[REF_B]], i[[SZ]] 0, i[[SZ]] 2779 780// OMP50: define internal void @__omp_offloading_{{.+}}_{{.+}}bar{{.+}}_l{{[0-9]+}}(i[[SZ]] noundef %{{.+}})781 782// OMP50: define {{.*}}@{{.*}}zee{{.*}}783 784// OMP50: [[LOCAL_THIS:%.+]] = alloca ptr785// OMP50: [[BP:%.+]] = alloca [1 x ptr]786// OMP50: [[P:%.+]] = alloca [1 x ptr]787// OMP50: [[LOCAL_THIS1:%.+]] = load ptr, ptr [[LOCAL_THIS]]788// OMP50: [[ARR_IDX:%.+]] = getelementptr inbounds [[S2]], ptr [[LOCAL_THIS1]], i[[SZ]] 0789// OMP50: [[ARR_IDX2:%.+]] = getelementptr inbounds [[S2]], ptr [[LOCAL_THIS1]], i[[SZ]] 0790 791// OMP50-DAG: [[BPADDR0:%.+]] = getelementptr inbounds [1 x ptr], ptr [[BP]], i32 0, i32 0792// OMP50-DAG: [[PADDR0:%.+]] = getelementptr inbounds [1 x ptr], ptr [[P]], i32 0, i32 0793// OMP50-DAG: store ptr [[ARR_IDX]], ptr [[BPADDR0]]794// OMP50-DAG: store ptr [[ARR_IDX2]], ptr [[PADDR0]]795 796// OMP50: [[BPR:%.+]] = getelementptr inbounds [1 x ptr], ptr [[BP]], i32 0, i32 0797// OMP50: [[PR:%.+]] = getelementptr inbounds [1 x ptr], ptr [[P]], i32 0, i32 0798// OMP50: [[RET:%.+]] = call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 {{.+}}, i32 {{.+}}, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])799// OMP50-NEXT: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0800// OMP50-NEXT: br i1 [[ERROR]], label %[[FAIL:[^,]+]], label %[[END:[^,]+]]801// OMP50: [[FAIL]]802// OMP50: call void [[HVT0:@.+]](ptr [[LOCAL_THIS1]])803// OMP50-NEXT: br label %[[END]]804// OMP50: [[END]]805 806void bar () {807#define pragma_target _Pragma("omp target")808pragma_target809{810 global = 0;811#pragma omp parallel shared(global)812 global = 1;813}814}815 816class S2 {817 int a, b, c;818 819public:820 void zee() {821#pragma omp critical822 #pragma omp target map(this[0])823 a++;824 }825};826 827#ifdef _DOMP51828void thread_limit_target(int TargetTL, int TeamsTL) {829 830#pragma omp target831{}832// OMP51: call i32 @__tgt_target_kernel({{.*}}, i64 -1, i32 -1, i32 0,833 834#pragma omp target835#pragma omp teams836{}837// OMP51: call i32 @__tgt_target_kernel({{.*}}, i64 -1, i32 0, i32 0,838 839#pragma omp target thread_limit(TargetTL)840{}841// OMP51: [[TL:%.*]] = load {{.*}} %TargetTL.addr842// OMP51: store {{.*}} [[TL]], {{.*}} [[CEA:%.*]]843// OMP51: load {{.*}} [[CEA]]844// OMP51: [[CE:%.*]] = load {{.*}} [[CEA]]845// OMP51: call ptr @__kmpc_omp_task_alloc({{.*@.omp_task_entry.*}})846// OMP51: call i32 [[OMP_TASK_ENTRY]]847 848#pragma omp target thread_limit(TargetTL)849#pragma omp teams850{}851// OMP51: [[TL:%.*]] = load {{.*}} %TargetTL.addr852// OMP51: store {{.*}} [[TL]], {{.*}} [[CEA:%.*]]853// OMP51: load {{.*}} [[CEA]]854// OMP51: call ptr @__kmpc_omp_task_alloc({{.*@.omp_task_entry.*}})855// OMP51: call i32 [[OMP_TASK_ENTRY]]856 857#pragma omp target858#pragma omp teams thread_limit(TeamsTL)859{}860// OMP51: load {{.*}} %TeamsTL.addr861// OMP51: [[TeamsL:%.*]] = load {{.*}} %TeamsTL.addr862// OMP51: call i32 @__tgt_target_kernel({{.*}}, i64 -1, i32 0, i32 [[TeamsL]],863 864#pragma omp target thread_limit(TargetTL)865#pragma omp teams thread_limit(TeamsTL)866{}867// OMP51: load {{.*}} %TeamsTL.addr868// OMP51: [[TeamsL:%.*]] = load {{.*}} %TeamsTL.addr869// OMP51: call ptr @__kmpc_omp_task_alloc({{.*@.omp_task_entry.*}})870// OMP51: call i32 [[OMP_TASK_ENTRY]]871 872}873#endif874// Check that the offloading functions are called after setting thread_limit in the task entry functions875 876// OMP51: define internal {{.*}}i32 [[OMP_TASK_ENTRY:@.+]](i32 {{.*}}%0, ptr noalias noundef %1)877// OMP51: call void @__kmpc_set_thread_limit(ptr @{{.+}}, i32 %{{.+}}, i32 %{{.+}})878// OMP51: call i32 @__tgt_target_kernel({{.*}}, i64 -1, i32 -1,879 880// OMP51: define internal {{.*}}i32 [[OMP_TASK_ENTRY:@.+]](i32 {{.*}}%0, ptr noalias noundef %1)881// OMP51: call void @__kmpc_set_thread_limit(ptr @{{.+}}, i32 %{{.+}}, i32 %{{.+}})882// OMP51: call i32 @__tgt_target_kernel({{.*}}, i64 -1, i32 0,883 884// OMP51: define internal {{.*}}i32 [[OMP_TASK_ENTRY:@.+]](i32 {{.*}}%0, ptr noalias noundef %1)885// OMP51: call void @__kmpc_set_thread_limit(ptr @{{.+}}, i32 %{{.+}}, i32 %{{.+}})886// OMP51: call i32 @__tgt_target_kernel({{.*}}, i64 -1, i32 0,887 888 889int main () {890 S2 bar;891 bar.zee();892}893 894#endif895