1 | // Only test codegen on target side, as private clause does not require any action on the host side |
2 | // Test target codegen - host bc file has to be created first. |
3 | // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm-bc %s -o %t-ppc-host.bc |
4 | // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --check-prefix TCHECK --check-prefix TCHECK-64 |
5 | // RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o %t %s |
6 | // RUN: %clang_cc1 -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix TCHECK --check-prefix TCHECK-64 |
7 | // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm-bc %s -o %t-x86-host.bc |
8 | // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s --check-prefix TCHECK --check-prefix TCHECK-32 |
9 | // RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o %t %s |
10 | // RUN: %clang_cc1 -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix TCHECK --check-prefix TCHECK-32 |
11 | |
12 | // RUN: %clang_cc1 -verify -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm-bc %s -o %t-ppc-host.bc |
13 | // RUN: %clang_cc1 -verify -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck --check-prefix SIMD-ONLY0 %s |
14 | // RUN: %clang_cc1 -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o %t %s |
15 | // RUN: %clang_cc1 -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s |
16 | // RUN: %clang_cc1 -verify -fopenmp-simd -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm-bc %s -o %t-x86-host.bc |
17 | // RUN: %clang_cc1 -verify -fopenmp-simd -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck --check-prefix SIMD-ONLY0 %s |
18 | // RUN: %clang_cc1 -fopenmp-simd -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o %t %s |
19 | // RUN: %clang_cc1 -fopenmp-simd -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s |
20 | // SIMD-ONLY0-NOT: {{__kmpc|__tgt}} |
21 | |
22 | // expected-no-diagnostics |
23 | #ifndef HEADER |
24 | #define HEADER |
25 | |
26 | template<typename tx, typename ty> |
27 | struct TT{ |
28 | tx X; |
29 | ty Y; |
30 | TT<tx, ty> operator*(const TT<tx, ty> &) { return *this; } |
31 | }; |
32 | |
33 | // TCHECK: [[S1:%.+]] = type { double } |
34 | |
35 | int foo(int n) { |
36 | int a = 0; |
37 | short aa = 0; |
38 | float b[10]; |
39 | float bn[n]; |
40 | double c[5][10]; |
41 | double cn[5][n]; |
42 | TT<long long, char> d; |
43 | |
44 | #pragma omp target reduction(*:a) |
45 | { |
46 | } |
47 | |
48 | // TCHECK: define weak void @__omp_offloading_{{.+}}(i32*{{.+}} %{{.+}}) |
49 | // TCHECK: [[A:%.+]] = alloca i{{[0-9]+}}*, |
50 | // TCHECK: store {{.+}}, {{.+}} [[A]], |
51 | // TCHECK: load i32*, i32** [[A]], |
52 | // TCHECK: ret void |
53 | |
54 | #pragma omp target reduction(+:a) |
55 | { |
56 | a = 1; |
57 | } |
58 | |
59 | // TCHECK: define weak void @__omp_offloading_{{.+}}(i32*{{.+}} %{{.+}}) |
60 | // TCHECK: [[A:%.+]] = alloca i{{[0-9]+}}*, |
61 | // TCHECK: store {{.+}}, {{.+}} [[A]], |
62 | // TCHECK: [[REF:%.+]] = load i32*, i32** [[A]], |
63 | // TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[REF]], |
64 | // TCHECK: ret void |
65 | |
66 | #pragma omp target reduction(-:a, aa) |
67 | { |
68 | a = 1; |
69 | aa = 1; |
70 | } |
71 | |
72 | // TCHECK: define weak void @__omp_offloading_{{.+}}(i32*{{.+}} [[A:%.+]], i16*{{.+}} [[AA:%.+]]) |
73 | // TCHECK: [[A:%.+]] = alloca i{{[0-9]+}}*, |
74 | // TCHECK: [[AA:%.+]] = alloca i{{[0-9]+}}*, |
75 | // TCHECK: store {{.+}}, {{.+}} [[A]], |
76 | // TCHECK: store {{.+}}, {{.+}} [[AA]], |
77 | // TCHECK: [[A_REF:%.+]] = load i32*, i32** [[A]], |
78 | // TCHECK: [[AA_REF:%.+]] = load i16*, i16** [[AA]], |
79 | // TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[A_REF]], |
80 | // TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[AA_REF]], |
81 | // TCHECK: ret void |
82 | |
83 | return a; |
84 | } |
85 | |
86 | |
87 | template<typename tx> |
88 | tx ftemplate(int n) { |
89 | tx a = 0; |
90 | short aa = 0; |
91 | tx b[10]; |
92 | |
93 | #pragma omp target reduction(+:a,aa,b) |
94 | { |
95 | a = 1; |
96 | aa = 1; |
97 | b[2] = 1; |
98 | } |
99 | |
100 | return a; |
101 | } |
102 | |
103 | static |
104 | int fstatic(int n) { |
105 | int a = 0; |
106 | short aa = 0; |
107 | char aaa = 0; |
108 | int b[10]; |
109 | |
110 | #pragma omp target reduction(-:a,aa,aaa,b) |
111 | { |
112 | a = 1; |
113 | aa = 1; |
114 | aaa = 1; |
115 | b[2] = 1; |
116 | } |
117 | |
118 | return a; |
119 | } |
120 | |
121 | // TCHECK: define weak void @__omp_offloading_{{.+}}(i32*{{.+}}, i16*{{.+}}, i8*{{.+}}, [10 x i32]*{{.+}}) |
122 | // TCHECK: [[A:%.+]] = alloca i{{[0-9]+}}*, |
123 | // TCHECK: [[A2:%.+]] = alloca i{{[0-9]+}}*, |
124 | // TCHECK: [[A3:%.+]] = alloca i{{[0-9]+}}*, |
125 | // TCHECK: [[B:%.+]] = alloca [10 x i{{[0-9]+}}]*, |
126 | // TCHECK: store {{.+}}, {{.+}} [[A]], |
127 | // TCHECK: store {{.+}}, {{.+}} [[A2]], |
128 | // TCHECK: store {{.+}}, {{.+}} [[A3]], |
129 | // TCHECK: store {{.+}}, {{.+}} [[B]], |
130 | // TCHECK: [[A_REF:%.+]] = load i32*, i32** [[A]], |
131 | // TCHECK: [[AA_REF:%.+]] = load i16*, i16** [[AA]], |
132 | // TCHECK: [[A3_REF:%.+]] = load i8*, i8** [[A3]], |
133 | // TCHECK: [[B_REF:%.+]] = load {{.+}}*, {{.+}}** [[B]], |
134 | // TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[A_REF]], |
135 | // TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[AA_REF]], |
136 | // TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[A3_REF]], |
137 | // TCHECK: [[B_GEP:%.+]] = getelementptr inbounds [10 x i{{[0-9]+}}], [10 x i{{[0-9]+}}]* [[B_REF]], i{{[0-9]+}} 0, i{{[0-9]+}} 2 |
138 | // TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[B_GEP]], |
139 | // TCHECK: ret void |
140 | |
141 | struct S1 { |
142 | double a; |
143 | |
144 | int r1(int n){ |
145 | int b = n+1; |
146 | short int c[2][n]; |
147 | |
148 | #pragma omp target reduction(max:b,c) |
149 | { |
150 | this->a = (double)b + 1.5; |
151 | c[1][1] = ++a; |
152 | } |
153 | |
154 | return c[1][1] + (int)b; |
155 | } |
156 | |
157 | // TCHECK: define weak void @__omp_offloading_{{.+}}([[S1]]* [[TH:%.+]], i32*{{.+}}, i{{[0-9]+}} [[VLA:%.+]], i{{[0-9]+}} [[VLA1:%.+]], i16*{{.+}}) |
158 | // TCHECK: [[TH_ADDR:%.+]] = alloca [[S1]]*, |
159 | // TCHECK: [[B_ADDR:%.+]] = alloca i{{[0-9]+}}*, |
160 | // TCHECK: [[VLA_ADDR:%.+]] = alloca i{{[0-9]+}}, |
161 | // TCHECK: [[VLA_ADDR2:%.+]] = alloca i{{[0-9]+}}, |
162 | // TCHECK: [[C_ADDR:%.+]] = alloca i{{[0-9]+}}*, |
163 | // TCHECK: store [[S1]]* [[TH]], [[S1]]** [[TH_ADDR]], |
164 | // TCHECK: store i{{[0-9]+}}* {{.+}}, i{{[0-9]+}}** [[B_ADDR]], |
165 | // TCHECK: store i{{[0-9]+}} [[VLA]], i{{[0-9]+}}* [[VLA_ADDR]], |
166 | // TCHECK: store i{{[0-9]+}} [[VLA1]], i{{[0-9]+}}* [[VLA_ADDR2]], |
167 | // TCHECK: store i{{[0-9]+}}* {{.+}}, i{{[0-9]+}}** [[C_ADDR]], |
168 | // TCHECK: [[TH_ADDR_REF:%.+]] = load [[S1]]*, [[S1]]** [[TH_ADDR]], |
169 | // TCHECK: [[B_REF:%.+]] = load i{{[0-9]+}}*, i{{[0-9]+}}** [[B_ADDR]], |
170 | // TCHECK: [[VLA_ADDR_REF:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[VLA_ADDR]], |
171 | // TCHECK: [[VLA_ADDR_REF2:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[VLA_ADDR2]], |
172 | // TCHECK: [[C_REF:%.+]] = load i{{[0-9]+}}*, i{{[0-9]+}}** [[C_ADDR]], |
173 | |
174 | // this->a = (double)b + 1.5; |
175 | // TCHECK: [[B_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[B_REF]], |
176 | // TCHECK: [[B_CONV:%.+]] = sitofp i{{[0-9]+}} [[B_VAL]] to double |
177 | // TCHECK: [[NEW_A_VAL:%.+]] = fadd double [[B_CONV]], 1.5{{.+}}+00 |
178 | // TCHECK: [[A_FIELD:%.+]] = getelementptr inbounds [[S1]], [[S1]]* [[TH_ADDR_REF]], i{{[0-9]+}} 0, i{{[0-9]+}} 0 |
179 | // TCHECK: store double [[NEW_A_VAL]], double* [[A_FIELD]], |
180 | |
181 | // c[1][1] = ++a; |
182 | // TCHECK: [[A_FIELD4:%.+]] = getelementptr inbounds [[S1]], [[S1]]* [[TH_ADDR_REF]], i{{[0-9]+}} 0, i{{[0-9]+}} 0 |
183 | // TCHECK: [[A_FIELD4_VAL:%.+]] = load double, double* [[A_FIELD4]], |
184 | // TCHECK: [[A_FIELD_INC:%.+]] = fadd double [[A_FIELD4_VAL]], 1.0{{.+}}+00 |
185 | // TCHECK: store double [[A_FIELD_INC]], double* [[A_FIELD4]], |
186 | // TCHECK: [[A_FIELD_INC_CONV:%.+]] = fptosi double [[A_FIELD_INC]] to i{{[0-9]+}} |
187 | // TCHECK: [[C_IND:%.+]] = mul{{.+}} i{{[0-9]+}} 1, [[VLA_ADDR_REF2]] |
188 | // TCHECK: [[C_1_REF:%.+]] = getelementptr inbounds i{{[0-9]+}}, i{{[0-9]+}}* [[C_REF]], i{{[0-9]+}} [[C_IND]] |
189 | // TCHECK: [[C_1_1_REF:%.+]] = getelementptr inbounds i{{[0-9]+}}, i{{[0-9]+}}* [[C_1_REF]], i{{[0-9]+}} 1 |
190 | // TCHECK: store i{{[0-9]+}} [[A_FIELD_INC_CONV]], i{{[0-9]+}}* [[C_1_1_REF]], |
191 | |
192 | // finish |
193 | // TCHECK: ret void |
194 | }; |
195 | |
196 | |
197 | int bar(int n){ |
198 | int a = 0; |
199 | a += foo(n); |
200 | S1 S; |
201 | a += S.r1(n); |
202 | a += fstatic(n); |
203 | a += ftemplate<int>(n); |
204 | |
205 | return a; |
206 | } |
207 | |
208 | // template |
209 | // TCHECK: define weak void @__omp_offloading_{{.+}}(i{{[0-9]+}}*{{.+}}, i{{[0-9]+}}*{{.+}}, [10 x i32]*{{.+}}) |
210 | // TCHECK: [[A:%.+]] = alloca i{{[0-9]+}}*, |
211 | // TCHECK: [[A2:%.+]] = alloca i{{[0-9]+}}*, |
212 | // TCHECK: [[B:%.+]] = alloca [10 x i{{[0-9]+}}]*, |
213 | // TCHECK: store {{.+}}, {{.+}} [[A]], |
214 | // TCHECK: store {{.+}}, {{.+}} [[A2]], |
215 | // TCHECK: store {{.+}}, {{.+}} [[B]], |
216 | // TCHECK: [[A_REF:%.+]] = load i32*, i32** [[A]], |
217 | // TCHECK: [[AA_REF:%.+]] = load i16*, i16** [[AA]], |
218 | // TCHECK: [[B_REF:%.+]] = load {{.+}}*, {{.+}}** [[B]], |
219 | // TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[A_REF]], |
220 | // TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[AA_REF]], |
221 | // TCHECK: [[B_GEP:%.+]] = getelementptr inbounds [10 x i{{[0-9]+}}], [10 x i{{[0-9]+}}]* [[B_REF]], i{{[0-9]+}} 0, i{{[0-9]+}} 2 |
222 | // TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[B_GEP]], |
223 | // TCHECK: ret void |
224 | |
225 | #endif |
226 | |