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