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