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