1 | // 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 |
2 | // 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 |
3 | // 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 -emit-pch -o %t |
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 -include-pch %t -o - | FileCheck %s |
5 | |
6 | // 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 |
7 | // 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 %s --check-prefix SIMD-ONLY |
8 | // 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 -emit-pch -o %t |
9 | // 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 -include-pch %t -verify -o - | FileCheck %s --check-prefix SIMD-ONLY |
10 | |
11 | // expected-no-diagnostics |
12 | |
13 | // SIMD-ONLY-NOT: {{__kmpc|__tgt}} |
14 | |
15 | // CHECK-NOT: define {{.*}}{{baz1|baz4|maini1|Base|virtual_}} |
16 | // CHECK-DAG: Bake |
17 | // CHECK-NOT: @{{hhh|ggg|fff|eee}} = |
18 | // CHECK-DAG: @aaa = external global i32, |
19 | // CHECK-DAG: @bbb = global i32 0, |
20 | // CHECK-DAG: weak constant %struct.__tgt_offload_entry { i8* bitcast (i32* @bbb to i8*), |
21 | // CHECK-DAG: @ccc = external global i32, |
22 | // CHECK-DAG: @ddd = global i32 0, |
23 | // CHECK-DAG: @hhh_decl_tgt_link_ptr = common global i32* null |
24 | // CHECK-DAG: @ggg_decl_tgt_link_ptr = common global i32* null |
25 | // CHECK-DAG: @fff_decl_tgt_link_ptr = common global i32* null |
26 | // CHECK-DAG: @eee_decl_tgt_link_ptr = common global i32* null |
27 | // CHECK-DAG: @{{.*}}maini1{{.*}}aaa = internal global i64 23, |
28 | // CHECK-DAG: @b = global i32 15, |
29 | // CHECK-DAG: @d = global i32 0, |
30 | // CHECK-DAG: @c = external global i32, |
31 | // CHECK-DAG: @globals = global %struct.S zeroinitializer, |
32 | // CHECK-DAG: [[STAT:@.+stat]] = internal global %struct.S zeroinitializer, |
33 | // CHECK-DAG: [[STAT_REF:@.+]] = internal constant %struct.S* [[STAT]] |
34 | // CHECK-DAG: @out_decl_target = global i32 0, |
35 | // CHECK-DAG: @llvm.used = appending global [6 x i8*] [i8* bitcast (void ()* @__omp_offloading__{{.+}}_globals_l[[@LINE+80]]_ctor to i8*), i8* bitcast (void ()* @__omp_offloading__{{.+}}_stat_l[[@LINE+81]]_ctor to i8*), |
36 | // CHECK-DAG: @llvm.compiler.used = appending global [1 x i8*] [i8* bitcast (%struct.S** [[STAT_REF]] to i8*)], |
37 | |
38 | // CHECK-DAG: define {{.*}}i32 @{{.*}}{{foo|bar|baz2|baz3|FA|f_method}}{{.*}}() |
39 | // CHECK-DAG: define {{.*}}void @{{.*}}TemplateClass{{.*}}(%class.TemplateClass* %{{.*}}) |
40 | // CHECK-DAG: define {{.*}}i32 @{{.*}}TemplateClass{{.*}}f_method{{.*}}(%class.TemplateClass* %{{.*}}) |
41 | // CHECK-DAG: define {{.*}}void @__omp_offloading__{{.*}}_globals_l[[@LINE+74]]_ctor() |
42 | |
43 | #ifndef HEADER |
44 | #define HEADER |
45 | |
46 | #pragma omp declare target |
47 | extern int bbb; |
48 | #pragma omp end declare target |
49 | #pragma omp declare target |
50 | extern int bbb; |
51 | #pragma omp end declare target |
52 | |
53 | #pragma omp declare target |
54 | extern int aaa; |
55 | int bbb = 0; |
56 | extern int ccc; |
57 | int ddd = 0; |
58 | #pragma omp end declare target |
59 | |
60 | #pragma omp declare target |
61 | extern int bbb; |
62 | #pragma omp end declare target |
63 | |
64 | extern int eee; |
65 | int fff = 0; |
66 | extern int ggg; |
67 | int hhh = 0; |
68 | #pragma omp declare target link(eee, fff, ggg, hhh) |
69 | |
70 | int out_decl_target = 0; |
71 | #pragma omp declare target |
72 | void lambda () { |
73 | #ifdef __cpp_lambdas |
74 | (void)[&] { (void)out_decl_target; }; |
75 | #else |
76 | #pragma clang __debug captured |
77 | { |
78 | (void)out_decl_target; |
79 | } |
80 | #endif |
81 | }; |
82 | #pragma omp end declare target |
83 | |
84 | template <typename T> |
85 | class TemplateClass { |
86 | T a; |
87 | public: |
88 | TemplateClass() {} |
89 | T f_method() const { return a; } |
90 | }; |
91 | |
92 | int foo(); |
93 | |
94 | static int baz1() { return 0; } |
95 | |
96 | int baz2(); |
97 | |
98 | int baz4() { return 5; } |
99 | |
100 | template <typename T> |
101 | T FA() { |
102 | TemplateClass<T> s; |
103 | return s.f_method(); |
104 | } |
105 | |
106 | #pragma omp declare target |
107 | struct S { |
108 | int a; |
109 | S(int a) : a(a) {} |
110 | }; |
111 | |
112 | int foo() { return 0; } |
113 | int b = 15; |
114 | int d; |
115 | S globals(d); |
116 | static S stat(d); |
117 | #pragma omp end declare target |
118 | int c; |
119 | |
120 | int bar() { return 1 + foo() + bar() + baz1() + baz2(); } |
121 | |
122 | int maini1() { |
123 | int a; |
124 | static long aa = 32 + bbb + ccc + fff + ggg; |
125 | // CHECK-DAG: define weak void @__omp_offloading_{{.*}}maini1{{.*}}_l[[@LINE+1]](i32* dereferenceable{{.*}}, i64 {{.*}}, i64 {{.*}}) |
126 | #pragma omp target map(tofrom \ |
127 | : a, b) |
128 | { |
129 | S s(a); |
130 | static long aaa = 23; |
131 | a = foo() + bar() + b + c + d + aa + aaa + FA<int>(); |
132 | } |
133 | return baz4(); |
134 | } |
135 | |
136 | int baz3() { return 2 + baz2(); } |
137 | int baz2() { |
138 | // CHECK-DAG: define weak void @__omp_offloading_{{.*}}baz2{{.*}}_l[[@LINE+1]](i64 {{.*}}) |
139 | #pragma omp target parallel |
140 | ++c; |
141 | return 2 + baz3(); |
142 | } |
143 | |
144 | extern int create() throw(); |
145 | |
146 | static __typeof(create) __t_create __attribute__((__weakref__("__create"))); |
147 | |
148 | int baz5() { |
149 | bool a; |
150 | // CHECK-DAG: define weak void @__omp_offloading_{{.*}}baz5{{.*}}_l[[@LINE+1]](i64 {{.*}}) |
151 | #pragma omp target |
152 | a = __extension__(void *) & __t_create != 0; |
153 | return a; |
154 | } |
155 | |
156 | template <typename T> |
157 | struct Base { |
158 | virtual ~Base() {} |
159 | }; |
160 | |
161 | template class Base<int>; |
162 | |
163 | template <typename T> |
164 | struct Bake { |
165 | virtual ~Bake() {} |
166 | }; |
167 | |
168 | #pragma omp declare target |
169 | template class Bake<int>; |
170 | #pragma omp end declare target |
171 | |
172 | struct BaseNonT { |
173 | virtual ~BaseNonT() {} |
174 | }; |
175 | |
176 | #pragma omp declare target |
177 | struct BakeNonT { |
178 | virtual ~BakeNonT() {} |
179 | }; |
180 | #pragma omp end declare target |
181 | |
182 | template <typename T> |
183 | struct B { |
184 | virtual void virtual_foo(); |
185 | }; |
186 | |
187 | void new_bar() { new B<int>(); } |
188 | |
189 | template <typename T> |
190 | void B<T>::virtual_foo() { |
191 | #pragma omp target |
192 | {} |
193 | } |
194 | |
195 | struct A { |
196 | virtual void emitted() {} |
197 | }; |
198 | |
199 | template <typename T> |
200 | struct C : public A { |
201 | virtual void emitted(); |
202 | }; |
203 | |
204 | template <typename T> |
205 | void C<T>::emitted() { |
206 | #pragma omp target |
207 | {} |
208 | } |
209 | |
210 | int main() { |
211 | A *X = new C<int>(); |
212 | X->emitted(); |
213 | return 0; |
214 | } |
215 | |
216 | // CHECK-DAG: define {{.*}}void @__omp_offloading_{{.*}}virtual_foo{{.*}}_l[[@LINE-25]]() |
217 | // CHECK-DAG: define {{.*}}void @__omp_offloading_{{.*}}emitted{{.*}}_l[[@LINE-11]]() |
218 | |
219 | // CHECK-DAG: declare extern_weak signext i32 @__create() |
220 | |
221 | // CHECK-NOT: define {{.*}}{{baz1|baz4|maini1|Base|virtual_}} |
222 | |
223 | // CHECK-DAG: !{i32 1, !"aaa", i32 0, i32 {{[0-9]+}}} |
224 | // CHECK-DAG: !{i32 1, !"ccc", i32 0, i32 {{[0-9]+}}} |
225 | // CHECK-DAG: !{{{.+}}virtual_foo |
226 | |
227 | #endif // HEADER |
228 | |