1 | // Test host codegen. |
2 | // RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CHECK --check-prefix CHECK-64 |
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 %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CHECK --check-prefix CHECK-64 |
5 | // RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CHECK --check-prefix CHECK-32 |
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 %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CHECK --check-prefix CHECK-32 |
8 | |
9 | // RUN: %clang_cc1 -verify -fopenmp-simd -fopenmp-version=45 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY0 %s |
10 | // RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=45 -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-simd -fopenmp-version=45 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY0 %s |
12 | // RUN: %clang_cc1 -verify -fopenmp-simd -fopenmp-version=45 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY0 %s |
13 | // RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=45 -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-simd -fopenmp-version=45 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY0 %s |
15 | // SIMD-ONLY0-NOT: {{__kmpc|__tgt}} |
16 | |
17 | // Test target codegen - host bc file has to be created first. |
18 | // RUN: %clang_cc1 -verify -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 |
19 | // RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix TCHECK --check-prefix TCHECK-64 |
20 | // 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-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o %t %s |
21 | // RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix TCHECK --check-prefix TCHECK-64 |
22 | // RUN: %clang_cc1 -verify -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 |
23 | // RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix TCHECK --check-prefix TCHECK-32 |
24 | // 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-device -fopenmp-host-ir-file-path %t-x86-host.bc -o %t %s |
25 | // RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix TCHECK --check-prefix TCHECK-32 |
26 | |
27 | // RUN: %clang_cc1 -verify -fopenmp-simd -fopenmp-version=45 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm-bc %s -o %t-ppc-host.bc |
28 | // RUN: %clang_cc1 -verify -fopenmp-simd -fopenmp-version=45 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY1 %s |
29 | // RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=45 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o %t %s |
30 | // RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=45 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY1 %s |
31 | // RUN: %clang_cc1 -verify -fopenmp-simd -fopenmp-version=45 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm-bc %s -o %t-x86-host.bc |
32 | // RUN: %clang_cc1 -verify -fopenmp-simd -fopenmp-version=45 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY1 %s |
33 | // RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=45 -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o %t %s |
34 | // RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=45 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY1 %s |
35 | // SIMD-ONLY1-NOT: {{__kmpc|__tgt}} |
36 | |
37 | // expected-no-diagnostics |
38 | #ifndef HEADER |
39 | #define HEADER |
40 | |
41 | // CHECK-DAG: [[TT:%.+]] = type { i64, i8 } |
42 | // CHECK-DAG: [[S1:%.+]] = type { double } |
43 | // CHECK-DAG: [[ENTTY:%.+]] = type { i8*, i8*, i[[SZ:32|64]], i32, i32 } |
44 | // CHECK-DAG: [[DEVTY:%.+]] = type { i8*, i8*, [[ENTTY]]*, [[ENTTY]]* } |
45 | // CHECK-DAG: [[DSCTY:%.+]] = type { i32, [[DEVTY]]*, [[ENTTY]]*, [[ENTTY]]* } |
46 | |
47 | // TCHECK: [[ENTTY:%.+]] = type { i8*, i8*, i{{32|64}}, i32, i32 } |
48 | |
49 | // CHECK-DAG: $[[REGFN:\.omp_offloading\..+]] = comdat |
50 | |
51 | // We have 8 target regions, but only 7 that actually will generate offloading |
52 | // code, only 6 will have mapped arguments, and only 4 have all-constant map |
53 | // sizes. |
54 | |
55 | // CHECK-DAG: [[SIZET2:@.+]] = private unnamed_addr constant [3 x i[[SZ]]] [i[[SZ]] 2, i[[SZ]] 4, i[[SZ]] 4] |
56 | // CHECK-DAG: [[MAPT2:@.+]] = private unnamed_addr constant [3 x i64] [i64 800, i64 800, i64 800] |
57 | // CHECK-DAG: [[SIZET3:@.+]] = private unnamed_addr constant [2 x i[[SZ]]] [i[[SZ]] 4, i[[SZ]] 2] |
58 | // CHECK-DAG: [[MAPT3:@.+]] = private unnamed_addr constant [2 x i64] [i64 800, i64 800] |
59 | // CHECK-DAG: [[MAPT4:@.+]] = private unnamed_addr constant [9 x i64] [i64 800, i64 547, i64 288, i64 547, i64 547, i64 288, i64 288, i64 547, i64 547] |
60 | // CHECK-DAG: [[SIZET5:@.+]] = private unnamed_addr constant [3 x i[[SZ]]] [i[[SZ]] 4, i[[SZ]] 2, i[[SZ]] 40] |
61 | // CHECK-DAG: [[MAPT5:@.+]] = private unnamed_addr constant [3 x i64] [i64 800, i64 800, i64 547] |
62 | // CHECK-DAG: [[SIZET6:@.+]] = private unnamed_addr constant [4 x i[[SZ]]] [i[[SZ]] 4, i[[SZ]] 2, i[[SZ]] 1, i[[SZ]] 40] |
63 | // CHECK-DAG: [[MAPT6:@.+]] = private unnamed_addr constant [4 x i64] [i64 800, i64 800, i64 800, i64 547] |
64 | // CHECK-DAG: [[MAPT7:@.+]] = private unnamed_addr constant [6 x i64] [i64 32, i64 281474976711171, i64 800, i64 288, i64 288, i64 547] |
65 | // CHECK-DAG: @{{.*}} = weak constant i8 0 |
66 | // CHECK-DAG: @{{.*}} = weak constant i8 0 |
67 | // CHECK-DAG: @{{.*}} = weak constant i8 0 |
68 | // CHECK-DAG: @{{.*}} = weak constant i8 0 |
69 | // CHECK-DAG: @{{.*}} = weak constant i8 0 |
70 | // CHECK-DAG: @{{.*}} = weak constant i8 0 |
71 | // CHECK-DAG: @{{.*}} = weak constant i8 0 |
72 | |
73 | // TCHECK: @{{.+}} = weak constant [[ENTTY]] |
74 | // TCHECK: @{{.+}} = weak constant [[ENTTY]] |
75 | // TCHECK: @{{.+}} = weak constant [[ENTTY]] |
76 | // TCHECK: @{{.+}} = weak constant [[ENTTY]] |
77 | // TCHECK: @{{.+}} = weak constant [[ENTTY]] |
78 | // TCHECK: @{{.+}} = weak constant [[ENTTY]] |
79 | // TCHECK: @{{.+}} = weak constant [[ENTTY]] |
80 | // TCHECK-NOT: @{{.+}} = weak constant [[ENTTY]] |
81 | |
82 | // Check if offloading descriptor is created. |
83 | // CHECK: [[ENTBEGIN:@.+]] = external constant [[ENTTY]] |
84 | // CHECK: [[ENTEND:@.+]] = external constant [[ENTTY]] |
85 | // CHECK: [[DEVBEGIN:@.+]] = extern_weak constant i8 |
86 | // CHECK: [[DEVEND:@.+]] = extern_weak constant i8 |
87 | // CHECK: [[IMAGES:@.+]] = internal unnamed_addr constant [1 x [[DEVTY]]] [{{.+}} { i8* [[DEVBEGIN]], i8* [[DEVEND]], [[ENTTY]]* [[ENTBEGIN]], [[ENTTY]]* [[ENTEND]] }], comdat($[[REGFN]]) |
88 | // CHECK: [[DESC:@.+]] = internal constant [[DSCTY]] { i32 1, [[DEVTY]]* getelementptr inbounds ([1 x [[DEVTY]]], [1 x [[DEVTY]]]* [[IMAGES]], i32 0, i32 0), [[ENTTY]]* [[ENTBEGIN]], [[ENTTY]]* [[ENTEND]] }, comdat($[[REGFN]]) |
89 | |
90 | // Check target registration is registered as a Ctor. |
91 | // CHECK: appending global [1 x { i32, void ()*, i8* }] [{ i32, void ()*, i8* } { i32 0, void ()* @[[REGFN]], i8* bitcast (void ()* @[[REGFN]] to i8*) }] |
92 | |
93 | |
94 | template<typename tx, typename ty> |
95 | struct TT{ |
96 | tx X; |
97 | ty Y; |
98 | }; |
99 | |
100 | // CHECK-LABEL: get_val |
101 | long long get_val() { return 0; } |
102 | |
103 | // CHECK: define {{.*}}[[FOO:@.+]]( |
104 | int foo(int n) { |
105 | int a = 0; |
106 | short aa = 0; |
107 | float b[10]; |
108 | float bn[n]; |
109 | double c[5][10]; |
110 | double cn[5][n]; |
111 | TT<long long, char> d; |
112 | |
113 | // CHECK: [[RET:%.+]] = call i32 @__tgt_target_nowait(i64 -1, i8* @{{[^,]+}}, i32 0, i8** null, i8** null, i[[SZ]]* null, i64* null) |
114 | // CHECK-NEXT: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0 |
115 | // CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:[^,]+]], label %[[END:[^,]+]] |
116 | // CHECK: [[FAIL]] |
117 | // CHECK: call void [[HVT0:@.+]]() |
118 | // CHECK-NEXT: br label %[[END]] |
119 | // CHECK: [[END]] |
120 | #pragma omp target simd nowait |
121 | for (int i = 3; i < 32; i += 5) { |
122 | } |
123 | |
124 | // CHECK: call void [[HVT1:@.+]](i[[SZ]] {{[^,]+}}, i{{32|64}}{{[*]*}} {{[^)]+}}) |
125 | long long k = get_val(); |
126 | #pragma omp target simd if(target: 0) linear(k : 3) |
127 | for (int i = 10; i > 1; i--) { |
128 | a += 1; |
129 | } |
130 | |
131 | // CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target(i64 -1, i8* @{{[^,]+}}, i32 3, i8** [[BP:%[^,]+]], i8** [[P:%[^,]+]], i[[SZ]]* getelementptr inbounds ([3 x i[[SZ]]], [3 x i[[SZ]]]* [[SIZET2]], i32 0, i32 0), i64* getelementptr inbounds ([3 x i64], [3 x i64]* [[MAPT2]], i32 0, i32 0)) |
132 | // CHECK-DAG: [[BP]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[BPR:%[^,]+]], i32 0, i32 0 |
133 | // CHECK-DAG: [[P]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[PR:%[^,]+]], i32 0, i32 0 |
134 | // CHECK-DAG: [[BPADDR0:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[BPR]], i32 0, i32 0 |
135 | // CHECK-DAG: [[PADDR0:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[PR]], i32 0, i32 0 |
136 | // CHECK-DAG: [[CBPADDR0:%.+]] = bitcast i8** [[BPADDR0]] to i[[SZ]]* |
137 | // CHECK-DAG: [[CPADDR0:%.+]] = bitcast i8** [[PADDR0]] to i[[SZ]]* |
138 | // CHECK-DAG: store i[[SZ]] [[VAL0:%.+]], i[[SZ]]* [[CBPADDR0]], |
139 | // CHECK-DAG: store i[[SZ]] [[VAL0]], i[[SZ]]* [[CPADDR0]], |
140 | // CHECK-DAG: [[BPADDR1:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[BPR]], i32 0, i32 1 |
141 | // CHECK-DAG: [[PADDR1:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[PR]], i32 0, i32 1 |
142 | // CHECK-DAG: [[CBPADDR1:%.+]] = bitcast i8** [[BPADDR1]] to i[[SZ]]* |
143 | // CHECK-DAG: [[CPADDR1:%.+]] = bitcast i8** [[PADDR1]] to i[[SZ]]* |
144 | // CHECK-DAG: store i[[SZ]] [[VAL1:%.+]], i[[SZ]]* [[CBPADDR1]], |
145 | // CHECK-DAG: store i[[SZ]] [[VAL1]], i[[SZ]]* [[CPADDR1]], |
146 | // CHECK-DAG: [[BPADDR2:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[BPR]], i32 0, i32 1 |
147 | // CHECK-DAG: [[PADDR2:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[PR]], i32 0, i32 1 |
148 | // CHECK-DAG: [[CBPADDR2:%.+]] = bitcast i8** [[BPADDR2]] to i[[SZ]]* |
149 | // CHECK-DAG: [[CPADDR2:%.+]] = bitcast i8** [[PADDR2]] to i[[SZ]]* |
150 | // CHECK-DAG: store i[[SZ]] [[VAL2:%.+]], i[[SZ]]* [[CBPADDR2]], |
151 | // CHECK-DAG: store i[[SZ]] [[VAL2]], i[[SZ]]* [[CPADDR2]], |
152 | |
153 | // CHECK-NEXT: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0 |
154 | // CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:[^,]+]], label %[[END:[^,]+]] |
155 | // CHECK: [[FAIL]] |
156 | // CHECK: call void [[HVT2:@.+]](i[[SZ]] {{[^,]+}}, i[[SZ]] {{[^)]+}}) |
157 | // CHECK-NEXT: br label %[[END]] |
158 | // CHECK: [[END]] |
159 | int lin = 12; |
160 | #pragma omp target simd if(target: 1) linear(lin, a : get_val()) |
161 | for (unsigned long long it = 2000; it >= 600; it-=400) { |
162 | aa += 1; |
163 | } |
164 | |
165 | // CHECK: [[IF:%.+]] = icmp sgt i32 {{[^,]+}}, 10 |
166 | // CHECK: br i1 [[IF]], label %[[IFTHEN:[^,]+]], label %[[IFELSE:[^,]+]] |
167 | // CHECK: [[IFTHEN]] |
168 | // CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target(i64 -1, i8* @{{[^,]+}}, i32 2, i8** [[BPR:%[^,]+]], i8** [[PR:%[^,]+]], i[[SZ]]* getelementptr inbounds ([2 x i[[SZ]]], [2 x i[[SZ]]]* [[SIZET3]], i32 0, i32 0), i64* getelementptr inbounds ([2 x i64], [2 x i64]* [[MAPT3]], i32 0, i32 0)) |
169 | // CHECK-DAG: [[BPR]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[BP:%[^,]+]], i32 0, i32 0 |
170 | // CHECK-DAG: [[PR]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[P:%[^,]+]], i32 0, i32 0 |
171 | |
172 | // CHECK-DAG: [[BPADDR0:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[BP]], i32 0, i32 0 |
173 | // CHECK-DAG: [[PADDR0:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[P]], i32 0, i32 0 |
174 | // CHECK-DAG: [[CBPADDR0:%.+]] = bitcast i8** [[BPADDR0]] to i[[SZ]]* |
175 | // CHECK-DAG: [[CPADDR0:%.+]] = bitcast i8** [[PADDR0]] to i[[SZ]]* |
176 | // CHECK-DAG: store i[[SZ]] [[VAL0:%.+]], i[[SZ]]* [[CBPADDR0]], |
177 | // CHECK-DAG: store i[[SZ]] [[VAL0]], i[[SZ]]* [[CPADDR0]], |
178 | |
179 | // CHECK-DAG: [[BPADDR1:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[BP]], i32 0, i32 1 |
180 | // CHECK-DAG: [[PADDR1:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[P]], i32 0, i32 1 |
181 | // CHECK-DAG: [[CBPADDR1:%.+]] = bitcast i8** [[BPADDR1]] to i[[SZ]]* |
182 | // CHECK-DAG: [[CPADDR1:%.+]] = bitcast i8** [[PADDR1]] to i[[SZ]]* |
183 | // CHECK-DAG: store i[[SZ]] [[VAL1:%.+]], i[[SZ]]* [[CBPADDR1]], |
184 | // CHECK-DAG: store i[[SZ]] [[VAL1]], i[[SZ]]* [[CPADDR1]], |
185 | // CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0 |
186 | // CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]] |
187 | // CHECK: [[FAIL]] |
188 | // CHECK: call void [[HVT3:@.+]]({{[^,]+}}, {{[^,]+}}) |
189 | // CHECK-NEXT: br label %[[END]] |
190 | // CHECK: [[END]] |
191 | // CHECK-NEXT: br label %[[IFEND:.+]] |
192 | // CHECK: [[IFELSE]] |
193 | // CHECK: call void [[HVT3]]({{[^,]+}}, {{[^,]+}}) |
194 | // CHECK-NEXT: br label %[[IFEND]] |
195 | // CHECK: [[IFEND]] |
196 | |
197 | #pragma omp target simd if(target: n>10) |
198 | for (short it = 6; it <= 20; it-=-4) { |
199 | a += 1; |
200 | aa += 1; |
201 | } |
202 | |
203 | // We capture 3 VLA sizes in this target region |
204 | // CHECK-64: [[A_VAL:%.+]] = load i32, i32* %{{.+}}, |
205 | // CHECK-64: [[A_ADDR:%.+]] = bitcast i[[SZ]]* [[A_CADDR:%.+]] to i32* |
206 | // CHECK-64: store i32 [[A_VAL]], i32* [[A_ADDR]], |
207 | // CHECK-64: [[A_CVAL:%.+]] = load i[[SZ]], i[[SZ]]* [[A_CADDR]], |
208 | |
209 | // CHECK-32: [[A_VAL:%.+]] = load i32, i32* %{{.+}}, |
210 | // CHECK-32: store i32 [[A_VAL]], i32* [[A_CADDR:%.+]], |
211 | // CHECK-32: [[A_CVAL:%.+]] = load i[[SZ]], i[[SZ]]* [[A_CADDR]], |
212 | |
213 | // CHECK: [[IF:%.+]] = icmp sgt i32 {{[^,]+}}, 20 |
214 | // CHECK: br i1 [[IF]], label %[[TRY:[^,]+]], label %[[FAIL:[^,]+]] |
215 | // CHECK: [[TRY]] |
216 | // CHECK: [[BNSIZE:%.+]] = mul nuw i[[SZ]] [[VLA0:%.+]], 4 |
217 | // CHECK: [[CNELEMSIZE2:%.+]] = mul nuw i[[SZ]] 5, [[VLA1:%.+]] |
218 | // CHECK: [[CNSIZE:%.+]] = mul nuw i[[SZ]] [[CNELEMSIZE2]], 8 |
219 | |
220 | // CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target(i64 -1, i8* @{{[^,]+}}, i32 9, i8** [[BPR:%[^,]+]], i8** [[PR:%[^,]+]], i[[SZ]]* [[SR:%[^,]+]], i64* getelementptr inbounds ([9 x i64], [9 x i64]* [[MAPT4]], i32 0, i32 0)) |
221 | // CHECK-DAG: [[BPR]] = getelementptr inbounds [9 x i8*], [9 x i8*]* [[BP:%[^,]+]], i32 0, i32 0 |
222 | // CHECK-DAG: [[PR]] = getelementptr inbounds [9 x i8*], [9 x i8*]* [[P:%[^,]+]], i32 0, i32 0 |
223 | // CHECK-DAG: [[SR]] = getelementptr inbounds [9 x i[[SZ]]], [9 x i[[SZ]]]* [[S:%[^,]+]], i32 0, i32 0 |
224 | |
225 | // CHECK-DAG: [[SADDR0:%.+]] = getelementptr inbounds [9 x i[[SZ]]], [9 x i[[SZ]]]* [[S]], i32 0, i32 [[IDX0:[0-9]+]] |
226 | // CHECK-DAG: [[BPADDR0:%.+]] = getelementptr inbounds [9 x i8*], [9 x i8*]* [[BP]], i32 0, i32 [[IDX0]] |
227 | // CHECK-DAG: [[PADDR0:%.+]] = getelementptr inbounds [9 x i8*], [9 x i8*]* [[P]], i32 0, i32 [[IDX0]] |
228 | // CHECK-DAG: [[SADDR1:%.+]] = getelementptr inbounds [9 x i[[SZ]]], [9 x i[[SZ]]]* [[S]], i32 0, i32 [[IDX1:[0-9]+]] |
229 | // CHECK-DAG: [[BPADDR1:%.+]] = getelementptr inbounds [9 x i8*], [9 x i8*]* [[BP]], i32 0, i32 [[IDX1]] |
230 | // CHECK-DAG: [[PADDR1:%.+]] = getelementptr inbounds [9 x i8*], [9 x i8*]* [[P]], i32 0, i32 [[IDX1]] |
231 | // CHECK-DAG: [[SADDR2:%.+]] = getelementptr inbounds [9 x i[[SZ]]], [9 x i[[SZ]]]* [[S]], i32 0, i32 [[IDX2:[0-9]+]] |
232 | // CHECK-DAG: [[BPADDR2:%.+]] = getelementptr inbounds [9 x i8*], [9 x i8*]* [[BP]], i32 0, i32 [[IDX2]] |
233 | // CHECK-DAG: [[PADDR2:%.+]] = getelementptr inbounds [9 x i8*], [9 x i8*]* [[P]], i32 0, i32 [[IDX2]] |
234 | // CHECK-DAG: [[SADDR3:%.+]] = getelementptr inbounds [9 x i[[SZ]]], [9 x i[[SZ]]]* [[S]], i32 0, i32 [[IDX3:[0-9]+]] |
235 | // CHECK-DAG: [[BPADDR3:%.+]] = getelementptr inbounds [9 x i8*], [9 x i8*]* [[BP]], i32 0, i32 [[IDX3]] |
236 | // CHECK-DAG: [[PADDR3:%.+]] = getelementptr inbounds [9 x i8*], [9 x i8*]* [[P]], i32 0, i32 [[IDX3]] |
237 | // CHECK-DAG: [[SADDR4:%.+]] = getelementptr inbounds [9 x i[[SZ]]], [9 x i[[SZ]]]* [[S]], i32 0, i32 [[IDX4:[0-9]+]] |
238 | // CHECK-DAG: [[BPADDR4:%.+]] = getelementptr inbounds [9 x i8*], [9 x i8*]* [[BP]], i32 0, i32 [[IDX4]] |
239 | // CHECK-DAG: [[PADDR4:%.+]] = getelementptr inbounds [9 x i8*], [9 x i8*]* [[P]], i32 0, i32 [[IDX4]] |
240 | // CHECK-DAG: [[SADDR5:%.+]] = getelementptr inbounds [9 x i[[SZ]]], [9 x i[[SZ]]]* [[S]], i32 0, i32 [[IDX5:[0-9]+]] |
241 | // CHECK-DAG: [[BPADDR5:%.+]] = getelementptr inbounds [9 x i8*], [9 x i8*]* [[BP]], i32 0, i32 [[IDX5]] |
242 | // CHECK-DAG: [[PADDR5:%.+]] = getelementptr inbounds [9 x i8*], [9 x i8*]* [[P]], i32 0, i32 [[IDX5]] |
243 | // CHECK-DAG: [[SADDR6:%.+]] = getelementptr inbounds [9 x i[[SZ]]], [9 x i[[SZ]]]* [[S]], i32 0, i32 [[IDX6:[0-9]+]] |
244 | // CHECK-DAG: [[BPADDR6:%.+]] = getelementptr inbounds [9 x i8*], [9 x i8*]* [[BP]], i32 0, i32 [[IDX6]] |
245 | // CHECK-DAG: [[PADDR6:%.+]] = getelementptr inbounds [9 x i8*], [9 x i8*]* [[P]], i32 0, i32 [[IDX6]] |
246 | // CHECK-DAG: [[SADDR7:%.+]] = getelementptr inbounds [9 x i[[SZ]]], [9 x i[[SZ]]]* [[S]], i32 0, i32 [[IDX7:[0-9]+]] |
247 | // CHECK-DAG: [[BPADDR7:%.+]] = getelementptr inbounds [9 x i8*], [9 x i8*]* [[BP]], i32 0, i32 [[IDX7]] |
248 | // CHECK-DAG: [[PADDR7:%.+]] = getelementptr inbounds [9 x i8*], [9 x i8*]* [[P]], i32 0, i32 [[IDX7]] |
249 | // CHECK-DAG: [[SADDR8:%.+]] = getelementptr inbounds [9 x i[[SZ]]], [9 x i[[SZ]]]* [[S]], i32 0, i32 [[IDX8:[0-9]+]] |
250 | // CHECK-DAG: [[BPADDR8:%.+]] = getelementptr inbounds [9 x i8*], [9 x i8*]* [[BP]], i32 0, i32 [[IDX8]] |
251 | // CHECK-DAG: [[PADDR8:%.+]] = getelementptr inbounds [9 x i8*], [9 x i8*]* [[P]], i32 0, i32 [[IDX8]] |
252 | |
253 | // The names below are not necessarily consistent with the names used for the |
254 | // addresses above as some are repeated. |
255 | // CHECK-DAG: store i[[SZ]] [[VLA0]], i[[SZ]]* [[CBPADDR0:%.+]], |
256 | // CHECK-DAG: store i[[SZ]] [[VLA0]], i[[SZ]]* [[CPADDR0:%.+]], |
257 | // CHECK-DAG: [[CBPADDR0]] = bitcast i8** {{%[^,]+}} to i[[SZ]]* |
258 | // CHECK-DAG: [[CPADDR0]] = bitcast i8** {{%[^,]+}} to i[[SZ]]* |
259 | // CHECK-DAG: store i[[SZ]] {{4|8}}, i[[SZ]]* {{%[^,]+}} |
260 | |
261 | // CHECK-DAG: store i[[SZ]] [[VLA1]], i[[SZ]]* [[CBPADDR1:%.+]], |
262 | // CHECK-DAG: store i[[SZ]] [[VLA1]], i[[SZ]]* [[CPADDR1:%.+]], |
263 | // CHECK-DAG: [[CBPADDR1]] = bitcast i8** {{%[^,]+}} to i[[SZ]]* |
264 | // CHECK-DAG: [[CPADDR1]] = bitcast i8** {{%[^,]+}} to i[[SZ]]* |
265 | // CHECK-DAG: store i[[SZ]] {{4|8}}, i[[SZ]]* {{%[^,]+}} |
266 | |
267 | // CHECK-DAG: store i[[SZ]] 5, i[[SZ]]* [[CBPADDR2:%.+]], |
268 | // CHECK-DAG: store i[[SZ]] 5, i[[SZ]]* [[CPADDR2:%.+]], |
269 | // CHECK-DAG: [[CBPADDR2]] = bitcast i8** {{%[^,]+}} to i[[SZ]]* |
270 | // CHECK-DAG: [[CPADDR2]] = bitcast i8** {{%[^,]+}} to i[[SZ]]* |
271 | // CHECK-DAG: store i[[SZ]] {{4|8}}, i[[SZ]]* {{%[^,]+}} |
272 | |
273 | // CHECK-DAG: store i[[SZ]] [[A_CVAL]], i[[SZ]]* [[CBPADDR3:%.+]], |
274 | // CHECK-DAG: store i[[SZ]] [[A_CVAL]], i[[SZ]]* [[CPADDR3:%.+]], |
275 | // CHECK-DAG: [[CBPADDR3]] = bitcast i8** {{%[^,]+}} to i[[SZ]]* |
276 | // CHECK-DAG: [[CPADDR3]] = bitcast i8** {{%[^,]+}} to i[[SZ]]* |
277 | // CHECK-DAG: store i[[SZ]] 4, i[[SZ]]* {{%[^,]+}} |
278 | |
279 | // CHECK-DAG: store [10 x float]* %{{.+}}, [10 x float]** [[CBPADDR4:%.+]], |
280 | // CHECK-DAG: store [10 x float]* %{{.+}}, [10 x float]** [[CPADDR4:%.+]], |
281 | // CHECK-DAG: [[CBPADDR4]] = bitcast i8** {{%[^,]+}} to [10 x float]** |
282 | // CHECK-DAG: [[CPADDR4]] = bitcast i8** {{%[^,]+}} to [10 x float]** |
283 | // CHECK-DAG: store i[[SZ]] 40, i[[SZ]]* {{%[^,]+}} |
284 | |
285 | // CHECK-DAG: store float* %{{.+}}, float** [[CBPADDR5:%.+]], |
286 | // CHECK-DAG: store float* %{{.+}}, float** [[CPADDR5:%.+]], |
287 | // CHECK-DAG: [[CBPADDR5]] = bitcast i8** {{%[^,]+}} to float** |
288 | // CHECK-DAG: [[CPADDR5]] = bitcast i8** {{%[^,]+}} to float** |
289 | // CHECK-DAG: store i[[SZ]] [[BNSIZE]], i[[SZ]]* {{%[^,]+}} |
290 | |
291 | // CHECK-DAG: store [5 x [10 x double]]* %{{.+}}, [5 x [10 x double]]** [[CBPADDR6:%.+]], |
292 | // CHECK-DAG: store [5 x [10 x double]]* %{{.+}}, [5 x [10 x double]]** [[CPADDR6:%.+]], |
293 | // CHECK-DAG: [[CBPADDR6]] = bitcast i8** {{%[^,]+}} to [5 x [10 x double]]** |
294 | // CHECK-DAG: [[CPADDR6]] = bitcast i8** {{%[^,]+}} to [5 x [10 x double]]** |
295 | // CHECK-DAG: store i[[SZ]] 400, i[[SZ]]* {{%[^,]+}} |
296 | |
297 | // CHECK-DAG: store double* %{{.+}}, double** [[CBPADDR7:%.+]], |
298 | // CHECK-DAG: store double* %{{.+}}, double** [[CPADDR7:%.+]], |
299 | // CHECK-DAG: [[CBPADDR7]] = bitcast i8** {{%[^,]+}} to double** |
300 | // CHECK-DAG: [[CPADDR7]] = bitcast i8** {{%[^,]+}} to double** |
301 | // CHECK-DAG: store i[[SZ]] [[CNSIZE]], i[[SZ]]* {{%[^,]+}} |
302 | |
303 | // CHECK-DAG: store [[TT]]* %{{.+}}, [[TT]]** [[CBPADDR8:%.+]], |
304 | // CHECK-DAG: store [[TT]]* %{{.+}}, [[TT]]** [[CPADDR8:%.+]], |
305 | // CHECK-DAG: [[CBPADDR8]] = bitcast i8** {{%[^,]+}} to [[TT]]** |
306 | // CHECK-DAG: [[CPADDR8]] = bitcast i8** {{%[^,]+}} to [[TT]]** |
307 | // CHECK-DAG: store i[[SZ]] {{12|16}}, i[[SZ]]* {{%[^,]+}} |
308 | |
309 | // CHECK-NEXT: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0 |
310 | // CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:[^,]+]], label %[[END:[^,]+]] |
311 | |
312 | // CHECK: [[FAIL]] |
313 | // CHECK: call void [[HVT4:@.+]]({{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}}) |
314 | // CHECK-NEXT: br label %[[END]] |
315 | // CHECK: [[END]] |
316 | #pragma omp target simd if(target: n>20) |
317 | for (unsigned char it = 'z'; it >= 'a'; it+=-1) { |
318 | a += 1; |
319 | b[2] += 1.0; |
320 | bn[3] += 1.0; |
321 | c[1][2] += 1.0; |
322 | cn[1][3] += 1.0; |
323 | d.X += 1; |
324 | d.Y += 1; |
325 | } |
326 | |
327 | return a; |
328 | } |
329 | |
330 | // Check that the offloading functions are emitted and that the arguments are |
331 | // correct and loaded correctly for the target regions in foo(). |
332 | |
333 | // CHECK: define internal void [[HVT0]]() |
334 | // CHECK: !llvm.loop |
335 | // CHECK: ret void |
336 | // CHECK-NEXT: } |
337 | |
338 | |
339 | // CHECK: define internal void [[HVT1]](i[[SZ]] %{{.+}}, i{{32|64}}{{[*]*.*}} %{{.+}}) |
340 | // CHECK: [[AA_ADDR:%.+]] = alloca i[[SZ]], align |
341 | // CHECK: store i[[SZ]] %{{.+}}, i[[SZ]]* [[AA_ADDR]], align |
342 | // CHECK-64: [[AA_CADDR:%.+]] = bitcast i[[SZ]]* [[AA_ADDR]] to i32* |
343 | // CHECK-64: [[AA:%.+]] = load i32, i32* [[AA_CADDR]], align |
344 | // CHECK-32: [[AA:%.+]] = load i32, i32* [[AA_ADDR]], align |
345 | // CHECK: !llvm.access.group |
346 | // CHECK: !llvm.loop |
347 | // CHECK: ret void |
348 | // CHECK-NEXT: } |
349 | |
350 | // CHECK: define internal void [[HVT2]](i[[SZ]] %{{.+}}, i[[SZ]] %{{.+}}, i[[SZ]] %{{.+}}) |
351 | // CHECK: [[AA_ADDR:%.+]] = alloca i[[SZ]], align |
352 | // CHECK: store i[[SZ]] %{{.+}}, i[[SZ]]* [[AA_ADDR]], align |
353 | // CHECK: [[AA_CADDR:%.+]] = bitcast i[[SZ]]* [[AA_ADDR]] to i16* |
354 | // CHECK: [[AA:%.+]] = load i16, i16* [[AA_CADDR]], align |
355 | // CHECK: !llvm.loop |
356 | // CHECK: ret void |
357 | // CHECK-NEXT: } |
358 | |
359 | // CHECK: define internal void [[HVT3]] |
360 | // CHECK: [[A_ADDR:%.+]] = alloca i[[SZ]], align |
361 | // CHECK: [[AA_ADDR:%.+]] = alloca i[[SZ]], align |
362 | // CHECK-DAG: store i[[SZ]] %{{.+}}, i[[SZ]]* [[A_ADDR]], align |
363 | // CHECK-DAG: store i[[SZ]] %{{.+}}, i[[SZ]]* [[AA_ADDR]], align |
364 | // CHECK-64-DAG:[[A_CADDR:%.+]] = bitcast i[[SZ]]* [[A_ADDR]] to i32* |
365 | // CHECK-DAG: [[AA_CADDR:%.+]] = bitcast i[[SZ]]* [[AA_ADDR]] to i16* |
366 | // CHECK: !llvm.loop |
367 | // CHECK: ret void |
368 | // CHECK-NEXT: } |
369 | |
370 | // CHECK: define internal void [[HVT4]] |
371 | // Create local storage for each capture. |
372 | // CHECK: [[LOCAL_A:%.+]] = alloca i[[SZ]] |
373 | // CHECK: [[LOCAL_B:%.+]] = alloca [10 x float]* |
374 | // CHECK: [[LOCAL_VLA1:%.+]] = alloca i[[SZ]] |
375 | // CHECK: [[LOCAL_BN:%.+]] = alloca float* |
376 | // CHECK: [[LOCAL_C:%.+]] = alloca [5 x [10 x double]]* |
377 | // CHECK: [[LOCAL_VLA2:%.+]] = alloca i[[SZ]] |
378 | // CHECK: [[LOCAL_VLA3:%.+]] = alloca i[[SZ]] |
379 | // CHECK: [[LOCAL_CN:%.+]] = alloca double* |
380 | // CHECK: [[LOCAL_D:%.+]] = alloca [[TT]]* |
381 | // CHECK-DAG: store i[[SZ]] [[ARG_A:%.+]], i[[SZ]]* [[LOCAL_A]] |
382 | // CHECK-DAG: store [10 x float]* [[ARG_B:%.+]], [10 x float]** [[LOCAL_B]] |
383 | // CHECK-DAG: store i[[SZ]] [[ARG_VLA1:%.+]], i[[SZ]]* [[LOCAL_VLA1]] |
384 | // CHECK-DAG: store float* [[ARG_BN:%.+]], float** [[LOCAL_BN]] |
385 | // CHECK-DAG: store [5 x [10 x double]]* [[ARG_C:%.+]], [5 x [10 x double]]** [[LOCAL_C]] |
386 | // CHECK-DAG: store i[[SZ]] [[ARG_VLA2:%.+]], i[[SZ]]* [[LOCAL_VLA2]] |
387 | // CHECK-DAG: store i[[SZ]] [[ARG_VLA3:%.+]], i[[SZ]]* [[LOCAL_VLA3]] |
388 | // CHECK-DAG: store double* [[ARG_CN:%.+]], double** [[LOCAL_CN]] |
389 | // CHECK-DAG: store [[TT]]* [[ARG_D:%.+]], [[TT]]** [[LOCAL_D]] |
390 | |
391 | // CHECK-64-DAG:[[CONV_AP:%.+]] = bitcast i[[SZ]]* [[LOCAL_A]] to i32* |
392 | // CHECK-DAG: [[REF_B:%.+]] = load [10 x float]*, [10 x float]** [[LOCAL_B]], |
393 | // CHECK-DAG: [[VAL_VLA1:%.+]] = load i[[SZ]], i[[SZ]]* [[LOCAL_VLA1]], |
394 | // CHECK-DAG: [[REF_BN:%.+]] = load float*, float** [[LOCAL_BN]], |
395 | // CHECK-DAG: [[REF_C:%.+]] = load [5 x [10 x double]]*, [5 x [10 x double]]** [[LOCAL_C]], |
396 | // CHECK-DAG: [[VAL_VLA2:%.+]] = load i[[SZ]], i[[SZ]]* [[LOCAL_VLA2]], |
397 | // CHECK-DAG: [[VAL_VLA3:%.+]] = load i[[SZ]], i[[SZ]]* [[LOCAL_VLA3]], |
398 | // CHECK-DAG: [[REF_CN:%.+]] = load double*, double** [[LOCAL_CN]], |
399 | // CHECK-DAG: [[REF_D:%.+]] = load [[TT]]*, [[TT]]** [[LOCAL_D]], |
400 | |
401 | |
402 | template<typename tx> |
403 | tx ftemplate(int n) { |
404 | tx a = 0; |
405 | short aa = 0; |
406 | tx b[10]; |
407 | |
408 | #pragma omp target simd if(target: n>40) |
409 | for (long long i = -10; i < 10; i += 3) { |
410 | a += 1; |
411 | aa += 1; |
412 | b[2] += 1; |
413 | } |
414 | |
415 | return a; |
416 | } |
417 | |
418 | static |
419 | int fstatic(int n) { |
420 | int a = 0; |
421 | short aa = 0; |
422 | char aaa = 0; |
423 | int b[10]; |
424 | |
425 | #pragma omp target simd if(target: n>50) |
426 | for (unsigned i=100; i<10; i+=10) { |
427 | a += 1; |
428 | aa += 1; |
429 | aaa += 1; |
430 | b[2] += 1; |
431 | } |
432 | |
433 | return a; |
434 | } |
435 | |
436 | struct S1 { |
437 | double a; |
438 | |
439 | int r1(int n){ |
440 | int b = n+1; |
441 | short int c[2][n]; |
442 | |
443 | #pragma omp target simd if(target: n>60) |
444 | for (unsigned long long it = 2000; it >= 600; it -= 400) { |
445 | this->a = (double)b + 1.5; |
446 | c[1][1] = ++a; |
447 | } |
448 | |
449 | return c[1][1] + (int)b; |
450 | } |
451 | }; |
452 | |
453 | // CHECK: define {{.*}}@{{.*}}bar{{.*}} |
454 | int bar(int n){ |
455 | int a = 0; |
456 | |
457 | // CHECK: call {{.*}}i32 [[FOO]](i32 {{.*}}) |
458 | a += foo(n); |
459 | |
460 | S1 S; |
461 | // CHECK: call {{.*}}i32 [[FS1:@.+]]([[S1]]* {{.*}}, i32 {{.*}}) |
462 | a += S.r1(n); |
463 | |
464 | // CHECK: call {{.*}}i32 [[FSTATIC:@.+]](i32 {{.*}}) |
465 | a += fstatic(n); |
466 | |
467 | // CHECK: call {{.*}}i32 [[FTEMPLATE:@.+]](i32 {{.*}}) |
468 | a += ftemplate<int>(n); |
469 | |
470 | return a; |
471 | } |
472 | |
473 | // |
474 | // CHECK: define {{.*}}[[FS1]] |
475 | // |
476 | // CHECK: i8* @llvm.stacksave() |
477 | // CHECK-64: [[B_ADDR:%.+]] = bitcast i[[SZ]]* [[B_CADDR:%.+]] to i32* |
478 | // CHECK-64: store i32 %{{.+}}, i32* [[B_ADDR]], |
479 | // CHECK-64: [[B_CVAL:%.+]] = load i[[SZ]], i[[SZ]]* [[B_CADDR]], |
480 | |
481 | // CHECK-32: store i32 %{{.+}}, i32* %__vla_expr |
482 | // CHECK-32: store i32 %{{.+}}, i32* [[B_ADDR:%.+]], |
483 | // CHECK-32: [[B_CVAL:%.+]] = load i[[SZ]], i[[SZ]]* [[B_ADDR]], |
484 | |
485 | // CHECK: [[IF:%.+]] = icmp sgt i32 {{[^,]+}}, 60 |
486 | // CHECK: br i1 [[IF]], label %[[TRY:[^,]+]], label %[[FAIL:[^,]+]] |
487 | // CHECK: [[TRY]] |
488 | // We capture 2 VLA sizes in this target region |
489 | // CHECK: [[CELEMSIZE2:%.+]] = mul nuw i[[SZ]] 2, [[VLA0:%.+]] |
490 | // CHECK: [[CSIZE:%.+]] = mul nuw i[[SZ]] [[CELEMSIZE2]], 2 |
491 | |
492 | // CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target(i64 -1, i8* @{{[^,]+}}, i32 6, i8** [[BPR:%[^,]+]], i8** [[PR:%[^,]+]], i[[SZ]]* [[SR:%[^,]+]], i64* getelementptr inbounds ([6 x i64], [6 x i64]* [[MAPT7]], i32 0, i32 0)) |
493 | // CHECK-DAG: [[BPR]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BP:%.+]], i32 0, i32 0 |
494 | // CHECK-DAG: [[PR]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[P:%.+]], i32 0, i32 0 |
495 | // CHECK-DAG: [[SR]] = getelementptr inbounds [6 x i[[SZ]]], [6 x i[[SZ]]]* [[S:%.+]], i32 0, i32 0 |
496 | // CHECK-DAG: [[SADDR0:%.+]] = getelementptr inbounds [6 x i[[SZ]]], [6 x i[[SZ]]]* [[S]], i32 [[IDX0:[0-9]+]] |
497 | // CHECK-DAG: [[BPADDR0:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BP]], i32 [[IDX0]] |
498 | // CHECK-DAG: [[PADDR0:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[P]], i32 [[IDX0]] |
499 | // CHECK-DAG: [[SADDR1:%.+]] = getelementptr inbounds [6 x i[[SZ]]], [6 x i[[SZ]]]* [[S]], i32 [[IDX1:[0-9]+]] |
500 | // CHECK-DAG: [[BPADDR1:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BP]], i32 [[IDX1]] |
501 | // CHECK-DAG: [[PADDR1:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[P]], i32 [[IDX1]] |
502 | // CHECK-DAG: [[SADDR2:%.+]] = getelementptr inbounds [6 x i[[SZ]]], [6 x i[[SZ]]]* [[S]], i32 [[IDX2:[0-9]+]] |
503 | // CHECK-DAG: [[BPADDR2:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BP]], i32 [[IDX2]] |
504 | // CHECK-DAG: [[PADDR2:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[P]], i32 [[IDX2]] |
505 | // CHECK-DAG: [[SADDR3:%.+]] = getelementptr inbounds [6 x i[[SZ]]], [6 x i[[SZ]]]* [[S]], i32 [[IDX3:[0-9]+]] |
506 | // CHECK-DAG: [[BPADDR3:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BP]], i32 [[IDX3]] |
507 | // CHECK-DAG: [[PADDR3:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[P]], i32 [[IDX3]] |
508 | // CHECK-DAG: [[SADDR4:%.+]] = getelementptr inbounds [6 x i[[SZ]]], [6 x i[[SZ]]]* [[S]], i32 [[IDX4:[0-9]+]] |
509 | // CHECK-DAG: [[BPADDR4:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BP]], i32 [[IDX4]] |
510 | // CHECK-DAG: [[PADDR4:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[P]], i32 [[IDX4]] |
511 | // CHECK-DAG: [[SADDR5:%.+]] = getelementptr inbounds [6 x i[[SZ]]], [6 x i[[SZ]]]* [[S]], i32 [[IDX5:[0-9]+]] |
512 | // CHECK-DAG: [[BPADDR5:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BP]], i32 [[IDX5]] |
513 | // CHECK-DAG: [[PADDR5:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[P]], i32 [[IDX5]] |
514 | |
515 | // The names below are not necessarily consistent with the names used for the |
516 | // addresses above as some are repeated. |
517 | // CHECK-DAG: store i[[SZ]] [[VLA0]], i[[SZ]]* [[CBPADDR0:%.+]], |
518 | // CHECK-DAG: store i[[SZ]] [[VLA0]], i[[SZ]]* [[CPADDR0:%.+]], |
519 | // CHECK-DAG: [[CBPADDR0]] = bitcast i8** {{%[^,]+}} to i[[SZ]]* |
520 | // CHECK-DAG: [[CPADDR0]] = bitcast i8** {{%[^,]+}} to i[[SZ]]* |
521 | // CHECK-DAG: store i[[SZ]] {{4|8}}, i[[SZ]]* {{%[^,]+}} |
522 | |
523 | // CHECK-DAG: store i[[SZ]] 2, i[[SZ]]* [[CBPADDR1:%.+]], |
524 | // CHECK-DAG: store i[[SZ]] 2, i[[SZ]]* [[CPADDR1:%.+]], |
525 | // CHECK-DAG: [[CBPADDR1]] = bitcast i8** {{%[^,]+}} to i[[SZ]]* |
526 | // CHECK-DAG: [[CPADDR1]] = bitcast i8** {{%[^,]+}} to i[[SZ]]* |
527 | // CHECK-DAG: store i[[SZ]] {{4|8}}, i[[SZ]]* {{%[^,]+}} |
528 | |
529 | // CHECK-DAG: store i[[SZ]] [[B_CVAL]], i[[SZ]]* [[CBPADDR2:%.+]], |
530 | // CHECK-DAG: store i[[SZ]] [[B_CVAL]], i[[SZ]]* [[CPADDR2:%.+]], |
531 | // CHECK-DAG: [[CBPADDR2]] = bitcast i8** {{%[^,]+}} to i[[SZ]]* |
532 | // CHECK-DAG: [[CPADDR2]] = bitcast i8** {{%[^,]+}} to i[[SZ]]* |
533 | // CHECK-DAG: store i[[SZ]] 4, i[[SZ]]* {{%[^,]+}} |
534 | |
535 | // CHECK-DAG: store [[S1]]* %{{.+}}, [[S1]]** [[CBPADDR3:%.+]], |
536 | // CHECK-DAG: store double* %{{.+}}, double** [[CPADDR3:%.+]], |
537 | // CHECK-DAG: [[CBPADDR3]] = bitcast i8** {{%[^,]+}} to [[S1]]** |
538 | // CHECK-DAG: [[CPADDR3]] = bitcast i8** {{%[^,]+}} to double** |
539 | // CHECK-DAG: store i[[SZ]] 8, i[[SZ]]* {{%[^,]+}} |
540 | |
541 | // CHECK-DAG: store [[S1]]* %{{.+}}, [[S1]]** [[CBPADDR4:%.+]], |
542 | // CHECK-DAG: store double* %{{.+}}, double** [[CPADDR4:%.+]], |
543 | // CHECK-DAG: [[CBPADDR4]] = bitcast i8** {{%[^,]+}} to [[S1]]** |
544 | // CHECK-DAG: [[CPADDR4]] = bitcast i8** {{%[^,]+}} to double** |
545 | // CHECK-DAG: store i[[SZ]] 8, i[[SZ]]* {{%[^,]+}} |
546 | |
547 | // CHECK-DAG: store i16* %{{.+}}, i16** [[CBPADDR5:%.+]], |
548 | // CHECK-DAG: store i16* %{{.+}}, i16** [[CPADDR5:%.+]], |
549 | // CHECK-DAG: [[CBPADDR5]] = bitcast i8** {{%[^,]+}} to i16** |
550 | // CHECK-DAG: [[CPADDR5]] = bitcast i8** {{%[^,]+}} to i16** |
551 | // CHECK-DAG: store i[[SZ]] [[CSIZE]], i[[SZ]]* {{%[^,]+}} |
552 | |
553 | // CHECK-NEXT: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0 |
554 | // CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:[^,]+]], label %[[END:[^,]+]] |
555 | |
556 | // CHECK: [[FAIL]] |
557 | // CHECK: call void [[HVT7:@.+]]({{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}}) |
558 | // CHECK-NEXT: br label %[[END]] |
559 | // CHECK: [[END]] |
560 | |
561 | // |
562 | // CHECK: define {{.*}}[[FSTATIC]] |
563 | // |
564 | // CHECK: [[IF:%.+]] = icmp sgt i32 {{[^,]+}}, 50 |
565 | // CHECK: br i1 [[IF]], label %[[IFTHEN:[^,]+]], label %[[IFELSE:[^,]+]] |
566 | // CHECK: [[IFTHEN]] |
567 | // CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target(i64 -1, i8* @{{[^,]+}}, i32 4, i8** [[BPR:%[^,]+]], i8** [[PR:%[^,]+]], i[[SZ]]* getelementptr inbounds ([4 x i[[SZ]]], [4 x i[[SZ]]]* [[SIZET6]], i32 0, i32 0), i64* getelementptr inbounds ([4 x i64], [4 x i64]* [[MAPT6]], i32 0, i32 0)) |
568 | // CHECK-DAG: [[BPR]] = getelementptr inbounds [4 x i8*], [4 x i8*]* [[BP:%.+]], i32 0, i32 0 |
569 | // CHECK-DAG: [[PR]] = getelementptr inbounds [4 x i8*], [4 x i8*]* [[P:%.+]], i32 0, i32 0 |
570 | |
571 | // CHECK-DAG: [[BPADDR0:%.+]] = getelementptr inbounds [4 x i8*], [4 x i8*]* [[BP]], i32 0, i32 0 |
572 | // CHECK-DAG: [[PADDR0:%.+]] = getelementptr inbounds [4 x i8*], [4 x i8*]* [[P]], i32 0, i32 0 |
573 | // CHECK-DAG: [[CBPADDR0:%.+]] = bitcast i8** [[BPADDR0]] to i[[SZ]]* |
574 | // CHECK-DAG: [[CPADDR0:%.+]] = bitcast i8** [[PADDR0]] to i[[SZ]]* |
575 | // CHECK-DAG: store i[[SZ]] [[VAL0:%.+]], i[[SZ]]* [[CBPADDR0]], |
576 | // CHECK-DAG: store i[[SZ]] [[VAL0]], i[[SZ]]* [[CPADDR0]], |
577 | |
578 | // CHECK-DAG: [[BPADDR1:%.+]] = getelementptr inbounds [4 x i8*], [4 x i8*]* [[BP]], i32 0, i32 1 |
579 | // CHECK-DAG: [[PADDR1:%.+]] = getelementptr inbounds [4 x i8*], [4 x i8*]* [[P]], i32 0, i32 1 |
580 | // CHECK-DAG: [[CBPADDR1:%.+]] = bitcast i8** [[BPADDR1]] to i[[SZ]]* |
581 | // CHECK-DAG: [[CPADDR1:%.+]] = bitcast i8** [[PADDR1]] to i[[SZ]]* |
582 | // CHECK-DAG: store i[[SZ]] [[VAL1:%.+]], i[[SZ]]* [[CBPADDR1]], |
583 | // CHECK-DAG: store i[[SZ]] [[VAL1]], i[[SZ]]* [[CPADDR1]], |
584 | |
585 | // CHECK-DAG: [[BPADDR2:%.+]] = getelementptr inbounds [4 x i8*], [4 x i8*]* [[BP]], i32 0, i32 2 |
586 | // CHECK-DAG: [[PADDR2:%.+]] = getelementptr inbounds [4 x i8*], [4 x i8*]* [[P]], i32 0, i32 2 |
587 | // CHECK-DAG: [[CBPADDR2:%.+]] = bitcast i8** [[BPADDR2]] to i[[SZ]]* |
588 | // CHECK-DAG: [[CPADDR2:%.+]] = bitcast i8** [[PADDR2]] to i[[SZ]]* |
589 | // CHECK-DAG: store i[[SZ]] [[VAL2:%.+]], i[[SZ]]* [[CBPADDR2]], |
590 | // CHECK-DAG: store i[[SZ]] [[VAL2]], i[[SZ]]* [[CPADDR2]], |
591 | |
592 | // CHECK-DAG: [[BPADDR3:%.+]] = getelementptr inbounds [4 x i8*], [4 x i8*]* [[BP]], i32 0, i32 3 |
593 | // CHECK-DAG: [[PADDR3:%.+]] = getelementptr inbounds [4 x i8*], [4 x i8*]* [[P]], i32 0, i32 3 |
594 | // CHECK-DAG: [[CBPADDR3:%.+]] = bitcast i8** [[BPADDR3]] to [10 x i32]** |
595 | // CHECK-DAG: [[CPADDR3:%.+]] = bitcast i8** [[PADDR3]] to [10 x i32]** |
596 | // CHECK-DAG: store [10 x i32]* [[VAL3:%.+]], [10 x i32]** [[CBPADDR3]], |
597 | // CHECK-DAG: store [10 x i32]* [[VAL3]], [10 x i32]** [[CPADDR3]], |
598 | |
599 | // CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0 |
600 | // CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]] |
601 | // CHECK: [[FAIL]] |
602 | // CHECK: call void [[HVT6:@.+]]({{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}}) |
603 | // CHECK-NEXT: br label %[[END]] |
604 | // CHECK: [[END]] |
605 | // CHECK-NEXT: br label %[[IFEND:.+]] |
606 | // CHECK: [[IFELSE]] |
607 | // CHECK: call void [[HVT6]]({{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}}) |
608 | // CHECK-NEXT: br label %[[IFEND]] |
609 | // CHECK: [[IFEND]] |
610 | |
611 | // |
612 | // CHECK: define {{.*}}[[FTEMPLATE]] |
613 | // |
614 | // CHECK: [[IF:%.+]] = icmp sgt i32 {{[^,]+}}, 40 |
615 | // CHECK: br i1 [[IF]], label %[[IFTHEN:[^,]+]], label %[[IFELSE:[^,]+]] |
616 | // CHECK: [[IFTHEN]] |
617 | // CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target(i64 -1, i8* @{{[^,]+}}, i32 3, i8** [[BPR:%[^,]+]], i8** [[PR:%[^,]+]], i[[SZ]]* getelementptr inbounds ([3 x i[[SZ]]], [3 x i[[SZ]]]* [[SIZET5]], i32 0, i32 0), i64* getelementptr inbounds ([3 x i64], [3 x i64]* [[MAPT5]], i32 0, i32 0)) |
618 | // CHECK-DAG: [[BPR]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[BP:%.+]], i32 0, i32 0 |
619 | // CHECK-DAG: [[PR]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[P:%.+]], i32 0, i32 0 |
620 | |
621 | // CHECK-DAG: [[BPADDR0:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[BP]], i32 0, i32 0 |
622 | // CHECK-DAG: [[PADDR0:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[P]], i32 0, i32 0 |
623 | // CHECK-DAG: [[CBPADDR0:%.+]] = bitcast i8** [[BPADDR0]] to i[[SZ]]* |
624 | // CHECK-DAG: [[CPADDR0:%.+]] = bitcast i8** [[PADDR0]] to i[[SZ]]* |
625 | // CHECK-DAG: store i[[SZ]] [[VAL0:%.+]], i[[SZ]]* [[CBPADDR0]], |
626 | // CHECK-DAG: store i[[SZ]] [[VAL0]], i[[SZ]]* [[CPADDR0]], |
627 | |
628 | // CHECK-DAG: [[BPADDR1:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[BP]], i32 0, i32 1 |
629 | // CHECK-DAG: [[PADDR1:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[P]], i32 0, i32 1 |
630 | // CHECK-DAG: [[CBPADDR1:%.+]] = bitcast i8** [[BPADDR1]] to i[[SZ]]* |
631 | // CHECK-DAG: [[CPADDR1:%.+]] = bitcast i8** [[PADDR1]] to i[[SZ]]* |
632 | // CHECK-DAG: store i[[SZ]] [[VAL1:%.+]], i[[SZ]]* [[CBPADDR1]], |
633 | // CHECK-DAG: store i[[SZ]] [[VAL1]], i[[SZ]]* [[CPADDR1]], |
634 | |
635 | // CHECK-DAG: [[BPADDR2:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[BP]], i32 0, i32 2 |
636 | // CHECK-DAG: [[PADDR2:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[P]], i32 0, i32 2 |
637 | // CHECK-DAG: [[CBPADDR2:%.+]] = bitcast i8** [[BPADDR2]] to [10 x i32]** |
638 | // CHECK-DAG: [[CPADDR2:%.+]] = bitcast i8** [[PADDR2]] to [10 x i32]** |
639 | // CHECK-DAG: store [10 x i32]* [[VAL2:%.+]], [10 x i32]** [[CBPADDR2]], |
640 | // CHECK-DAG: store [10 x i32]* [[VAL2]], [10 x i32]** [[CPADDR2]], |
641 | |
642 | // CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0 |
643 | // CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]] |
644 | // CHECK: [[FAIL]] |
645 | // CHECK: call void [[HVT5:@.+]]({{[^,]+}}, {{[^,]+}}, {{[^,]+}}) |
646 | // CHECK-NEXT: br label %[[END]] |
647 | // CHECK: [[END]] |
648 | // CHECK-NEXT: br label %[[IFEND:.+]] |
649 | // CHECK: [[IFELSE]] |
650 | // CHECK: call void [[HVT:@.+]]({{[^,]+}}, {{[^,]+}}, {{[^,]+}}) |
651 | // CHECK-NEXT: br label %[[IFEND]] |
652 | // CHECK: [[IFEND]] |
653 | |
654 | // Check that the offloading functions are emitted and that the arguments are |
655 | // correct and loaded correctly for the target regions of the callees of bar(). |
656 | |
657 | // CHECK: define internal void [[HVT7]] |
658 | // Create local storage for each capture. |
659 | // CHECK: [[LOCAL_THIS:%.+]] = alloca [[S1]]* |
660 | // CHECK: [[LOCAL_B:%.+]] = alloca i[[SZ]] |
661 | // CHECK: [[LOCAL_VLA1:%.+]] = alloca i[[SZ]] |
662 | // CHECK: [[LOCAL_VLA2:%.+]] = alloca i[[SZ]] |
663 | // CHECK: [[LOCAL_C:%.+]] = alloca i16* |
664 | // CHECK-DAG: store [[S1]]* [[ARG_THIS:%.+]], [[S1]]** [[LOCAL_THIS]] |
665 | // CHECK-DAG: store i[[SZ]] [[ARG_B:%.+]], i[[SZ]]* [[LOCAL_B]] |
666 | // CHECK-DAG: store i[[SZ]] [[ARG_VLA1:%.+]], i[[SZ]]* [[LOCAL_VLA1]] |
667 | // CHECK-DAG: store i[[SZ]] [[ARG_VLA2:%.+]], i[[SZ]]* [[LOCAL_VLA2]] |
668 | // CHECK-DAG: store i16* [[ARG_C:%.+]], i16** [[LOCAL_C]] |
669 | // Store captures in the context. |
670 | // CHECK-DAG: [[REF_THIS:%.+]] = load [[S1]]*, [[S1]]** [[LOCAL_THIS]], |
671 | // CHECK-64-DAG:[[CONV_BP:%.+]] = bitcast i[[SZ]]* [[LOCAL_B]] to i32* |
672 | // CHECK-DAG: [[VAL_VLA1:%.+]] = load i[[SZ]], i[[SZ]]* [[LOCAL_VLA1]], |
673 | // CHECK-DAG: [[VAL_VLA2:%.+]] = load i[[SZ]], i[[SZ]]* [[LOCAL_VLA2]], |
674 | // CHECK-DAG: [[REF_C:%.+]] = load i16*, i16** [[LOCAL_C]], |
675 | |
676 | |
677 | // CHECK: define internal void [[HVT6]] |
678 | // Create local storage for each capture. |
679 | // CHECK: [[LOCAL_A:%.+]] = alloca i[[SZ]] |
680 | // CHECK: [[LOCAL_AA:%.+]] = alloca i[[SZ]] |
681 | // CHECK: [[LOCAL_AAA:%.+]] = alloca i[[SZ]] |
682 | // CHECK: [[LOCAL_B:%.+]] = alloca [10 x i32]* |
683 | // CHECK-DAG: store i[[SZ]] [[ARG_A:%.+]], i[[SZ]]* [[LOCAL_A]] |
684 | // CHECK-DAG: store i[[SZ]] [[ARG_AA:%.+]], i[[SZ]]* [[LOCAL_AA]] |
685 | // CHECK-DAG: store i[[SZ]] [[ARG_AAA:%.+]], i[[SZ]]* [[LOCAL_AAA]] |
686 | // CHECK-DAG: store [10 x i32]* [[ARG_B:%.+]], [10 x i32]** [[LOCAL_B]] |
687 | // Store captures in the context. |
688 | // CHECK-64-DAG:[[CONV_AP:%.+]] = bitcast i[[SZ]]* [[LOCAL_A]] to i32* |
689 | // CHECK-DAG: [[CONV_AAP:%.+]] = bitcast i[[SZ]]* [[LOCAL_AA]] to i16* |
690 | // CHECK-DAG: [[CONV_AAAP:%.+]] = bitcast i[[SZ]]* [[LOCAL_AAA]] to i8* |
691 | // CHECK-DAG: [[REF_B:%.+]] = load [10 x i32]*, [10 x i32]** [[LOCAL_B]], |
692 | |
693 | // CHECK: define internal void [[HVT5]] |
694 | // Create local storage for each capture. |
695 | // CHECK: [[LOCAL_A:%.+]] = alloca i[[SZ]] |
696 | // CHECK: [[LOCAL_AA:%.+]] = alloca i[[SZ]] |
697 | // CHECK: [[LOCAL_B:%.+]] = alloca [10 x i32]* |
698 | // CHECK-DAG: store i[[SZ]] [[ARG_A:%.+]], i[[SZ]]* [[LOCAL_A]] |
699 | // CHECK-DAG: store i[[SZ]] [[ARG_AA:%.+]], i[[SZ]]* [[LOCAL_AA]] |
700 | // CHECK-DAG: store [10 x i32]* [[ARG_B:%.+]], [10 x i32]** [[LOCAL_B]] |
701 | // Store captures in the context. |
702 | // CHECK-64-DAG:[[CONV_AP:%.+]] = bitcast i[[SZ]]* [[LOCAL_A]] to i32* |
703 | // CHECK-DAG: [[CONV_AAP:%.+]] = bitcast i[[SZ]]* [[LOCAL_AA]] to i16* |
704 | // CHECK-DAG: [[REF_B:%.+]] = load [10 x i32]*, [10 x i32]** [[LOCAL_B]], |
705 | |
706 | #endif |
707 | |