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