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