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 definition is here}}
44 __host__ int hdh() { return 0; }            // expected-error {{redefinition of 'hdh'}}
45 
46 __host__ int hhd() { return 0; }            // expected-note {{previous definition is here}}
47 __host__ __device__ int hhd() { return 0; } // expected-error {{redefinition of 'hhd'}}
48 // expected-warning@-1 {{attribute declaration must precede definition}}
49 // expected-note@-3 {{previous definition is here}}
50 
51 __host__ __device__ int hdd() { return 0; } // expected-note {{previous definition is here}}
52 __device__ int hdd() { return 0; }          // expected-error {{redefinition of 'hdd'}}
53 
54 __device__ int dhd() { return 0; }          // expected-note {{previous definition is here}}
55 __host__ __device__ int dhd() { return 0; } // expected-error {{redefinition of 'dhd'}}
56 // expected-warning@-1 {{attribute declaration must precede definition}}
57 // expected-note@-3 {{previous definition is here}}
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 definition is here}}
69 extern "C" __host__ int chhd1() { return 0; }            // expected-error {{redefinition of 'chhd1'}}
70 
71 extern "C" __host__ int chhd2() { return 0; }            // expected-note {{previous definition is here}}
72 extern "C" __host__ __device__ int chhd2() { return 0; } // expected-error {{redefinition of 'chhd2'}}
73 // expected-warning@-1 {{attribute declaration must precede definition}}
74 // expected-note@-3 {{previous definition is here}}
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 
174   HostFnPtr fp_h = h;
175   HostReturnTy ret_h = h();
176   HostFnPtr fp_ch = ch;
177   HostReturnTy ret_ch = ch();
178 
179   CurrentFnPtr fp_dh = dh;
180   CurrentReturnTy ret_dh = dh();
181   CurrentFnPtr fp_cdh = cdh;
182   CurrentReturnTy ret_cdh = cdh();
183 
184   g(); // expected-error {{call to global function g not configured}}
185 }
186 
187 // Test for address of overloaded function resolution in the global context.
188 HostFnPtr fp_h = h;
189 HostFnPtr fp_ch = ch;
190 CurrentFnPtr fp_dh = dh;
191 CurrentFnPtr fp_cdh = cdh;
192 GlobalFnPtr fp_g = g;
193 
194 
195 // Test overloading of destructors
196 // Can't mix H and unattributed destructors
197 struct d_h {
198   ~d_h() {} // expected-note {{previous declaration is here}}
199   __host__ ~d_h() {} // expected-error {{destructor cannot be redeclared}}
200 };
201 
202 // HD is OK
203 struct d_hd {
204   __host__ __device__ ~d_hd() {}
205 };
206 
207 // Test overloading of member functions
208 struct m_h {
209   void operator delete(void *ptr); // expected-note {{previous declaration is here}}
210   __host__ void operator delete(void *ptr); // expected-error {{class member cannot be redeclared}}
211 };
212 
213 // D/H overloading is OK
214 struct m_dh {
215   __device__ void operator delete(void *ptr);
216   __host__ void operator delete(void *ptr);
217 };
218 
219 // HD by itself is OK
220 struct m_hd {
221   __device__ __host__ void operator delete(void *ptr);
222 };
223 
224 struct m_hhd {
225   __host__ void operator delete(void *ptr) {} // expected-note {{previous declaration is here}}
226   __host__ __device__ void operator delete(void *ptr) {} // expected-error {{class member cannot be redeclared}}
227 };
228 
229 struct m_hdh {
230   __host__ __device__ void operator delete(void *ptr) {} // expected-note {{previous declaration is here}}
231   __host__ void operator delete(void *ptr) {} // expected-error {{class member cannot be redeclared}}
232 };
233 
234 struct m_dhd {
235   __device__ void operator delete(void *ptr) {} // expected-note {{previous declaration is here}}
236   __host__ __device__ void operator delete(void *ptr) {} // expected-error {{class member cannot be redeclared}}
237 };
238 
239 struct m_hdd {
240   __host__ __device__ void operator delete(void *ptr) {} // expected-note {{previous declaration is here}}
241   __device__ void operator delete(void *ptr) {} // expected-error {{class member cannot be redeclared}}
242 };
243 
244 // __global__ functions can't be overloaded based on attribute
245 // difference.
246 struct G {
247   friend void friend_of_g(G &arg);
248 private:
249   int x;
250 };
251 __global__ void friend_of_g(G &arg) { int x = arg.x; } // expected-note {{previous definition is here}}
252 void friend_of_g(G &arg) { int x = arg.x; } // expected-error {{redefinition of 'friend_of_g'}}
253 
254 // HD functions are sometimes allowed to call H or D functions -- this
255 // is an artifact of the source-to-source splitting performed by nvcc
256 // that we need to mimic. During device mode compilation in nvcc, host
257 // functions aren't present at all, so don't participate in
258 // overloading. But in clang, H and D functions are present in both
259 // compilation modes. Clang normally uses the target attribute as a
260 // tiebreaker between overloads with otherwise identical priority, but
261 // in order to match nvcc's behavior, we sometimes need to wholly
262 // discard overloads that would not be present during compilation
263 // under nvcc.
264 
265 template <typename T> TemplateReturnTy template_vs_function(T arg) {
266   return TemplateReturnTy();
267 }
268 __device__ DeviceReturnTy template_vs_function(float arg) {
269   return DeviceReturnTy();
270 }
271 
272 // Here we expect to call the templated function during host compilation, even
273 // if -fcuda-disable-target-call-checks is passed, and even though C++ overload
274 // rules prefer the non-templated function.
275 __host__ __device__ void test_host_device_calls_template(void) {
276 #ifdef __CUDA_ARCH__
277   typedef DeviceReturnTy ExpectedReturnTy;
278 #else
279   typedef TemplateReturnTy ExpectedReturnTy;
280 #endif
281 
282   ExpectedReturnTy ret1 = template_vs_function(1.0f);
283   ExpectedReturnTy ret2 = template_vs_function(2.0);
284 }
285 
286 // Calls from __host__ and __device__ functions should always call the
287 // overloaded function that matches their mode.
288 __host__ void test_host_calls_template_fn() {
289   TemplateReturnTy ret1 = template_vs_function(1.0f);
290   TemplateReturnTy ret2 = template_vs_function(2.0);
291 }
292 
293 __device__ void test_device_calls_template_fn() {
294   DeviceReturnTy ret1 = template_vs_function(1.0f);
295   DeviceReturnTy ret2 = template_vs_function(2.0);
296 }
297 
298 // If we have a mix of HD and H-only or D-only candidates in the overload set,
299 // normal C++ overload resolution rules apply first.
300 template <typename T> TemplateReturnTy template_vs_hd_function(T arg) {
301   return TemplateReturnTy();
302 }
303 __host__ __device__ HostDeviceReturnTy template_vs_hd_function(float arg) {
304   return HostDeviceReturnTy();
305 }
306 
307 __host__ __device__ void test_host_device_calls_hd_template() {
308   HostDeviceReturnTy ret1 = template_vs_hd_function(1.0f);
309   TemplateReturnTy ret2 = template_vs_hd_function(1);
310 }
311 
312 __host__ void test_host_calls_hd_template() {
313   HostDeviceReturnTy ret1 = template_vs_hd_function(1.0f);
314   TemplateReturnTy ret2 = template_vs_hd_function(1);
315 }
316 
317 __device__ void test_device_calls_hd_template() {
318   HostDeviceReturnTy ret1 = template_vs_hd_function(1.0f);
319   // Host-only function template is not callable with strict call checks,
320   // so for device side HD function will be the only choice.
321   HostDeviceReturnTy ret2 = template_vs_hd_function(1);
322 }
323 
324 // Check that overloads still work the same way on both host and
325 // device side when the overload set contains only functions from one
326 // side of compilation.
327 __device__ DeviceReturnTy device_only_function(int arg) { return DeviceReturnTy(); }
328 __device__ DeviceReturnTy2 device_only_function(float arg) { return DeviceReturnTy2(); }
329 __host__ HostReturnTy host_only_function(int arg) { return HostReturnTy(); }
330 __host__ HostReturnTy2 host_only_function(float arg) { return HostReturnTy2(); }
331 
332 __host__ __device__ void test_host_device_single_side_overloading() {
333   DeviceReturnTy ret1 = device_only_function(1);
334   DeviceReturnTy2 ret2 = device_only_function(1.0f);
335   HostReturnTy ret3 = host_only_function(1);
336   HostReturnTy2 ret4 = host_only_function(1.0f);
337 }
338 
339 // Verify that we allow overloading function templates.
340 template <typename T> __host__ T template_overload(const T &a) { return a; };
341 template <typename T> __device__ T template_overload(const T &a) { return a; };
342 
343 __host__ void test_host_template_overload() {
344   template_overload(1); // OK. Attribute-based overloading picks __host__ variant.
345 }
346 __device__ void test_device_template_overload() {
347   template_overload(1); // OK. Attribute-based overloading picks __device__ variant.
348 }
349