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