1 // Test target codegen - host bc file has to be created first.
2 // RUN: %clang_cc1 -fopenmp -x c++ -triple x86_64-unknown-linux -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-host.bc
3 // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple nvptx64-unknown-unknown -aux-triple x86_64-unknown-linux -fopenmp-targets=nvptx64-nvidia-cuda %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-host.bc -fsyntax-only
4 // RUN: %clang_cc1 -fopenmp -x c++ -triple powerpc64le-unknown-linux-gnu -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-host.bc
5 // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple nvptx64-unknown-unknown -aux-triple powerpc64le-unknown-linux-gnu -fopenmp-targets=nvptx64-nvidia-cuda %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-host.bc -fsyntax-only
6
7 struct T {
8 char a;
9 #ifndef _ARCH_PPC
10 // expected-note@+1 {{'f' defined here}}
11 __float128 f;
12 #else
13 // expected-note@+1 {{'f' defined here}}
14 long double f;
15 #endif
16 char c;
TT17 T() : a(12), f(15) {}
18 #ifndef _ARCH_PPC
19 // expected-error@+5 {{'f' requires 128 bit size '__float128' type support, but device 'nvptx64-unknown-unknown' does not support it}}
20 #else
21 // expected-error@+3 {{'f' requires 128 bit size 'long double' type support, but device 'nvptx64-unknown-unknown' does not support it}}
22 #endif
operator +T23 T &operator+(T &b) {
24 f += b.a;
25 return *this;
26 }
27 };
28
29 struct T1 {
30 char a;
31 __int128 f;
32 __int128 f1;
33 char c;
T1T134 T1() : a(12), f(15) {}
operator /T135 T1 &operator/(T1 &b) {
36 f /= b.a;
37 return *this;
38 }
39 };
40
41 #ifndef _ARCH_PPC
42 // expected-error@+2 {{'boo' requires 128 bit size '__float128' type support, but device 'nvptx64-unknown-unknown' does not support it}}
43 // expected-note@+1 2{{'boo' defined here}}
boo(__float128 A)44 void boo(__float128 A) { return; }
45 #else
46 // expected-error@+2 {{'boo' requires 128 bit size 'long double' type support, but device 'nvptx64-unknown-unknown' does not support it}}
47 // expected-note@+1 2{{'boo' defined here}}
boo(long double A)48 void boo(long double A) { return; }
49 #endif
50 #pragma omp declare target
51 T a = T();
52 T f = a;
foo(T a=T ())53 void foo(T a = T()) {
54 a = a + f; // expected-note {{called by 'foo'}}
55 #ifndef _ARCH_PPC
56 // expected-error@+5 {{'boo' requires 128 bit size '__float128' type support, but device 'nvptx64-unknown-unknown' does not support it}}
57 #else
58 // expected-error@+3 {{'boo' requires 128 bit size 'long double' type support, but device 'nvptx64-unknown-unknown' does not support it}}
59 #endif
60 // expected-note@+1 {{called by 'foo'}}
61 boo(0);
62 return;
63 }
bar()64 T bar() {
65 return T();
66 }
67
baz()68 void baz() {
69 T t = bar();
70 }
71 T1 a1 = T1();
72 T1 f1 = a1;
foo1(T1 a=T1 ())73 void foo1(T1 a = T1()) {
74 a = a / f1;
75 return;
76 }
bar1()77 T1 bar1() {
78 return T1();
79 }
baz1()80 void baz1() {
81 T1 t = bar1();
82 }
83
dead_inline_declare_target()84 inline void dead_inline_declare_target() {
85 long double *a, b = 0;
86 a = &b;
87 }
dead_static_declare_target()88 static void dead_static_declare_target() {
89 long double *a, b = 0;
90 a = &b;
91 }
92 template<bool>
dead_template_declare_target()93 void dead_template_declare_target() {
94 long double *a, b = 0;
95 a = &b;
96 }
97
98 // expected-note@+2 {{'ld_return1a' defined here}}
99 // expected-error@+1 {{'ld_return1a' requires 128 bit size 'long double' type support, but device 'nvptx64-unknown-unknown' does not support it}}
ld_return1a()100 long double ld_return1a() { return 0; }
101 // expected-note@+2 {{'ld_arg1a' defined here}}
102 // expected-error@+1 {{'ld_arg1a' requires 128 bit size 'long double' type support, but device 'nvptx64-unknown-unknown' does not support it}}
ld_arg1a(long double ld)103 void ld_arg1a(long double ld) {}
104
105 typedef long double ld_ty;
106 // expected-note@+2 {{'ld_return1b' defined here}}
107 // expected-error@+1 {{'ld_return1b' requires 128 bit size 'ld_ty' (aka 'long double') type support, but device 'nvptx64-unknown-unknown' does not support it}}
ld_return1b()108 ld_ty ld_return1b() { return 0; }
109 // expected-note@+2 {{'ld_arg1b' defined here}}
110 // expected-error@+1 {{'ld_arg1b' requires 128 bit size 'ld_ty' (aka 'long double') type support, but device 'nvptx64-unknown-unknown' does not support it}}
ld_arg1b(ld_ty ld)111 void ld_arg1b(ld_ty ld) {}
112
ld_return1c()113 static long double ld_return1c() { return 0; }
ld_arg1c(long double ld)114 static void ld_arg1c(long double ld) {}
115
ld_return1d()116 inline long double ld_return1d() { return 0; }
ld_arg1d(long double ld)117 inline void ld_arg1d(long double ld) {}
118
119 // expected-note@+1 {{'ld_return1e' defined here}}
ld_return1e()120 static long double ld_return1e() { return 0; }
121 // expected-note@+1 {{'ld_arg1e' defined here}}
ld_arg1e(long double ld)122 static void ld_arg1e(long double ld) {}
123
124 // expected-note@+1 {{'ld_return1f' defined here}}
ld_return1f()125 inline long double ld_return1f() { return 0; }
126 // expected-note@+1 {{'ld_arg1f' defined here}}
ld_arg1f(long double ld)127 inline void ld_arg1f(long double ld) {}
128
ld_use1()129 inline void ld_use1() {
130 long double ld = 0;
131 ld += 1;
132 }
ld_use2()133 static void ld_use2() {
134 long double ld = 0;
135 ld += 1;
136 }
137
ld_use3()138 inline void ld_use3() {
139 // expected-note@+1 {{'ld' defined here}}
140 long double ld = 0;
141 // expected-error@+1 {{'ld' requires 128 bit size 'long double' type support, but device 'nvptx64-unknown-unknown' does not support it}}
142 ld += 1;
143 }
ld_use4()144 static void ld_use4() {
145 // expected-note@+1 {{'ld' defined here}}
146 long double ld = 0;
147 // expected-error@+1 {{'ld' requires 128 bit size 'long double' type support, but device 'nvptx64-unknown-unknown' does not support it}}
148 ld += 1;
149 }
150
external()151 void external() {
152 // expected-error@+1 {{'ld_return1e' requires 128 bit size 'long double' type support, but device 'nvptx64-unknown-unknown' does not support it}}
153 void *p1 = reinterpret_cast<void*>(&ld_return1e);
154 // expected-error@+1 {{'ld_arg1e' requires 128 bit size 'long double' type support, but device 'nvptx64-unknown-unknown' does not support it}}
155 void *p2 = reinterpret_cast<void*>(&ld_arg1e);
156 // expected-error@+1 {{'ld_return1f' requires 128 bit size 'long double' type support, but device 'nvptx64-unknown-unknown' does not support it}}
157 void *p3 = reinterpret_cast<void*>(&ld_return1f);
158 // expected-error@+1 {{'ld_arg1f' requires 128 bit size 'long double' type support, but device 'nvptx64-unknown-unknown' does not support it}}
159 void *p4 = reinterpret_cast<void*>(&ld_arg1f);
160 // TODO: The error message "called by" is not great.
161 // expected-note@+1 {{called by 'external'}}
162 void *p5 = reinterpret_cast<void*>(&ld_use3);
163 // expected-note@+1 {{called by 'external'}}
164 void *p6 = reinterpret_cast<void*>(&ld_use4);
165 }
166
167 #ifndef _ARCH_PPC
168 // expected-note@+2 {{'ld_return2a' defined here}}
169 // expected-error@+1 {{'ld_return2a' requires 128 bit size '__float128' type support, but device 'nvptx64-unknown-unknown' does not support it}}
ld_return2a()170 __float128 ld_return2a() { return 0; }
171 // expected-note@+2 {{'ld_arg2a' defined here}}
172 // expected-error@+1 {{'ld_arg2a' requires 128 bit size '__float128' type support, but device 'nvptx64-unknown-unknown' does not support it}}
ld_arg2a(__float128 ld)173 void ld_arg2a(__float128 ld) {}
174
175 typedef __float128 fp128_ty;
176 // expected-note@+2 {{'ld_return2b' defined here}}
177 // expected-error@+1 {{'ld_return2b' requires 128 bit size 'fp128_ty' (aka '__float128') type support, but device 'nvptx64-unknown-unknown' does not support it}}
ld_return2b()178 fp128_ty ld_return2b() { return 0; }
179 // expected-note@+2 {{'ld_arg2b' defined here}}
180 // expected-error@+1 {{'ld_arg2b' requires 128 bit size 'fp128_ty' (aka '__float128') type support, but device 'nvptx64-unknown-unknown' does not support it}}
ld_arg2b(fp128_ty ld)181 void ld_arg2b(fp128_ty ld) {}
182 #endif
183
184 #pragma omp end declare target
185
186 // TODO: There should not be an error here, dead_inline is never emitted.
187 // expected-note@+1 {{'f' defined here}}
dead_inline(long double f)188 inline long double dead_inline(long double f) {
189 #pragma omp target map(f)
190 // expected-error@+1 {{'f' requires 128 bit size 'long double' type support, but device 'nvptx64-unknown-unknown' does not support it}}
191 f = 1;
192 return f;
193 }
194
195 // TODO: There should not be an error here, dead_static is never emitted.
196 // expected-note@+1 {{'f' defined here}}
dead_static(long double f)197 static long double dead_static(long double f) {
198 #pragma omp target map(f)
199 // expected-error@+1 {{'f' requires 128 bit size 'long double' type support, but device 'nvptx64-unknown-unknown' does not support it}}
200 f = 1;
201 return f;
202 }
203
204 template<typename T>
dead_template(long double f)205 long double dead_template(long double f) {
206 #pragma omp target map(f)
207 f = 1;
208 return f;
209 }
210
211 #ifndef _ARCH_PPC
212 // expected-note@+1 {{'f' defined here}}
foo2(__float128 f)213 __float128 foo2(__float128 f) {
214 #pragma omp target map(f)
215 // expected-error@+1 {{'f' requires 128 bit size '__float128' type support, but device 'nvptx64-unknown-unknown' does not support it}}
216 f = 1;
217 return f;
218 }
219 #else
220 // expected-note@+1 {{'f' defined here}}
foo3(long double f)221 long double foo3(long double f) {
222 #pragma omp target map(f)
223 // expected-error@+1 {{'f' requires 128 bit size 'long double' type support, but device 'nvptx64-unknown-unknown' does not support it}}
224 f = 1;
225 return f;
226 }
227 #endif
228
foo3()229 T foo3() {
230 T S;
231 #pragma omp target map(S)
232 S.a = 1;
233 return S;
234 }
235
236 // Allow all sorts of stuff on host
237 #ifndef _ARCH_PPC
238 __float128 q, b;
239 __float128 c = q + b;
240 #else
241 long double q, b;
242 long double c = q + b;
243 #endif
244
hostFoo()245 void hostFoo() {
246 boo(c - b);
247 }
248
249 long double qa, qb;
250 decltype(qa + qb) qc;
251 double qd[sizeof(-(-(qc * 2)))];
252
253 struct A { };
254
255 template <bool>
256 struct A_type { typedef A type; };
257
258 template <class Sp, class Tp>
259 struct B {
260 enum { value = bool(Sp::value) || bool(Tp::value) };
261 typedef typename A_type<value>::type type;
262 };
263
264 void bar(_ExtInt(66) a) {
265 auto b = a;
266 }
267