1 // REQUIRES: x86-registered-target
2 // REQUIRES: nvptx-registered-target
3 
4 // Make sure we handle target overloads correctly.
5 // RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu \
6 // RUN:    -fsyntax-only -fcuda-target-overloads -verify %s
7 // RUN: %clang_cc1 -triple nvptx64-nvidia-cuda \
8 // RUN:    -fsyntax-only -fcuda-target-overloads -fcuda-is-device -verify %s
9 
10 // Check target overloads handling with disabled call target checks.
11 // RUN: %clang_cc1 -DNOCHECKS -triple x86_64-unknown-linux-gnu -fsyntax-only \
12 // RUN:    -fcuda-disable-target-call-checks -fcuda-target-overloads -verify %s
13 // RUN: %clang_cc1 -DNOCHECKS -triple nvptx64-nvidia-cuda -fsyntax-only \
14 // RUN:    -fcuda-disable-target-call-checks -fcuda-target-overloads \
15 // RUN:    -fcuda-is-device -verify %s
16 
17 #include "Inputs/cuda.h"
18 
19 typedef int (*fp_t)(void);
20 typedef void (*gp_t)(void);
21 
22 // Host and unattributed functions can't be overloaded
23 __host__ int hh(void) { return 1; } // expected-note {{previous definition is here}}
24 int hh(void) { return 1; } // expected-error {{redefinition of 'hh'}}
25 
26 // H/D overloading is OK
27 __host__ int dh(void) { return 2; }
28 __device__ int dh(void) { return 2; }
29 
30 // H/HD and D/HD are not allowed
31 __host__ __device__ int hdh(void) { return 5; } // expected-note {{previous definition is here}}
32 __host__ int hdh(void) { return 4; } // expected-error {{redefinition of 'hdh'}}
33 
34 __host__ int hhd(void) { return 4; } // expected-note {{previous definition is here}}
35 __host__ __device__ int hhd(void) { return 5; } // expected-error {{redefinition of 'hhd'}}
36 // expected-warning@-1 {{attribute declaration must precede definition}}
37 // expected-note@-3 {{previous definition is here}}
38 
39 __host__ __device__ int hdd(void) { return 7; } // expected-note {{previous definition is here}}
40 __device__ int hdd(void) { return 6; } // expected-error {{redefinition of 'hdd'}}
41 
42 __device__ int dhd(void) { return 6; } // expected-note {{previous definition is here}}
43 __host__ __device__ int dhd(void) { return 7; } // expected-error {{redefinition of 'dhd'}}
44 // expected-warning@-1 {{attribute declaration must precede definition}}
45 // expected-note@-3 {{previous definition is here}}
46 
47 // Same tests for extern "C" functions
48 extern "C" __host__ int chh(void) {return 11;} // expected-note {{previous definition is here}}
49 extern "C" int chh(void) {return 11;} // expected-error {{redefinition of 'chh'}}
50 
51 // H/D overloading is OK
52 extern "C" __device__ int cdh(void) {return 10;}
53 extern "C" __host__ int cdh(void) {return 11;}
54 
55 // H/HD and D/HD overloading is not allowed.
56 extern "C" __host__ __device__ int chhd1(void) {return 12;} // expected-note {{previous definition is here}}
57 extern "C" __host__ int chhd1(void) {return 13;} // expected-error {{redefinition of 'chhd1'}}
58 
59 extern "C" __host__ int chhd2(void) {return 13;} // expected-note {{previous definition is here}}
60 extern "C" __host__ __device__ int chhd2(void) {return 12;} // expected-error {{redefinition of 'chhd2'}}
61 // expected-warning@-1 {{attribute declaration must precede definition}}
62 // expected-note@-3 {{previous definition is here}}
63 
64 // Helper functions to verify calling restrictions.
65 __device__ int d(void) { return 8; }
66 __host__ int h(void) { return 9; }
67 __global__ void g(void) {}
68 extern "C" __device__ int cd(void) {return 10;}
69 extern "C" __host__ int ch(void) {return 11;}
70 
71 __host__ void hostf(void) {
72   fp_t dp = d;
73   // expected-error@-1 {{reference to __device__ function 'd' in __host__ function}}
74   // expected-note@65 {{'d' declared here}}
75   fp_t cdp = cd;
76   // expected-error@-1 {{reference to __device__ function 'cd' in __host__ function}}
77   // expected-note@68 {{'cd' declared here}}
78   fp_t hp = h;
79   fp_t chp = ch;
80   fp_t dhp = dh;
81   fp_t cdhp = cdh;
82   gp_t gp = g;
83 
84   d();
85   // expected-error@-1 {{no matching function for call to 'd'}}
86   // expected-note@65 {{candidate function not viable: call to __device__ function from __host__ function}}
87   cd();
88   // expected-error@-1 {{no matching function for call to 'cd'}}
89   // expected-note@68 {{candidate function not viable: call to __device__ function from __host__ function}}
90   h();
91   ch();
92   dh();
93   cdh();
94   g(); // expected-error {{call to global function g not configured}}
95   g<<<0,0>>>();
96 }
97 
98 
99 __device__ void devicef(void) {
100   fp_t dp = d;
101   fp_t cdp = cd;
102   fp_t hp = h;
103   // expected-error@-1 {{reference to __host__ function 'h' in __device__ function}}
104   // expected-note@66 {{'h' declared here}}
105   fp_t chp = ch;
106   // expected-error@-1 {{reference to __host__ function 'ch' in __device__ function}}
107   // expected-note@69 {{'ch' declared here}}
108   fp_t dhp = dh;
109   fp_t cdhp = cdh;
110   gp_t gp = g; // expected-error {{reference to __global__ function 'g' in __device__ function}}
111                // expected-note@67 {{'g' declared here}}
112 
113   d();
114   cd();
115   h(); // expected-error {{no matching function for call to 'h'}}
116   // expected-note@66 {{candidate function not viable: call to __host__ function from __device__ function}}
117   ch(); // expected-error {{no matching function for call to 'ch'}}
118   // expected-note@69 {{candidate function not viable: call to __host__ function from __device__ function}}
119   dh();
120   cdh();
121   g(); // expected-error {{no matching function for call to 'g'}}
122   // expected-note@67 {{candidate function not viable: call to __global__ function from __device__ function}}
123   g<<<0,0>>>(); // expected-error {{reference to __global__ function 'g' in __device__ function}}
124   // expected-note@67 {{'g' declared here}}
125 }
126 
127 __global__ void globalf(void) {
128   fp_t dp = d;
129   fp_t cdp = cd;
130   fp_t hp = h;
131   // expected-error@-1 {{reference to __host__ function 'h' in __global__ function}}
132   // expected-note@66 {{'h' declared here}}
133   fp_t chp = ch;
134   // expected-error@-1 {{reference to __host__ function 'ch' in __global__ function}}
135   // expected-note@69 {{'ch' declared here}}
136   fp_t dhp = dh;
137   fp_t cdhp = cdh;
138   gp_t gp = g;
139   // expected-error@-1 {{reference to __global__ function 'g' in __global__ function}}
140   // expected-note@67 {{'g' declared here}}
141 
142   d();
143   cd();
144   h();
145   // expected-error@-1 {{no matching function for call to 'h'}}
146   // expected-note@66 {{candidate function not viable: call to __host__ function from __global__ function}}
147   ch();
148   // expected-error@-1 {{no matching function for call to 'ch'}}
149   // expected-note@69 {{candidate function not viable: call to __host__ function from __global__ function}}
150   dh();
151   cdh();
152   g(); // expected-error {{no matching function for call to 'g'}}
153   // expected-note@67 {{candidate function not viable: call to __global__ function from __global__ function}}
154   g<<<0,0>>>(); // expected-error {{reference to __global__ function 'g' in __global__ function}}
155   // expected-note@67 {{'g' declared here}}
156 }
157 
158 __host__ __device__ void hostdevicef(void) {
159   fp_t dp = d;
160   fp_t cdp = cd;
161   fp_t hp = h;
162   fp_t chp = ch;
163 #if !defined(NOCHECKS)
164 #if !defined(__CUDA_ARCH__)
165   // expected-error@-6 {{reference to __device__ function 'd' in __host__ __device__ function}}
166   // expected-note@65 {{'d' declared here}}
167   // expected-error@-7 {{reference to __device__ function 'cd' in __host__ __device__ function}}
168   // expected-note@68 {{'cd' declared here}}
169 #else
170   // expected-error@-9 {{reference to __host__ function 'h' in __host__ __device__ function}}
171   // expected-note@66 {{'h' declared here}}
172   // expected-error@-10 {{reference to __host__ function 'ch' in __host__ __device__ function}}
173   // expected-note@69 {{'ch' declared here}}
174 #endif
175 #endif
176   fp_t dhp = dh;
177   fp_t cdhp = cdh;
178   gp_t gp = g;
179 #if defined(__CUDA_ARCH__)
180   // expected-error@-2 {{reference to __global__ function 'g' in __host__ __device__ function}}
181   // expected-note@67 {{'g' declared here}}
182 #endif
183 
184   d();
185   cd();
186   h();
187   ch();
188 #if !defined(NOCHECKS)
189 #if !defined(__CUDA_ARCH__)
190   // expected-error@-6 {{no matching function for call to 'd'}}
191   // expected-note@65 {{candidate function not viable: call to __device__ function from __host__ __device__ function}}
192   // expected-error@-7 {{no matching function for call to 'cd'}}
193   // expected-note@68 {{candidate function not viable: call to __device__ function from __host__ __device__ function}}
194 #else
195   // expected-error@-9 {{no matching function for call to 'h'}}
196   // expected-note@66 {{candidate function not viable: call to __host__ function from __host__ __device__ function}}
197   // expected-error@-10 {{no matching function for call to 'ch'}}
198   // expected-note@69 {{candidate function not viable: call to __host__ function from __host__ __device__ function}}
199 #endif
200 #endif
201 
202   dh();
203   cdh();
204   g();
205   g<<<0,0>>>();
206 #if !defined(__CUDA_ARCH__)
207   // expected-error@-3 {{call to global function g not configured}}
208 #else
209   // expected-error@-5 {{no matching function for call to 'g'}}
210   // expected-note@67 {{candidate function not viable: call to __global__ function from __host__ __device__ function}}
211   // expected-error@-6 {{reference to __global__ function 'g' in __host__ __device__ function}}
212   // expected-note@67 {{'g' declared here}}
213 #endif  // __CUDA_ARCH__
214 }
215 
216 // Test for address of overloaded function resolution in the global context.
217 fp_t hp = h;
218 fp_t chp = ch;
219 fp_t dhp = dh;
220 fp_t cdhp = cdh;
221 gp_t gp = g;
222 
223 
224 // Test overloading of destructors
225 // Can't mix H and unattributed destructors
226 struct d_h {
227   ~d_h() {} // expected-note {{previous declaration is here}}
228   __host__ ~d_h() {} // expected-error {{destructor cannot be redeclared}}
229 };
230 
231 // H/D overloading is OK
232 struct d_dh {
233   __device__ ~d_dh() {}
234   __host__ ~d_dh() {}
235 };
236 
237 // HD is OK
238 struct d_hd {
239   __host__ __device__ ~d_hd() {}
240 };
241 
242 // Mixing H/D and HD is not allowed.
243 struct d_dhhd {
244   __device__ ~d_dhhd() {}
245   __host__ ~d_dhhd() {} // expected-note {{previous declaration is here}}
246   __host__ __device__ ~d_dhhd() {} // expected-error {{destructor cannot be redeclared}}
247 };
248 
249 struct d_hhd {
250   __host__ ~d_hhd() {} // expected-note {{previous declaration is here}}
251   __host__ __device__ ~d_hhd() {} // expected-error {{destructor cannot be redeclared}}
252 };
253 
254 struct d_hdh {
255   __host__ __device__ ~d_hdh() {} // expected-note {{previous declaration is here}}
256   __host__ ~d_hdh() {} // expected-error {{destructor cannot be redeclared}}
257 };
258 
259 struct d_dhd {
260   __device__ ~d_dhd() {} // expected-note {{previous declaration is here}}
261   __host__ __device__ ~d_dhd() {} // expected-error {{destructor cannot be redeclared}}
262 };
263 
264 struct d_hdd {
265   __host__ __device__ ~d_hdd() {} // expected-note {{previous declaration is here}}
266   __device__ ~d_hdd() {} // expected-error {{destructor cannot be redeclared}}
267 };
268 
269 // Test overloading of member functions
270 struct m_h {
271   void operator delete(void *ptr); // expected-note {{previous declaration is here}}
272   __host__ void operator delete(void *ptr); // expected-error {{class member cannot be redeclared}}
273 };
274 
275 // D/H overloading is OK
276 struct m_dh {
277   __device__ void operator delete(void *ptr);
278   __host__ void operator delete(void *ptr);
279 };
280 
281 // HD by itself is OK
282 struct m_hd {
283   __device__ __host__ void operator delete(void *ptr);
284 };
285 
286 struct m_hhd {
287   __host__ void operator delete(void *ptr) {} // expected-note {{previous declaration is here}}
288   __host__ __device__ void operator delete(void *ptr) {} // expected-error {{class member cannot be redeclared}}
289 };
290 
291 struct m_hdh {
292   __host__ __device__ void operator delete(void *ptr) {} // expected-note {{previous declaration is here}}
293   __host__ void operator delete(void *ptr) {} // expected-error {{class member cannot be redeclared}}
294 };
295 
296 struct m_dhd {
297   __device__ void operator delete(void *ptr) {} // expected-note {{previous declaration is here}}
298   __host__ __device__ void operator delete(void *ptr) {} // expected-error {{class member cannot be redeclared}}
299 };
300 
301 struct m_hdd {
302   __host__ __device__ void operator delete(void *ptr) {} // expected-note {{previous declaration is here}}
303   __device__ void operator delete(void *ptr) {} // expected-error {{class member cannot be redeclared}}
304 };
305 
306 // __global__ functions can't be overloaded based on attribute
307 // difference.
308 struct G {
309   friend void friend_of_g(G &arg);
310 private:
311   int x;
312 };
313 __global__ void friend_of_g(G &arg) { int x = arg.x; } // expected-note {{previous definition is here}}
314 void friend_of_g(G &arg) { int x = arg.x; } // expected-error {{redefinition of 'friend_of_g'}}
315