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 OMP45
3 // 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 %s
4 // 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 OMP45
5 // 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 OMP45
6 // 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 %s
7 // 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 OMP45
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 OMP50
10 // RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
11 // 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 OMP50
12 // 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 OMP50
13 // RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s
14 // 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 OMP50
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 OMP51
17 // 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 %s
18 // 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 OMP51
19 // 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 OMP51
20 // 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 %s
21 // 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 OMP51
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 %s
24 // RUN: %clang_cc1 -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
25 // 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 %s
26 // 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 %s
27 // RUN: %clang_cc1 -fopenmp-simd -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s
28 // 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 %s
29 // SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
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.bc
33 // 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-64
34 // 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 %s
35 // 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-64
36 // 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.bc
37 // 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-32
38 // 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 %s
39 // 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-32
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.bc
42 // 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-64
43 // 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 %s
44 // 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-64
45 // 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.bc
46 // 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-32
47 // 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 %s
48 // 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-32
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.bc
51 // 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 %s
52 // 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 %s
53 // 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 %s
54 // 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.bc
55 // 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 %s
56 // 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 %s
57 // 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 %s
58 // SIMD-ONLY1-NOT: {{__kmpc|__tgt}}
60 // expected-no-diagnostics
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 { ptr, ptr, i[[SZ:32|64]], i32, i32 }
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 }
75 // TCHECK: [[ENTTY:%.+]] = type { ptr, ptr, i{{32|64}}, i32, i32 }
77 // We have 9 target regions, but only 8 that actually will generate offloading
78 // code and have mapped arguments, and only 6 have all-constant map sizes.
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 0
97 // CHECK-DAG: @{{.*}} = weak constant i8 0
98 // CHECK-DAG: @{{.*}} = weak constant i8 0
99 // CHECK-DAG: @{{.*}} = weak constant i8 0
100 // CHECK-DAG: @{{.*}} = weak constant i8 0
101 // CHECK-DAG: @{{.*}} = weak constant i8 0
102 // CHECK-DAG: @{{.*}} = weak constant i8 0
103 // CHECK-DAG: @{{.*}} = weak constant i8 0
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]]
117 template<typename tx
, typename ty
>
126 // CHECK: define {{.*}}[[FOO:@.+]](
128 // CHECK: [[OFFLOADBPTR:%.+]] = alloca [2 x ptr], align
129 // CHECK: [[OFFLOADPTR:%.+]] = alloca [2 x ptr], align
130 // CHECK: [[OFFLOADMAPPER:%.+]] = alloca [2 x ptr], align
137 TT
<long long, char> d
;
140 // CHECK: [[ADD:%.+]] = add nsw i32
141 // CHECK: store i32 [[ADD]], ptr [[DEVICE_CAP:%.+]],
142 // CHECK: [[DEV:%.+]] = load i32, ptr [[DEVICE_CAP]],
143 // CHECK: [[DEVICE:%.+]] = sext i32 [[DEV]] to i64
144 // 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]], 0
146 // CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:[^,]+]], label %[[END:[^,]+]]
148 // CHECK: call void [[HVT0:@.+]]()
149 // CHECK-NEXT: br label %[[END]]
151 #pragma omp target device(global + a)
155 // CHECK: [[BPRGEP:%.+]] = getelementptr inbounds [2 x ptr], ptr [[OFFLOADBPTR]], i32 0, i32 0
156 // CHECK: [[PRGEP:%.+]] = getelementptr inbounds [2 x ptr], ptr [[OFFLOADPTR]], i32 0, i32 0
157 // CHECK: [[BPRGEP:%.+]] = getelementptr inbounds [2 x ptr], ptr [[OFFLOADBPTR]], i32 0, i32 0
158 // CHECK: [[PRGEP:%.+]] = getelementptr inbounds [2 x ptr], ptr [[OFFLOADPTR]], i32 0, i32 0
159 // CHECK: [[DEVICE:%.+]] = sext i32 {{%.+}} to i64
160 // 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 1
163 // CHECK-32: [[SIZEGEP:%.+]] = getelementptr inbounds nuw [[KMP_PRIVATES_T]], ptr [[TASK_WITH_PRIVATES_GEP]], i32 0, i32 0
164 // 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 3
166 // 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 4
168 // 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 1
170 // 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 2
172 // 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 3
174 // 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) nowait
183 // CHECK: call void [[HVT1:@.+]](i[[SZ]] {{[^,]+}})
184 #pragma omp target if(0) firstprivate(global)
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 2
191 // CHECK-DAG: store ptr [[BP:%.+]], ptr [[BPARG]]
192 // CHECK-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
193 // CHECK-DAG: store ptr [[P:%.+]], ptr [[PARG]]
194 // CHECK-DAG: [[BP]] = getelementptr inbounds [1 x ptr], ptr [[BPR:%[^,]+]], i32 0, i32 0
195 // CHECK-DAG: [[P]] = getelementptr inbounds [1 x ptr], ptr [[PR:%[^,]+]], i32 0, i32 0
196 // 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]]
201 // CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0
202 // CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:[^,]+]], label %[[END:[^,]+]]
204 // CHECK: call void [[HVT2:@.+]](i[[SZ]] {{[^,]+}})
205 // CHECK-NEXT: br label %[[END]]
207 #pragma omp target if (1)
212 // CHECK: [[IF:%.+]] = icmp sgt i32 {{[^,]+}}, 10
213 // CHECK: br i1 [[IF]], label %[[IFTHEN:[^,]+]], label %[[IFELSE:[^,]+]]
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 2
217 // CHECK-DAG: store ptr [[BPR:%.+]], ptr [[BPARG]]
218 // CHECK-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
219 // CHECK-DAG: store ptr [[PR:%.+]], ptr [[PARG]]
220 // CHECK-DAG: [[BPR]] = getelementptr inbounds [2 x ptr], ptr [[BP:%[^,]+]], i32 0, i32 0
221 // CHECK-DAG: [[PR]] = getelementptr inbounds [2 x ptr], ptr [[P:%[^,]+]], i32 0, i32 0
223 // CHECK-DAG: [[BPADDR0:%.+]] = getelementptr inbounds [2 x ptr], ptr [[BP]], i32 0, i32 0
224 // CHECK-DAG: [[PADDR0:%.+]] = getelementptr inbounds [2 x ptr], ptr [[P]], i32 0, i32 0
225 // CHECK-DAG: store i[[SZ]] [[BP0:%[^,]+]], ptr [[BPADDR0]]
226 // CHECK-DAG: store i[[SZ]] [[P0:%[^,]+]], ptr [[PADDR0]]
228 // CHECK-DAG: [[BPADDR1:%.+]] = getelementptr inbounds [2 x ptr], ptr [[BP]], i32 0, i32 1
229 // CHECK-DAG: [[PADDR1:%.+]] = getelementptr inbounds [2 x ptr], ptr [[P]], i32 0, i32 1
230 // CHECK-DAG: store i[[SZ]] [[BP1:%[^,]+]], ptr [[BPADDR1]]
231 // CHECK-DAG: store i[[SZ]] [[P1:%[^,]+]], ptr [[PADDR1]]
232 // CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0
233 // CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]]
235 // CHECK: call void [[HVT3:@.+]]({{[^,]+}}, {{[^,]+}})
236 // CHECK-NEXT: br label %[[END]]
238 // CHECK-NEXT: br label %[[IFEND:.+]]
240 // CHECK: call void [[HVT3]]({{[^,]+}}, {{[^,]+}})
241 // CHECK-NEXT: br label %[[IFEND]]
244 #pragma omp target if (n > 10)
250 // We capture 3 VLA sizes in this target region
251 // 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]],
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]],
259 // CHECK: [[IF:%.+]] = icmp sgt i32 {{[^,]+}}, 20
260 // CHECK: br i1 [[IF]], label %[[TRY:[^,]+]], label %[[IFELSE:[^,]+]]
262 // CHECK-64: [[BNSIZE:%.+]] = mul nuw i64 [[VLA0:%.+]], 4
263 // CHECK-32: [[BNSZSIZE:%.+]] = mul nuw i32 [[VLA0:%.+]], 4
264 // CHECK-32: [[BNSIZE:%.+]] = sext i32 [[BNSZSIZE]] to i64
265 // CHECK: [[CNELEMSIZE2:%.+]] = mul nuw i[[SZ]] 5, [[VLA1:%.+]]
266 // CHECK-64: [[CNSIZE:%.+]] = mul nuw i64 [[CNELEMSIZE2]], 8
267 // CHECK-32: [[CNSZSIZE:%.+]] = mul nuw i32 [[CNELEMSIZE2]], 8
268 // CHECK-32: [[CNSIZE:%.+]] = sext i32 [[CNSZSIZE]] to i64
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 2
272 // CHECK-DAG: store ptr [[BPR:%.+]], ptr [[BPARG]]
273 // CHECK-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
274 // CHECK-DAG: store ptr [[PR:%.+]], ptr [[PARG]]
275 // CHECK-DAG: [[SARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 4
276 // CHECK-DAG: store ptr [[SZ4:%.+]], ptr [[SARG]]
277 // CHECK-DAG: [[BPR]] = getelementptr inbounds [9 x ptr], ptr [[BP:%[^,]+]], i32 0, i32 0
278 // CHECK-DAG: [[PR]] = getelementptr inbounds [9 x ptr], ptr [[P:%[^,]+]], i32 0, i32 0
279 // CHECK-DAG: [[SZ4]] = getelementptr inbounds [9 x i64], ptr [[PSZ:%[^,]+]], i32 0, i32 0
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]]
302 // The names below are not necessarily consistent with the names used for the
303 // addresses above as some are repeated.
304 // CHECK-DAG: store i[[SZ]] [[VLA0]], ptr [[BPADDR2]]
305 // CHECK-DAG: store i[[SZ]] [[VLA0]], ptr [[PADDR2]]
307 // CHECK-DAG: store i[[SZ]] [[VLA1]], ptr [[BPADDR6]]
308 // CHECK-DAG: store i[[SZ]] [[VLA1]], ptr [[PADDR6]]
310 // CHECK-DAG: store i[[SZ]] 5, ptr [[BPADDR5]]
311 // CHECK-DAG: store i[[SZ]] 5, ptr [[PADDR5]]
313 // CHECK-DAG: store i[[SZ]] [[A_CVAL]], ptr [[BPADDR0]]
314 // CHECK-DAG: store i[[SZ]] [[A_CVAL]], ptr [[PADDR0]]
316 // CHECK-DAG: store ptr %{{.+}}, ptr [[BPADDR1]]
317 // CHECK-DAG: store ptr %{{.+}}, ptr [[PADDR1]]
319 // CHECK-DAG: store ptr %{{.+}}, ptr [[BPADDR3]]
320 // CHECK-DAG: store ptr %{{.+}}, ptr [[PADDR3]]
321 // CHECK-DAG: store i64 [[BNSIZE]], ptr [[PSZ3]]
323 // CHECK-DAG: store ptr %{{.+}}, ptr [[BPADDR4]]
324 // CHECK-DAG: store ptr %{{.+}}, ptr [[PADDR4]]
326 // CHECK-DAG: store ptr %{{.+}}, ptr [[BPADDR7]]
327 // CHECK-DAG: store ptr %{{.+}}, ptr [[PADDR7]]
328 // CHECK-DAG: store i64 [[CNSIZE]], ptr [[PSZ7]]
330 // CHECK-DAG: store ptr %{{.+}}, ptr [[BPADDR8]]
331 // CHECK-DAG: store ptr %{{.+}}, ptr [[PADDR8]]
333 // CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0
334 // CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]]
336 // CHECK: call void [[HVT4:@.+]]({{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}})
337 // CHECK-NEXT: br label %[[END]]
339 // CHECK-NEXT: br label %[[IFEND:.+]]
341 // CHECK: call void [[HVT4]]({{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}})
342 // CHECK-NEXT: br label %[[IFEND]]
345 #pragma omp target if (n > 20)
359 // Check that the offloading functions are emitted and that the arguments are
360 // correct and loaded correctly for the target regions in foo().
362 // CHECK: define internal void [[HVT0]]()
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 2
368 // CHECK-DAG: store ptr [[BPR:%.+]], ptr [[BPARG]]
369 // CHECK-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
370 // CHECK-DAG: store ptr [[PR:%.+]], ptr [[PARG]]
371 // CHECK-DAG: [[SARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 4
372 // CHECK-DAG: store ptr [[SIZE:%.+]], ptr [[SARG]]
373 // CHECK-DAG: [[DEVICE]] = sext i32 [[DEV:%.+]] to i64
374 // CHECK-DAG: [[DEV]] = load i32, ptr [[DEVADDR:%.+]], align
375 // CHECK-DAG: [[DEVADDR]] = getelementptr inbounds nuw [[ANON_T]], ptr {{%.+}}, i32 0, i32 2
376 // CHECK-DAG: [[BPR]] = load ptr, ptr [[FPPTR_BPR:%.+]], align
377 // CHECK-DAG: [[PR]] = load ptr, ptr [[FPPTR_PR:%.+]], align
378 // CHECK-DAG: [[SIZE]] = load ptr, ptr [[FPPTR_SIZE:%.+]], align
379 // 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]], align
381 // CHECK-DAG: {{%.+}} = load ptr, ptr [[FPPTR_GLOBAL:%.+]], align
382 // CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0
383 // CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:[^,]+]], label %[[END:[^,]+]]
385 // CHECK: [[PLOCAL:%.+]] = load ptr, ptr [[PLOCALADDR]], align
386 // CHECK: [[GLOBAL:%.+]] = load i32, ptr {{@.+}}, align
387 // CHECK-32: store i32 [[GLOBAL]], ptr [[GLOBALCAST:%.+]], align
388 // CHECK-64: store i32 [[GLOBAL]], ptr [[GLOBALCAST:%.+]], align
389 // CHECK: [[GLOBAL:%.+]] = load i[[SZ]], ptr [[GLOBALCAST]], align
390 // CHECK: call void [[HVT0_]](ptr [[PLOCAL]], i[[SZ]] [[GLOBAL]])
391 // CHECK-NEXT: br label %[[END]]
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]], align
397 // CHECK: store i[[SZ]] %{{.+}}, ptr [[AA_ADDR]], align
398 // CHECK-64: load i32, ptr [[AA_ADDR]], align
399 // CHECK-32: load i32, ptr [[AA_ADDR]], align
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]], align
404 // CHECK: store i[[SZ]] %{{.+}}, ptr [[AA_ADDR]], align
405 // CHECK: load i16, ptr [[AA_ADDR]], align
407 // CHECK: define internal void [[HVT3]]
408 // Create stack storage and store argument in there.
409 // CHECK: [[A_ADDR:%.+]] = alloca i[[SZ]], align
410 // CHECK: [[AA_ADDR:%.+]] = alloca i[[SZ]], align
411 // CHECK-DAG: store i[[SZ]] %{{.+}}, ptr [[A_ADDR]], align
412 // CHECK-DAG: store i[[SZ]] %{{.+}}, ptr [[AA_ADDR]], align
413 // CHECK-64-DAG:load i32, ptr [[A_ADDR]], align
414 // CHECK-32-DAG:load i32, ptr [[A_ADDR]], align
415 // CHECK-DAG: load i16, ptr [[AA_ADDR]], align
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 ptr
421 // CHECK: [[LOCAL_VLA1:%.+]] = alloca i[[SZ]]
422 // CHECK: [[LOCAL_BN:%.+]] = alloca ptr
423 // CHECK: [[LOCAL_C:%.+]] = alloca ptr
424 // CHECK: [[LOCAL_VLA2:%.+]] = alloca i[[SZ]]
425 // CHECK: [[LOCAL_VLA3:%.+]] = alloca i[[SZ]]
426 // CHECK: [[LOCAL_CN:%.+]] = alloca ptr
427 // CHECK: [[LOCAL_D:%.+]] = alloca ptr
428 // 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]]
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]],
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]] 2
451 // CHECK-DAG: getelementptr inbounds float, ptr [[REF_BN]], i[[SZ]] 3
452 // CHECK-DAG: getelementptr inbounds [5 x [10 x double]], ptr [[REF_C]], i[[SZ]] 0, i[[SZ]] 1
453 // CHECK-DAG: getelementptr inbounds double, ptr [[REF_CN]], i[[SZ]] %{{.+}}
454 // CHECK-DAG: getelementptr inbounds nuw [[TT]], ptr [[REF_D]], i32 0, i32 0
456 template<typename tx
>
457 tx
ftemplate(int n
) {
462 #pragma omp target if(n>40)
479 #pragma omp target if(n>50)
497 #pragma omp target if(n>60)
499 this->a
= (double)b
+ 1.5;
503 return c
[1][1] + (int)b
;
507 // CHECK: define {{.*}}@{{.*}}bar{{.*}}
511 // CHECK: call {{.*}}i32 [[FOO]](i32 {{.*}})
515 // CHECK: call {{.*}}i32 [[FS1:@.+]](ptr {{.*}}, i32 {{.*}})
518 // CHECK: call {{.*}}i32 [[FSTATIC:@.+]](i32 {{.*}})
521 // CHECK: call {{.*}}i32 [[FTEMPLATE:@.+]](i32 {{.*}})
522 a
+= ftemplate
<int>(n
);
528 // CHECK: define {{.*}}[[FS1]]
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]],
534 // CHECK-32: store i32 %{{.+}}, ptr %__vla_expr
535 // CHECK-32: store i32 %{{.+}}, ptr [[B_CADDR:%.+]],
536 // CHECK-32: [[B_CVAL:%.+]] = load i[[SZ]], ptr [[B_CADDR]],
538 // CHECK: [[IF:%.+]] = icmp sgt i32 {{[^,]+}}, 60
539 // CHECK: br i1 [[IF]], label %[[TRY:[^,]+]], label %[[IFELSE:[^,]+]]
541 // We capture 2 VLA sizes in this target region
542 // CHECK: [[CELEMSIZE2:%.+]] = mul nuw i[[SZ]] 2, [[VLA0:%.+]]
543 // CHECK-64: [[CSIZE:%.+]] = mul nuw i64 [[CELEMSIZE2]], 2
544 // CHECK-32: [[CSZSIZE:%.+]] = mul nuw i32 [[CELEMSIZE2]], 2
545 // CHECK-32: [[CSIZE:%.+]] = sext i32 [[CSZSIZE]] to i64
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 2
549 // CHECK-DAG: store ptr [[BPR:%.+]], ptr [[BPARG]]
550 // CHECK-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
551 // CHECK-DAG: store ptr [[PR:%.+]], ptr [[PARG]]
552 // CHECK-DAG: [[SARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 4
553 // CHECK-DAG: store ptr [[SZ7:%.+]], ptr [[SARG]]
554 // CHECK-DAG: [[BPR]] = getelementptr inbounds [5 x ptr], ptr [[BP:%.+]], i32 0, i32 0
555 // CHECK-DAG: [[PR]] = getelementptr inbounds [5 x ptr], ptr [[P:%.+]], i32 0, i32 0
556 // CHECK-DAG: [[SZ7]] = getelementptr inbounds [5 x i64], ptr [[PSZ:%.+]], i32 0, i32 0
557 // 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]]
569 // The names below are not necessarily consistent with the names used for the
570 // 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]]
575 // CHECK-DAG: store i[[SZ]] [[VLA0]], ptr [[BPADDR3]]
576 // CHECK-DAG: store i[[SZ]] [[VLA0]], ptr [[PADDR3]]
578 // CHECK-DAG: store i[[SZ]] 2, ptr [[BPADDR2]]
579 // CHECK-DAG: store i[[SZ]] 2, ptr [[PADDR2]]
581 // CHECK-DAG: store i[[SZ]] [[B_CVAL]], ptr [[BPADDR1]]
582 // CHECK-DAG: store i[[SZ]] [[B_CVAL]], ptr [[PADDR1]]
584 // CHECK-DAG: store ptr [[THIS:%.+]], ptr [[BPADDR0]]
585 // CHECK-DAG: store ptr [[A:%.+]], ptr [[PADDR0]]
587 // CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0
588 // CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]]
590 // CHECK: call void [[HVT7:@.+]]({{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}})
591 // CHECK-NEXT: br label %[[END]]
593 // CHECK-NEXT: br label %[[IFEND:.+]]
595 // CHECK: call void [[HVT7]]({{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}})
596 // CHECK-NEXT: br label %[[IFEND]]
601 // CHECK: define {{.*}}[[FSTATIC]]
603 // CHECK: [[IF:%.+]] = icmp sgt i32 {{[^,]+}}, 50
604 // CHECK: br i1 [[IF]], label %[[IFTHEN:[^,]+]], label %[[IFELSE:[^,]+]]
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 2
608 // CHECK-DAG: store ptr [[BPR:%.+]], ptr [[BPARG]]
609 // CHECK-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
610 // CHECK-DAG: store ptr [[PR:%.+]], ptr [[PARG]]
611 // CHECK-DAG: [[BPR]] = getelementptr inbounds [4 x ptr], ptr [[BP:%.+]], i32 0, i32 0
612 // CHECK-DAG: [[PR]] = getelementptr inbounds [4 x ptr], ptr [[P:%.+]], i32 0, i32 0
614 // CHECK-DAG: [[BPADDR0:%.+]] = getelementptr inbounds [4 x ptr], ptr [[BP]], i32 0, i32 0
615 // CHECK-DAG: [[PADDR0:%.+]] = getelementptr inbounds [4 x ptr], ptr [[P]], i32 0, i32 0
616 // CHECK-DAG: store i[[SZ]] [[VAL0:%[^,]+]], ptr [[BPADDR0]]
617 // CHECK-DAG: store i[[SZ]] [[VAL0]], ptr [[PADDR0]]
619 // CHECK-DAG: [[BPADDR1:%.+]] = getelementptr inbounds [4 x ptr], ptr [[BP]], i32 0, i32 1
620 // CHECK-DAG: [[PADDR1:%.+]] = getelementptr inbounds [4 x ptr], ptr [[P]], i32 0, i32 1
621 // CHECK-DAG: store i[[SZ]] [[VAL1:%[^,]+]], ptr [[BPADDR1]]
622 // CHECK-DAG: store i[[SZ]] [[VAL1]], ptr [[PADDR1]]
624 // CHECK-DAG: [[BPADDR2:%.+]] = getelementptr inbounds [4 x ptr], ptr [[BP]], i32 0, i32 2
625 // CHECK-DAG: [[PADDR2:%.+]] = getelementptr inbounds [4 x ptr], ptr [[P]], i32 0, i32 2
626 // CHECK-DAG: store i[[SZ]] [[VAL2:%[^,]+]], ptr [[BPADDR2]]
627 // CHECK-DAG: store i[[SZ]] [[VAL2]], ptr [[PADDR2]]
629 // CHECK-DAG: [[BPADDR3:%.+]] = getelementptr inbounds [4 x ptr], ptr [[BP]], i32 0, i32 3
630 // CHECK-DAG: [[PADDR3:%.+]] = getelementptr inbounds [4 x ptr], ptr [[P]], i32 0, i32 3
631 // CHECK-DAG: store ptr [[VAL3:%[^,]+]], ptr [[BPADDR3]]
632 // CHECK-DAG: store ptr [[VAL3]], ptr [[PADDR3]]
634 // CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0
635 // CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]]
637 // CHECK: call void [[HVT6:@.+]]({{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}})
638 // CHECK-NEXT: br label %[[END]]
640 // CHECK-NEXT: br label %[[IFEND:.+]]
642 // CHECK: call void [[HVT6]]({{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}})
643 // CHECK-NEXT: br label %[[IFEND]]
648 // CHECK: define {{.*}}[[FTEMPLATE]]
650 // CHECK: [[IF:%.+]] = icmp sgt i32 {{[^,]+}}, 40
651 // CHECK: br i1 [[IF]], label %[[IFTHEN:[^,]+]], label %[[IFELSE:[^,]+]]
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 2
655 // CHECK-DAG: store ptr [[BPR:%.+]], ptr [[BPARG]]
656 // CHECK-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
657 // CHECK-DAG: store ptr [[PR:%.+]], ptr [[PARG]]
658 // CHECK-DAG: [[BPR]] = getelementptr inbounds [3 x ptr], ptr [[BP:%.+]], i32 0, i32 0
659 // CHECK-DAG: [[PR]] = getelementptr inbounds [3 x ptr], ptr [[P:%.+]], i32 0, i32 0
661 // CHECK-DAG: [[BPADDR0:%.+]] = getelementptr inbounds [3 x ptr], ptr [[BP]], i32 0, i32 0
662 // CHECK-DAG: [[PADDR0:%.+]] = getelementptr inbounds [3 x ptr], ptr [[P]], i32 0, i32 0
663 // CHECK-DAG: store i[[SZ]] [[VAL0:%[^,]+]], ptr [[BPADDR0]]
664 // CHECK-DAG: store i[[SZ]] [[VAL0]], ptr [[PADDR0]]
666 // CHECK-DAG: [[BPADDR1:%.+]] = getelementptr inbounds [3 x ptr], ptr [[BP]], i32 0, i32 1
667 // CHECK-DAG: [[PADDR1:%.+]] = getelementptr inbounds [3 x ptr], ptr [[P]], i32 0, i32 1
668 // CHECK-DAG: store i[[SZ]] [[VAL1:%[^,]+]], ptr [[BPADDR1]]
669 // CHECK-DAG: store i[[SZ]] [[VAL1]], ptr [[PADDR1]]
671 // CHECK-DAG: [[BPADDR2:%.+]] = getelementptr inbounds [3 x ptr], ptr [[BP]], i32 0, i32 2
672 // CHECK-DAG: [[PADDR2:%.+]] = getelementptr inbounds [3 x ptr], ptr [[P]], i32 0, i32 2
673 // CHECK-DAG: store ptr [[VAL2:%[^,]+]], ptr [[BPADDR2]]
674 // CHECK-DAG: store ptr [[VAL2]], ptr [[PADDR2]]
676 // CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0
677 // CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]]
679 // CHECK: call void [[HVT5:@.+]]({{[^,]+}}, {{[^,]+}}, {{[^,]+}})
680 // CHECK-NEXT: br label %[[END]]
682 // CHECK-NEXT: br label %[[IFEND:.+]]
684 // CHECK: call void [[HVT5]]({{[^,]+}}, {{[^,]+}}, {{[^,]+}})
685 // CHECK-NEXT: br label %[[IFEND]]
689 // OMP45: define internal void @__omp_offloading_{{.+}}_{{.+}}bar{{.+}}_l{{[0-9]+}}(i[[SZ]] noundef %{{.+}})
691 // OMP45: define {{.*}}@{{.*}}zee{{.*}}
693 // OMP45: [[LOCAL_THIS:%.+]] = alloca ptr
694 // OMP45: [[BP:%.+]] = alloca [1 x ptr]
695 // OMP45: [[P:%.+]] = alloca [1 x ptr]
696 // OMP45: [[LOCAL_THIS1:%.+]] = load ptr, ptr [[LOCAL_THIS]]
698 // OMP45: call void @__kmpc_critical(
699 // OMP45: [[ARR_IDX:%.+]] = getelementptr inbounds [[S2]], ptr [[LOCAL_THIS1]], i[[SZ]] 0
700 // OMP45: [[ARR_IDX2:%.+]] = getelementptr inbounds [[S2]], ptr [[LOCAL_THIS1]], i[[SZ]] 0
702 // OMP45-DAG: [[BPADDR0:%.+]] = getelementptr inbounds [1 x ptr], ptr [[BP]], i32 0, i32 0
703 // OMP45-DAG: [[PADDR0:%.+]] = getelementptr inbounds [1 x ptr], ptr [[P]], i32 0, i32 0
704 // OMP45-DAG: store ptr [[ARR_IDX]], ptr [[BPADDR0]]
705 // OMP45-DAG: store ptr [[ARR_IDX2]], ptr [[PADDR0]]
707 // OMP45: [[BPR:%.+]] = getelementptr inbounds [1 x ptr], ptr [[BP]], i32 0, i32 0
708 // OMP45: [[PR:%.+]] = getelementptr inbounds [1 x ptr], ptr [[P]], i32 0, i32 0
709 // 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]], 0
711 // OMP45-NEXT: br i1 [[ERROR]], label %[[FAIL:[^,]+]], label %[[END:[^,]+]]
713 // OMP45: call void [[HVT0:@.+]](ptr [[LOCAL_THIS1]])
714 // OMP45-NEXT: br label %[[END]]
716 // OMP45: call void @__kmpc_end_critical(
718 // Check that the offloading functions are emitted and that the arguments are
719 // correct and loaded correctly for the target regions of the callees of bar().
721 // CHECK: define internal void [[HVT7]]
722 // Create local storage for each capture.
723 // CHECK: [[LOCAL_THIS:%.+]] = alloca ptr
724 // 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 ptr
728 // 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]],
739 // CHECK-DAG: getelementptr inbounds nuw [[S1]], ptr [[REF_THIS]], i32 0, i32 0
740 // 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]] %{{.+}}
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 ptr
751 // 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]],
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]] 2
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 ptr
769 // 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]],
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]] 2
780 // OMP50: define internal void @__omp_offloading_{{.+}}_{{.+}}bar{{.+}}_l{{[0-9]+}}(i[[SZ]] noundef %{{.+}})
782 // OMP50: define {{.*}}@{{.*}}zee{{.*}}
784 // OMP50: [[LOCAL_THIS:%.+]] = alloca ptr
785 // 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]] 0
789 // OMP50: [[ARR_IDX2:%.+]] = getelementptr inbounds [[S2]], ptr [[LOCAL_THIS1]], i[[SZ]] 0
791 // OMP50-DAG: [[BPADDR0:%.+]] = getelementptr inbounds [1 x ptr], ptr [[BP]], i32 0, i32 0
792 // OMP50-DAG: [[PADDR0:%.+]] = getelementptr inbounds [1 x ptr], ptr [[P]], i32 0, i32 0
793 // OMP50-DAG: store ptr [[ARR_IDX]], ptr [[BPADDR0]]
794 // OMP50-DAG: store ptr [[ARR_IDX2]], ptr [[PADDR0]]
796 // OMP50: [[BPR:%.+]] = getelementptr inbounds [1 x ptr], ptr [[BP]], i32 0, i32 0
797 // OMP50: [[PR:%.+]] = getelementptr inbounds [1 x ptr], ptr [[P]], i32 0, i32 0
798 // 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]], 0
800 // OMP50-NEXT: br i1 [[ERROR]], label %[[FAIL:[^,]+]], label %[[END:[^,]+]]
802 // OMP50: call void [[HVT0:@.+]](ptr [[LOCAL_THIS1]])
803 // OMP50-NEXT: br label %[[END]]
807 #define pragma_target _Pragma("omp target")
811 #pragma omp parallel shared(global)
822 #pragma omp target map(this[0])
828 void thread_limit_target(int TargetTL
, int TeamsTL
) {
832 // OMP51: call i32 @__tgt_target_kernel({{.*}}, i64 -1, i32 -1, i32 0,
837 // OMP51: call i32 @__tgt_target_kernel({{.*}}, i64 -1, i32 0, i32 0,
839 #pragma omp target thread_limit(TargetTL)
841 // OMP51: [[TL:%.*]] = load {{.*}} %TargetTL.addr
842 // 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]]
848 #pragma omp target thread_limit(TargetTL)
851 // OMP51: [[TL:%.*]] = load {{.*}} %TargetTL.addr
852 // 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]]
858 #pragma omp teams thread_limit(TeamsTL)
860 // OMP51: load {{.*}} %TeamsTL.addr
861 // OMP51: [[TeamsL:%.*]] = load {{.*}} %TeamsTL.addr
862 // OMP51: call i32 @__tgt_target_kernel({{.*}}, i64 -1, i32 0, i32 [[TeamsL]],
864 #pragma omp target thread_limit(TargetTL)
865 #pragma omp teams thread_limit(TeamsTL)
867 // OMP51: load {{.*}} %TeamsTL.addr
868 // OMP51: [[TeamsL:%.*]] = load {{.*}} %TeamsTL.addr
869 // OMP51: call ptr @__kmpc_omp_task_alloc({{.*@.omp_task_entry.*}})
870 // OMP51: call i32 [[OMP_TASK_ENTRY]]
874 // Check that the offloading functions are called after setting thread_limit in the task entry functions
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,
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,
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,