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 // Opaque return types used to check that we pick the right overloads.
20 struct HostReturnTy {};
21 struct HostReturnTy2 {};
22 struct DeviceReturnTy {};
23 struct DeviceReturnTy2 {};
24 struct HostDeviceReturnTy {};
25 struct TemplateReturnTy {};
26 
27 typedef HostReturnTy (*HostFnPtr)();
28 typedef DeviceReturnTy (*DeviceFnPtr)();
29 typedef HostDeviceReturnTy (*HostDeviceFnPtr)();
30 typedef void (*GlobalFnPtr)();  // __global__ functions must return void.
31 
32 // CurrentReturnTy is {HostReturnTy,DeviceReturnTy} during {host,device}
33 // compilation.
34 #ifdef __CUDA_ARCH__
35 typedef DeviceReturnTy CurrentReturnTy;
36 #else
37 typedef HostReturnTy CurrentReturnTy;
38 #endif
39 
40 // CurrentFnPtr is a function pointer to a {host,device} function during
41 // {host,device} compilation.
42 typedef CurrentReturnTy (*CurrentFnPtr)();
43 
44 // Host and unattributed functions can't be overloaded.
45 __host__ void hh() {} // expected-note {{previous definition is here}}
46 void hh() {} // expected-error {{redefinition of 'hh'}}
47 
48 // H/D overloading is OK.
49 __host__ HostReturnTy dh() { return HostReturnTy(); }
50 __device__ DeviceReturnTy dh() { return DeviceReturnTy(); }
51 
52 // H/HD and D/HD are not allowed.
53 __host__ __device__ int hdh() { return 0; } // expected-note {{previous definition is here}}
54 __host__ int hdh() { return 0; }            // expected-error {{redefinition of 'hdh'}}
55 
56 __host__ int hhd() { return 0; }            // expected-note {{previous definition is here}}
57 __host__ __device__ int hhd() { return 0; } // expected-error {{redefinition of 'hhd'}}
58 // expected-warning@-1 {{attribute declaration must precede definition}}
59 // expected-note@-3 {{previous definition is here}}
60 
61 __host__ __device__ int hdd() { return 0; } // expected-note {{previous definition is here}}
62 __device__ int hdd() { return 0; }          // expected-error {{redefinition of 'hdd'}}
63 
64 __device__ int dhd() { return 0; }          // expected-note {{previous definition is here}}
65 __host__ __device__ int dhd() { return 0; } // expected-error {{redefinition of 'dhd'}}
66 // expected-warning@-1 {{attribute declaration must precede definition}}
67 // expected-note@-3 {{previous definition is here}}
68 
69 // Same tests for extern "C" functions.
70 extern "C" __host__ int chh() { return 0; } // expected-note {{previous definition is here}}
71 extern "C" int chh() { return 0; }          // expected-error {{redefinition of 'chh'}}
72 
73 // H/D overloading is OK.
74 extern "C" __device__ DeviceReturnTy cdh() { return DeviceReturnTy(); }
75 extern "C" __host__ HostReturnTy cdh() { return HostReturnTy(); }
76 
77 // H/HD and D/HD overloading is not allowed.
78 extern "C" __host__ __device__ int chhd1() { return 0; } // expected-note {{previous definition is here}}
79 extern "C" __host__ int chhd1() { return 0; }            // expected-error {{redefinition of 'chhd1'}}
80 
81 extern "C" __host__ int chhd2() { return 0; }            // expected-note {{previous definition is here}}
82 extern "C" __host__ __device__ int chhd2() { return 0; } // expected-error {{redefinition of 'chhd2'}}
83 // expected-warning@-1 {{attribute declaration must precede definition}}
84 // expected-note@-3 {{previous definition is here}}
85 
86 // Helper functions to verify calling restrictions.
87 __device__ DeviceReturnTy d() { return DeviceReturnTy(); }
88 // expected-note@-1 1+ {{'d' declared here}}
89 // expected-note@-2 1+ {{candidate function not viable: call to __device__ function from __host__ function}}
90 // expected-note@-3 0+ {{candidate function not viable: call to __device__ function from __host__ __device__ function}}
91 
92 __host__ HostReturnTy h() { return HostReturnTy(); }
93 // expected-note@-1 1+ {{'h' declared here}}
94 // expected-note@-2 1+ {{candidate function not viable: call to __host__ function from __device__ function}}
95 // expected-note@-3 0+ {{candidate function not viable: call to __host__ function from __host__ __device__ function}}
96 // expected-note@-4 1+ {{candidate function not viable: call to __host__ function from __global__ function}}
97 
98 __global__ void g() {}
99 // expected-note@-1 1+ {{'g' declared here}}
100 // expected-note@-2 1+ {{candidate function not viable: call to __global__ function from __device__ function}}
101 // expected-note@-3 0+ {{candidate function not viable: call to __global__ function from __host__ __device__ function}}
102 // expected-note@-4 1+ {{candidate function not viable: call to __global__ function from __global__ function}}
103 
104 extern "C" __device__ DeviceReturnTy cd() { return DeviceReturnTy(); }
105 // expected-note@-1 1+ {{'cd' declared here}}
106 // expected-note@-2 1+ {{candidate function not viable: call to __device__ function from __host__ function}}
107 // expected-note@-3 0+ {{candidate function not viable: call to __device__ function from __host__ __device__ function}}
108 
109 extern "C" __host__ HostReturnTy ch() { return HostReturnTy(); }
110 // expected-note@-1 1+ {{'ch' declared here}}
111 // expected-note@-2 1+ {{candidate function not viable: call to __host__ function from __device__ function}}
112 // expected-note@-3 0+ {{candidate function not viable: call to __host__ function from __host__ __device__ function}}
113 // expected-note@-4 1+ {{candidate function not viable: call to __host__ function from __global__ function}}
114 
115 __host__ void hostf() {
116   DeviceFnPtr fp_d = d;         // expected-error {{reference to __device__ function 'd' in __host__ function}}
117   DeviceReturnTy ret_d = d();   // expected-error {{no matching function for call to 'd'}}
118   DeviceFnPtr fp_cd = cd;       // expected-error {{reference to __device__ function 'cd' in __host__ function}}
119   DeviceReturnTy ret_cd = cd(); // expected-error {{no matching function for call to 'cd'}}
120 
121   HostFnPtr fp_h = h;
122   HostReturnTy ret_h = h();
123   HostFnPtr fp_ch = ch;
124   HostReturnTy ret_ch = ch();
125 
126   HostFnPtr fp_dh = dh;
127   HostReturnTy ret_dh = dh();
128   HostFnPtr fp_cdh = cdh;
129   HostReturnTy ret_cdh = cdh();
130 
131   GlobalFnPtr fp_g = g;
132   g(); // expected-error {{call to global function g not configured}}
133   g<<<0, 0>>>();
134 }
135 
136 __device__ void devicef() {
137   DeviceFnPtr fp_d = d;
138   DeviceReturnTy ret_d = d();
139   DeviceFnPtr fp_cd = cd;
140   DeviceReturnTy ret_cd = cd();
141 
142   HostFnPtr fp_h = h;         // expected-error {{reference to __host__ function 'h' in __device__ function}}
143   HostReturnTy ret_h = h();   // expected-error {{no matching function for call to 'h'}}
144   HostFnPtr fp_ch = ch;       // expected-error {{reference to __host__ function 'ch' in __device__ function}}
145   HostReturnTy ret_ch = ch(); // expected-error {{no matching function for call to 'ch'}}
146 
147   DeviceFnPtr fp_dh = dh;
148   DeviceReturnTy ret_dh = dh();
149   DeviceFnPtr fp_cdh = cdh;
150   DeviceReturnTy ret_cdh = cdh();
151 
152   GlobalFnPtr fp_g = g; // expected-error {{reference to __global__ function 'g' in __device__ function}}
153   g(); // expected-error {{no matching function for call to 'g'}}
154   g<<<0,0>>>(); // expected-error {{reference to __global__ function 'g' in __device__ function}}
155 }
156 
157 __global__ void globalf() {
158   DeviceFnPtr fp_d = d;
159   DeviceReturnTy ret_d = d();
160   DeviceFnPtr fp_cd = cd;
161   DeviceReturnTy ret_cd = cd();
162 
163   HostFnPtr fp_h = h;         // expected-error {{reference to __host__ function 'h' in __global__ function}}
164   HostReturnTy ret_h = h();   // expected-error {{no matching function for call to 'h'}}
165   HostFnPtr fp_ch = ch;       // expected-error {{reference to __host__ function 'ch' in __global__ function}}
166   HostReturnTy ret_ch = ch(); // expected-error {{no matching function for call to 'ch'}}
167 
168   DeviceFnPtr fp_dh = dh;
169   DeviceReturnTy ret_dh = dh();
170   DeviceFnPtr fp_cdh = cdh;
171   DeviceReturnTy ret_cdh = cdh();
172 
173   GlobalFnPtr fp_g = g; // expected-error {{reference to __global__ function 'g' in __global__ function}}
174   g(); // expected-error {{no matching function for call to 'g'}}
175   g<<<0,0>>>(); // expected-error {{reference to __global__ function 'g' in __global__ function}}
176 }
177 
178 __host__ __device__ void hostdevicef() {
179   DeviceFnPtr fp_d = d;
180   DeviceReturnTy ret_d = d();
181   DeviceFnPtr fp_cd = cd;
182   DeviceReturnTy ret_cd = cd();
183 #if !defined(NOCHECKS) && !defined(__CUDA_ARCH__)
184   // expected-error@-5 {{reference to __device__ function 'd' in __host__ __device__ function}}
185   // expected-error@-5 {{no matching function for call to 'd'}}
186   // expected-error@-5 {{reference to __device__ function 'cd' in __host__ __device__ function}}
187   // expected-error@-5 {{no matching function for call to 'cd'}}
188 #endif
189 
190   HostFnPtr fp_h = h;
191   HostReturnTy ret_h = h();
192   HostFnPtr fp_ch = ch;
193   HostReturnTy ret_ch = ch();
194 #if !defined(NOCHECKS) && defined(__CUDA_ARCH__)
195   // expected-error@-5 {{reference to __host__ function 'h' in __host__ __device__ function}}
196   // expected-error@-5 {{no matching function for call to 'h'}}
197   // expected-error@-5 {{reference to __host__ function 'ch' in __host__ __device__ function}}
198   // expected-error@-5 {{no matching function for call to 'ch'}}
199 #endif
200 
201   CurrentFnPtr fp_dh = dh;
202   CurrentReturnTy ret_dh = dh();
203   CurrentFnPtr fp_cdh = cdh;
204   CurrentReturnTy ret_cdh = cdh();
205 
206   GlobalFnPtr fp_g = g;
207 #if defined(__CUDA_ARCH__)
208   // expected-error@-2 {{reference to __global__ function 'g' in __host__ __device__ function}}
209 #endif
210   g();
211   g<<<0,0>>>();
212 #if !defined(__CUDA_ARCH__)
213   // expected-error@-3 {{call to global function g not configured}}
214 #else
215   // expected-error@-5 {{no matching function for call to 'g'}}
216   // expected-error@-5 {{reference to __global__ function 'g' in __host__ __device__ function}}
217 #endif  // __CUDA_ARCH__
218 }
219 
220 // Test for address of overloaded function resolution in the global context.
221 HostFnPtr fp_h = h;
222 HostFnPtr fp_ch = ch;
223 CurrentFnPtr fp_dh = dh;
224 CurrentFnPtr fp_cdh = cdh;
225 GlobalFnPtr fp_g = g;
226 
227 
228 // Test overloading of destructors
229 // Can't mix H and unattributed destructors
230 struct d_h {
231   ~d_h() {} // expected-note {{previous declaration is here}}
232   __host__ ~d_h() {} // expected-error {{destructor cannot be redeclared}}
233 };
234 
235 // H/D overloading is OK
236 struct d_dh {
237   __device__ ~d_dh() {}
238   __host__ ~d_dh() {}
239 };
240 
241 // HD is OK
242 struct d_hd {
243   __host__ __device__ ~d_hd() {}
244 };
245 
246 // Mixing H/D and HD is not allowed.
247 struct d_dhhd {
248   __device__ ~d_dhhd() {}
249   __host__ ~d_dhhd() {} // expected-note {{previous declaration is here}}
250   __host__ __device__ ~d_dhhd() {} // expected-error {{destructor cannot be redeclared}}
251 };
252 
253 struct d_hhd {
254   __host__ ~d_hhd() {} // expected-note {{previous declaration is here}}
255   __host__ __device__ ~d_hhd() {} // expected-error {{destructor cannot be redeclared}}
256 };
257 
258 struct d_hdh {
259   __host__ __device__ ~d_hdh() {} // expected-note {{previous declaration is here}}
260   __host__ ~d_hdh() {} // expected-error {{destructor cannot be redeclared}}
261 };
262 
263 struct d_dhd {
264   __device__ ~d_dhd() {} // expected-note {{previous declaration is here}}
265   __host__ __device__ ~d_dhd() {} // expected-error {{destructor cannot be redeclared}}
266 };
267 
268 struct d_hdd {
269   __host__ __device__ ~d_hdd() {} // expected-note {{previous declaration is here}}
270   __device__ ~d_hdd() {} // expected-error {{destructor cannot be redeclared}}
271 };
272 
273 // Test overloading of member functions
274 struct m_h {
275   void operator delete(void *ptr); // expected-note {{previous declaration is here}}
276   __host__ void operator delete(void *ptr); // expected-error {{class member cannot be redeclared}}
277 };
278 
279 // D/H overloading is OK
280 struct m_dh {
281   __device__ void operator delete(void *ptr);
282   __host__ void operator delete(void *ptr);
283 };
284 
285 // HD by itself is OK
286 struct m_hd {
287   __device__ __host__ void operator delete(void *ptr);
288 };
289 
290 struct m_hhd {
291   __host__ void operator delete(void *ptr) {} // expected-note {{previous declaration is here}}
292   __host__ __device__ void operator delete(void *ptr) {} // expected-error {{class member cannot be redeclared}}
293 };
294 
295 struct m_hdh {
296   __host__ __device__ void operator delete(void *ptr) {} // expected-note {{previous declaration is here}}
297   __host__ void operator delete(void *ptr) {} // expected-error {{class member cannot be redeclared}}
298 };
299 
300 struct m_dhd {
301   __device__ void operator delete(void *ptr) {} // expected-note {{previous declaration is here}}
302   __host__ __device__ void operator delete(void *ptr) {} // expected-error {{class member cannot be redeclared}}
303 };
304 
305 struct m_hdd {
306   __host__ __device__ void operator delete(void *ptr) {} // expected-note {{previous declaration is here}}
307   __device__ void operator delete(void *ptr) {} // expected-error {{class member cannot be redeclared}}
308 };
309 
310 // __global__ functions can't be overloaded based on attribute
311 // difference.
312 struct G {
313   friend void friend_of_g(G &arg);
314 private:
315   int x;
316 };
317 __global__ void friend_of_g(G &arg) { int x = arg.x; } // expected-note {{previous definition is here}}
318 void friend_of_g(G &arg) { int x = arg.x; } // expected-error {{redefinition of 'friend_of_g'}}
319 
320 // HD functions are sometimes allowed to call H or D functions -- this
321 // is an artifact of the source-to-source splitting performed by nvcc
322 // that we need to mimic. During device mode compilation in nvcc, host
323 // functions aren't present at all, so don't participate in
324 // overloading. But in clang, H and D functions are present in both
325 // compilation modes. Clang normally uses the target attribute as a
326 // tiebreaker between overloads with otherwise identical priority, but
327 // in order to match nvcc's behavior, we sometimes need to wholly
328 // discard overloads that would not be present during compilation
329 // under nvcc.
330 
331 template <typename T> TemplateReturnTy template_vs_function(T arg) {
332   return TemplateReturnTy();
333 }
334 __device__ DeviceReturnTy template_vs_function(float arg) {
335   return DeviceReturnTy();
336 }
337 
338 // Here we expect to call the templated function during host compilation, even
339 // if -fcuda-disable-target-call-checks is passed, and even though C++ overload
340 // rules prefer the non-templated function.
341 __host__ __device__ void test_host_device_calls_template(void) {
342 #ifdef __CUDA_ARCH__
343   typedef DeviceReturnTy ExpectedReturnTy;
344 #else
345   typedef TemplateReturnTy ExpectedReturnTy;
346 #endif
347 
348   ExpectedReturnTy ret1 = template_vs_function(1.0f);
349   ExpectedReturnTy ret2 = template_vs_function(2.0);
350 }
351 
352 // Calls from __host__ and __device__ functions should always call the
353 // overloaded function that matches their mode.
354 __host__ void test_host_calls_template_fn() {
355   TemplateReturnTy ret1 = template_vs_function(1.0f);
356   TemplateReturnTy ret2 = template_vs_function(2.0);
357 }
358 
359 __device__ void test_device_calls_template_fn() {
360   DeviceReturnTy ret1 = template_vs_function(1.0f);
361   DeviceReturnTy ret2 = template_vs_function(2.0);
362 }
363 
364 // If we have a mix of HD and H-only or D-only candidates in the overload set,
365 // normal C++ overload resolution rules apply first.
366 template <typename T> TemplateReturnTy template_vs_hd_function(T arg) {
367   return TemplateReturnTy();
368 }
369 __host__ __device__ HostDeviceReturnTy template_vs_hd_function(float arg) {
370   return HostDeviceReturnTy();
371 }
372 
373 __host__ __device__ void test_host_device_calls_hd_template() {
374   HostDeviceReturnTy ret1 = template_vs_hd_function(1.0f);
375 
376 #if defined(__CUDA_ARCH__) && !defined(NOCHECKS)
377   typedef HostDeviceReturnTy ExpectedReturnTy;
378 #else
379   typedef TemplateReturnTy ExpectedReturnTy;
380 #endif
381   ExpectedReturnTy ret2 = template_vs_hd_function(1);
382 }
383 
384 __host__ void test_host_calls_hd_template() {
385   HostDeviceReturnTy ret1 = template_vs_hd_function(1.0f);
386   TemplateReturnTy ret2 = template_vs_hd_function(1);
387 }
388 
389 __device__ void test_device_calls_hd_template() {
390   HostDeviceReturnTy ret1 = template_vs_hd_function(1.0f);
391   // Host-only function template is not callable with strict call checks,
392   // so for device side HD function will be the only choice.
393   HostDeviceReturnTy ret2 = template_vs_hd_function(1);
394 }
395 
396 // Check that overloads still work the same way on both host and
397 // device side when the overload set contains only functions from one
398 // side of compilation.
399 __device__ DeviceReturnTy device_only_function(int arg) { return DeviceReturnTy(); }
400 __device__ DeviceReturnTy2 device_only_function(float arg) { return DeviceReturnTy2(); }
401 __host__ HostReturnTy host_only_function(int arg) { return HostReturnTy(); }
402 __host__ HostReturnTy2 host_only_function(float arg) { return HostReturnTy2(); }
403 
404 __host__ __device__ void test_host_device_nochecks_overloading() {
405 #ifdef NOCHECKS
406   DeviceReturnTy ret1 = device_only_function(1);
407   DeviceReturnTy2 ret2 = device_only_function(1.0f);
408   HostReturnTy ret3 = host_only_function(1);
409   HostReturnTy2 ret4 = host_only_function(1.0f);
410 #endif
411 }
412