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;
operator *TT30 TT<tx, ty> operator*(const TT<tx, ty> &) { return *this; }
31 };
32
33 // TCHECK: [[S1:%.+]] = type { double }
34
foo(int n)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>
ftemplate(int n)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
fstatic(int n)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
r1S1144 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
bar(int n)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