1 // 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 -DLOAD
2 // RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - -DLOAD | FileCheck %s
3 // RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -emit-pch -o %t
4 // RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -include-pch %t -o - -DLOAD | FileCheck %s
5
6 // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - -DOMP5 | FileCheck %s --check-prefix HOST5
7 // 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 -DOMP5
8 // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - -DOMP5 | FileCheck %s --check-prefix DEV5
9
10 // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - -DOMP5 | FileCheck %s --check-prefix KMPC-ONLY
11
12 // RUN: %clang_cc1 -verify -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - -DOMP5 | FileCheck %s --check-prefix SIMD-ONLY
13 // 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 -DOMP5
14 // RUN: %clang_cc1 -verify -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - -DOMP5 | FileCheck %s --check-prefix SIMD-ONLY
15
16 // 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 -fopenmp-version=45
17 // RUN: %clang_cc1 -verify -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - -fopenmp-version=45 | FileCheck %s --check-prefix SIMD-ONLY
18 // RUN: %clang_cc1 -verify -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -emit-pch -o %t -fopenmp-version=45
19 // RUN: %clang_cc1 -verify -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -include-pch %t -verify -o - -fopenmp-version=45 | FileCheck %s --check-prefix SIMD-ONLY
20
21 // expected-no-diagnostics
22
23 // SIMD-ONLY-NOT: {{__kmpc|__tgt}}
24 // KMPC-ONLY-NOT: __tgt
25
26 // CHECK-NOT: define {{.*}}{{baz1|baz4|maini1|Base|virtual_}}
27 // CHECK-DAG: Bake
28 // CHECK-NOT: @{{hhh|ggg|fff|eee}} =
29 // CHECK-DAG: @aaa = external global i32,
30 // CHECK-DAG: @bbb ={{ hidden | }}global i32 0,
31 // CHECK-DAG: weak constant %struct.__tgt_offload_entry { i8* bitcast (i32* @bbb to i8*),
32 // CHECK-DAG: @ccc = external global i32,
33 // CHECK-DAG: @ddd ={{ hidden | }}global i32 0,
34 // CHECK-DAG: @hhh_decl_tgt_ref_ptr = weak global i32* null
35 // CHECK-DAG: @ggg_decl_tgt_ref_ptr = weak global i32* null
36 // CHECK-DAG: @fff_decl_tgt_ref_ptr = weak global i32* null
37 // CHECK-DAG: @eee_decl_tgt_ref_ptr = weak global i32* null
38 // CHECK-DAG: @{{.*}}maini1{{.*}}aaa = internal global i64 23,
39 // CHECK-DAG: @b ={{ hidden | }}global i32 15,
40 // CHECK-DAG: @d ={{ hidden | }}global i32 0,
41 // CHECK-DAG: @c = external global i32,
42 // CHECK-DAG: @globals ={{ hidden | }}global %struct.S zeroinitializer,
43 // CHECK-DAG: [[STAT:@.+stat]] = internal global %struct.S zeroinitializer,
44 // CHECK-DAG: [[STAT_REF:@.+]] = internal constant %struct.S* [[STAT]]
45 // CHECK-DAG: @out_decl_target ={{ hidden | }}global i32 0,
46 // CHECK-DAG: @llvm.used = appending global [2 x i8*] [i8* bitcast (void ()* @__omp_offloading__{{.+}}_globals_l[[@LINE+84]]_ctor to i8*), i8* bitcast (void ()* @__omp_offloading__{{.+}}_stat_l[[@LINE+85]]_ctor to i8*)],
47 // CHECK-DAG: @llvm.compiler.used = appending global [1 x i8*] [i8* bitcast (%struct.S** [[STAT_REF]] to i8*)],
48
49 // CHECK-DAG: define {{.*}}i32 @{{.*}}{{foo|bar|baz2|baz3|FA|f_method}}{{.*}}()
50 // CHECK-DAG: define {{.*}}void @{{.*}}TemplateClass{{.*}}(%class.TemplateClass* {{[^,]*}} %{{.*}})
51 // CHECK-DAG: define {{.*}}i32 @{{.*}}TemplateClass{{.*}}f_method{{.*}}(%class.TemplateClass* {{[^,]*}} %{{.*}})
52 // CHECK-DAG: define {{.*}}void @__omp_offloading__{{.*}}_globals_l[[@LINE+78]]_ctor()
53
54 #ifndef HEADER
55 #define HEADER
56
57 #pragma omp declare target
58 extern int bbb;
59 #pragma omp end declare target
60 #pragma omp declare target
61 extern int bbb;
62 #pragma omp end declare target
63
64 #pragma omp declare target
65 extern int aaa;
66 int bbb = 0;
67 extern int ccc;
68 int ddd = 0;
69 #pragma omp end declare target
70
71 #pragma omp declare target
72 extern int bbb;
73 #pragma omp end declare target
74
75 extern int eee;
76 int fff = 0;
77 extern int ggg;
78 int hhh = 0;
79 #pragma omp declare target link(eee, fff, ggg, hhh)
80
81 int out_decl_target = 0;
82 #ifdef OMP5
83 #pragma omp declare target(out_decl_target)
84 #endif
85
86 #pragma omp declare target
lambda()87 void lambda () {
88 #ifdef __cpp_lambdas
89 (void)[&] { (void)out_decl_target; };
90 #else
91 #pragma clang __debug captured
92 {
93 (void)out_decl_target;
94 }
95 #endif
96 };
97 #pragma omp end declare target
98
99 template <typename T>
100 class TemplateClass {
101 T a;
102 public:
TemplateClass()103 TemplateClass() {}
f_method() const104 T f_method() const { return a; }
105 };
106
107 int foo();
108
baz1()109 static int baz1() { return 0; }
110
111 int baz2();
112
baz4()113 int baz4() { return 5; }
114
115 template <typename T>
FA()116 T FA() {
117 TemplateClass<T> s;
118 return s.f_method();
119 }
120
121 #pragma omp declare target
122 struct S {
123 int a;
SS124 S(int a) : a(a) {}
125 };
126
foo()127 int foo() { return 0; }
128 int b = 15;
129 int d;
130 S globals(d);
131 static S stat(d);
132 #pragma omp end declare target
133 int c;
134
bar()135 int bar() { return 1 + foo() + bar() + baz1() + baz2(); }
136
maini1()137 int maini1() {
138 int a;
139 static long aa = 32 + bbb + ccc + fff + ggg;
140 // CHECK-DAG: define weak void @__omp_offloading_{{.*}}maini1{{.*}}_l[[@LINE+1]](i32* nonnull align {{[0-9]+}} dereferenceable({{[0-9]+}}) %{{.*}}, i64 {{.*}}, i64 {{.*}})
141 #pragma omp target map(tofrom \
142 : a, b)
143 {
144 S s(a);
145 static long aaa = 23;
146 a = foo() + bar() + b + c + d + aa + aaa + FA<int>();
147 }
148 return baz4();
149 }
150
baz3()151 int baz3() { return 2 + baz2(); }
baz2()152 int baz2() {
153 // CHECK-DAG: define weak void @__omp_offloading_{{.*}}baz2{{.*}}_l[[@LINE+1]](i64 {{.*}})
154 #pragma omp target parallel
155 ++c;
156 return 2 + baz3();
157 }
158
159 extern int create() throw();
160
161 static __typeof(create) __t_create __attribute__((__weakref__("__create")));
162
baz5()163 int baz5() {
164 bool a;
165 // CHECK-DAG: define weak void @__omp_offloading_{{.*}}baz5{{.*}}_l[[@LINE+1]](i64 {{.*}})
166 #pragma omp target
167 a = __extension__(void *) & __t_create != 0;
168 return a;
169 }
170
171 template <typename T>
172 struct Base {
~BaseBase173 virtual ~Base() {}
174 };
175
176 template class Base<int>;
177
178 template <typename T>
179 struct Bake {
~BakeBake180 virtual ~Bake() {}
181 };
182
183 #pragma omp declare target
184 template class Bake<int>;
185 #pragma omp end declare target
186
187 struct BaseNonT {
~BaseNonTBaseNonT188 virtual ~BaseNonT() {}
189 };
190
191 #pragma omp declare target
192 struct BakeNonT {
~BakeNonTBakeNonT193 virtual ~BakeNonT() {}
194 };
195 #pragma omp end declare target
196
197 template <typename T>
198 struct B {
199 virtual void virtual_foo();
200 };
201
new_bar()202 void new_bar() { new B<int>(); }
203
204 template <typename T>
virtual_foo()205 void B<T>::virtual_foo() {
206 #pragma omp target
207 {}
208 }
209
210 struct A {
emittedA211 virtual void emitted() {}
212 };
213
214 template <typename T>
215 struct C : public A {
216 virtual void emitted();
217 };
218
219 template <typename T>
emitted()220 void C<T>::emitted() {
221 #pragma omp target
222 {}
223 }
224
main()225 int main() {
226 A *X = new C<int>();
227 X->emitted();
228 return 0;
229 }
230
231 // CHECK-DAG: define {{.*}}void @__omp_offloading_{{.*}}virtual_foo{{.*}}_l[[@LINE-25]]()
232 // CHECK-DAG: define {{.*}}void @__omp_offloading_{{.*}}emitted{{.*}}_l[[@LINE-11]]()
233
234 template <typename T>
235 struct TTT {
emittedTTT236 virtual void emitted() {
237 #pragma omp target
238 {}
239 }
240 };
241
242 // CHECK-DAG: define {{.*}}void @__omp_offloading_{{.*}}emitted{{.*}}_l[[@LINE-5]]()
243
244 // CHECK-DAG: declare extern_weak signext i32 @__create()
245
246 // CHECK-NOT: define {{.*}}{{baz1|baz4|maini1|Base|virtual_}}
247
248 // CHECK-DAG: !{i32 1, !"aaa", i32 0, i32 {{[0-9]+}}}
249 // CHECK-DAG: !{i32 1, !"ccc", i32 0, i32 {{[0-9]+}}}
250 // CHECK-DAG: !{{{.+}}virtual_foo
251
252 #ifdef OMP5
host_fun()253 void host_fun() {}
254 #pragma omp declare target to(host_fun) device_type(host)
device_fun()255 void device_fun() {}
256 #pragma omp declare target to(device_fun) device_type(nohost)
257 // HOST5-NOT: define {{.*}}void {{.*}}device_fun{{.*}}
258 // HOST5: define {{.*}}void {{.*}}host_fun{{.*}}
259 // HOST5-NOT: define {{.*}}void {{.*}}device_fun{{.*}}
260
261 // DEV5-NOT: define {{.*}}void {{.*}}host_fun{{.*}}
262 // DEV5: define {{.*}}void {{.*}}device_fun{{.*}}
263 // DEV5-NOT: define {{.*}}void {{.*}}host_fun{{.*}}
264 #endif // OMP5
265
266 #endif // HEADER
267
268 #ifdef LOAD
269 #pragma omp declare target
new_bar1()270 void new_bar1() {
271 TTT<int> *X = new TTT<int>();
272 X->emitted();
273 }
274 #pragma omp end declare target
275 #endif
276