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: @flag = hidden global i8 undef,
30 // CHECK-DAG: @aaa = external global i32,
31 // CHECK-DAG: @bbb ={{ hidden | }}global i32 0,
32 // CHECK-DAG: weak constant %struct.__tgt_offload_entry { i8* bitcast (i32* @bbb to i8*),
33 // CHECK-DAG: @ccc = external global i32,
34 // CHECK-DAG: @ddd ={{ hidden | }}global i32 0,
35 // CHECK-DAG: @hhh_decl_tgt_ref_ptr = weak global i32* null
36 // CHECK-DAG: @ggg_decl_tgt_ref_ptr = weak global i32* null
37 // CHECK-DAG: @fff_decl_tgt_ref_ptr = weak global i32* null
38 // CHECK-DAG: @eee_decl_tgt_ref_ptr = weak global i32* null
39 // CHECK-DAG: @{{.*}}maini1{{.*}}aaa = internal global i64 23,
40 // CHECK-DAG: @pair = {{.*}}addrspace(3) global %struct.PAIR undef
41 // CHECK-DAG: @b ={{ hidden | }}global i32 15,
42 // CHECK-DAG: @d ={{ hidden | }}global i32 0,
43 // CHECK-DAG: @c = external global i32,
44 // CHECK-DAG: @globals ={{ hidden | }}global %struct.S zeroinitializer,
45 // CHECK-DAG: [[STAT:@.+stat]] = internal global %struct.S zeroinitializer,
46 // CHECK-DAG: [[STAT_REF:@.+]] = internal constant %struct.S* [[STAT]]
47 // CHECK-DAG: @out_decl_target ={{ hidden | }}global i32 0,
48 // 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*)],
49 // CHECK-DAG: @llvm.compiler.used = appending global [1 x i8*] [i8* bitcast (%struct.S** [[STAT_REF]] to i8*)],
50 
51 // CHECK-DAG: define {{.*}}i32 @{{.*}}{{foo|bar|baz2|baz3|FA|f_method}}{{.*}}()
52 // CHECK-DAG: define {{.*}}void @{{.*}}TemplateClass{{.*}}(%class.TemplateClass* {{[^,]*}} %{{.*}})
53 // CHECK-DAG: define {{.*}}i32 @{{.*}}TemplateClass{{.*}}f_method{{.*}}(%class.TemplateClass* {{[^,]*}} %{{.*}})
54 // CHECK-DAG: define {{.*}}void @__omp_offloading__{{.*}}_globals_l[[@LINE+78]]_ctor()
55 
56 #ifndef HEADER
57 #define HEADER
58 #pragma omp declare target
59 bool flag [[clang::loader_uninitialized]];
60 extern int bbb;
61 #pragma omp end declare target
62 #pragma omp declare target
63 extern int bbb;
64 #pragma omp end declare target
65 
66 #pragma omp declare target
67 extern int aaa;
68 int bbb = 0;
69 extern int ccc;
70 int ddd = 0;
71 #pragma omp end declare target
72 
73 #pragma omp declare target
74 extern int bbb;
75 #pragma omp end declare target
76 
77 extern int eee;
78 int fff = 0;
79 extern int ggg;
80 int hhh = 0;
81 #pragma omp declare target link(eee, fff, ggg, hhh)
82 
83 int out_decl_target = 0;
84 #ifdef OMP5
85 #pragma omp declare target(out_decl_target)
86 #endif
87 
88 #pragma omp declare target
lambda()89 void lambda () {
90 #ifdef __cpp_lambdas
91   (void)[&] { (void)out_decl_target; };
92 #else
93 #pragma clang __debug captured
94   {
95     (void)out_decl_target;
96   }
97 #endif
98 };
99 #pragma omp end declare target
100 
101 template <typename T>
102 class TemplateClass {
103   T a;
104 public:
TemplateClass()105   TemplateClass() {}
f_method() const106   T f_method() const { return a; }
107 };
108 
109 int foo();
110 
baz1()111 static int baz1() { return 0; }
112 
113 int baz2();
114 
baz4()115 int baz4() { return 5; }
116 
117 template <typename T>
FA()118 T FA() {
119   TemplateClass<T> s;
120   return s.f_method();
121 }
122 
123 #pragma omp declare target
124 struct S {
125   int a;
SS126   S(int a) : a(a) {}
127 };
128 
foo()129 int foo() { return 0; }
130 int b = 15;
131 int d;
132 S globals(d);
133 static S stat(d);
134 #pragma omp end declare target
135 int c;
136 
bar()137 int bar() { return 1 + foo() + bar() + baz1() + baz2(); }
138 
maini1()139 int maini1() {
140   int a;
141   static long aa = 32 + bbb + ccc + fff + ggg;
142 // CHECK-DAG: define weak void @__omp_offloading_{{.*}}maini1{{.*}}_l[[@LINE+1]](i32* nonnull align {{[0-9]+}} dereferenceable({{[0-9]+}}) %{{.*}}, i64 {{.*}}, i64 {{.*}})
143 #pragma omp target map(tofrom \
144                        : a, b)
145   {
146     S s(a);
147     static long aaa = 23;
148     a = foo() + bar() + b + c + d + aa + aaa + FA<int>();
149   }
150   return baz4();
151 }
152 
baz3()153 int baz3() { return 2 + baz2(); }
baz2()154 int baz2() {
155 // CHECK-DAG: define weak void @__omp_offloading_{{.*}}baz2{{.*}}_l[[@LINE+1]](i64 {{.*}})
156 #pragma omp target parallel
157   ++c;
158   return 2 + baz3();
159 }
160 
161 extern int create() throw();
162 
163 static __typeof(create) __t_create __attribute__((__weakref__("__create")));
164 
baz5()165 int baz5() {
166   bool a;
167 // CHECK-DAG: define weak void @__omp_offloading_{{.*}}baz5{{.*}}_l[[@LINE+1]](i64 {{.*}})
168 #pragma omp target
169   a = __extension__(void *) & __t_create != 0;
170   return a;
171 }
172 
173 template <typename T>
174 struct Base {
~BaseBase175   virtual ~Base() {}
176 };
177 
178 template class Base<int>;
179 
180 template <typename T>
181 struct Bake {
~BakeBake182   virtual ~Bake() {}
183 };
184 
185 #pragma omp declare target
186 template class Bake<int>;
187 #pragma omp end declare target
188 
189 struct BaseNonT {
~BaseNonTBaseNonT190   virtual ~BaseNonT() {}
191 };
192 
193 #pragma omp declare target
194 struct BakeNonT {
~BakeNonTBakeNonT195   virtual ~BakeNonT() {}
196 };
197 #pragma omp end declare target
198 
199 template <typename T>
200 struct B {
201   virtual void virtual_foo();
202 };
203 
new_bar()204 void new_bar() { new B<int>(); }
205 
206 template <typename T>
virtual_foo()207 void B<T>::virtual_foo() {
208 #pragma omp target
209   {}
210 }
211 
212 struct A {
emittedA213   virtual void emitted() {}
214 };
215 
216 template <typename T>
217 struct C : public A {
218   virtual void emitted();
219 };
220 
221 template <typename T>
emitted()222 void C<T>::emitted() {
223 #pragma omp target
224   {}
225 }
226 
main()227 int main() {
228   A *X = new C<int>();
229   X->emitted();
230   return 0;
231 }
232 
233 // CHECK-DAG: define {{.*}}void @__omp_offloading_{{.*}}virtual_foo{{.*}}_l[[@LINE-25]]()
234 // CHECK-DAG: define {{.*}}void @__omp_offloading_{{.*}}emitted{{.*}}_l[[@LINE-11]]()
235 
236 template <typename T>
237 struct TTT {
emittedTTT238   virtual void emitted() {
239 #pragma omp target
240   {}
241   }
242 };
243 
244 // CHECK-DAG: define {{.*}}void @__omp_offloading_{{.*}}emitted{{.*}}_l[[@LINE-5]]()
245 
246 // CHECK-DAG: declare extern_weak signext i32 @__create()
247 
248 // CHECK-NOT: define {{.*}}{{baz1|baz4|maini1|Base|virtual_}}
249 
250 // CHECK-DAG: !{i32 1, !"aaa", i32 0, i32 {{[0-9]+}}}
251 // CHECK-DAG: !{i32 1, !"ccc", i32 0, i32 {{[0-9]+}}}
252 // CHECK-DAG: !{{{.+}}virtual_foo
253 
254 #ifdef OMP5
host_fun()255 void host_fun() {}
256 #pragma omp declare target to(host_fun) device_type(host)
device_fun()257 void device_fun() {}
258 #pragma omp declare target to(device_fun) device_type(nohost)
259 // HOST5-NOT: define {{.*}}void {{.*}}device_fun{{.*}}
260 // HOST5: define {{.*}}void {{.*}}host_fun{{.*}}
261 // HOST5-NOT: define {{.*}}void {{.*}}device_fun{{.*}}
262 
263 // DEV5-NOT: define {{.*}}void {{.*}}host_fun{{.*}}
264 // DEV5: define {{.*}}void {{.*}}device_fun{{.*}}
265 // DEV5-NOT: define {{.*}}void {{.*}}host_fun{{.*}}
266 #endif // OMP5
267 
268 struct PAIR {
269   int a;
270   int b;
271 };
272 
273 #pragma omp declare target
274 PAIR pair __attribute__((address_space(3), loader_uninitialized));
275 #pragma omp end declare target
276 
277 #endif // HEADER
278 
279 #ifdef LOAD
280 #pragma omp declare target
new_bar1()281 void new_bar1() {
282   TTT<int> *X = new TTT<int>();
283   X->emitted();
284 }
285 #pragma omp end declare target
286 #endif
287