Clang Project

clang_source_code/test/SemaCUDA/function-overload.cu
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.
10struct HostReturnTy {};
11struct HostReturnTy2 {};
12struct DeviceReturnTy {};
13struct DeviceReturnTy2 {};
14struct HostDeviceReturnTy {};
15struct TemplateReturnTy {};
16
17typedef HostReturnTy (*HostFnPtr)();
18typedef DeviceReturnTy (*DeviceFnPtr)();
19typedef HostDeviceReturnTy (*HostDeviceFnPtr)();
20typedef void (*GlobalFnPtr)();  // __global__ functions must return void.
21
22// CurrentReturnTy is {HostReturnTy,DeviceReturnTy} during {host,device}
23// compilation.
24#ifdef __CUDA_ARCH__
25typedef DeviceReturnTy CurrentReturnTy;
26#else
27typedef HostReturnTy CurrentReturnTy;
28#endif
29
30// CurrentFnPtr is a function pointer to a {host,device} function during
31// {host,device} compilation.
32typedef CurrentReturnTy (*CurrentFnPtr)();
33
34// Host and unattributed functions can't be overloaded.
35__host__ void hh() {} // expected-note {{previous definition is here}}
36void 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.
60extern "C" __host__ int chh() { return 0; } // expected-note {{previous definition is here}}
61extern "C" int chh() { return 0; }          // expected-error {{redefinition of 'chh'}}
62
63// H/D overloading is OK.
64extern "C" __device__ DeviceReturnTy cdh() { return DeviceReturnTy(); }
65extern "C" __host__ HostReturnTy cdh() { return HostReturnTy(); }
66
67// H/HD and D/HD overloading is not allowed.
68extern "C" __host__ __device__ int chhd1() { return 0; } // expected-note {{previous declaration is here}}
69extern "C" __host__ int chhd1() { return 0; }
70// expected-error@-1 {{__host__ function 'chhd1' cannot overload __host__ __device__ function 'chhd1'}}
71
72extern "C" __host__ int chhd2() { return 0; } // expected-note {{previous declaration is here}}
73extern "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
94extern "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
99extern "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.
215HostFnPtr fp_h = h;
216HostFnPtr fp_ch = ch;
217CurrentFnPtr fp_dh = dh;
218CurrentFnPtr fp_cdh = cdh;
219GlobalFnPtr fp_g = g;
220
221
222// Test overloading of destructors
223// Can't mix H and unattributed destructors
224struct d_h {
225  ~d_h() {} // expected-note {{previous definition is here}}
226  __host__ ~d_h() {} // expected-error {{destructor cannot be redeclared}}
227};
228
229// HD is OK
230struct d_hd {
231  __host__ __device__ ~d_hd() {}
232};
233
234// Test overloading of member functions
235struct 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
241struct m_dh {
242  __device__ void operator delete(void *ptr);
243  __host__ void operator delete(void *ptr);
244};
245
246// HD by itself is OK
247struct m_hd {
248  __device__ __host__ void operator delete(void *ptr);
249};
250
251struct 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
257struct 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
263struct 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
269struct 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.
277struct G {
278  friend void friend_of_g(G &arg); // expected-note {{previous declaration is here}}
279private:
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'}}
285void 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
298template <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.
333template <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.
396template <typename T> __host__ T template_overload(const T &a) { return a; };
397template <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
CodeThinker, Enjoy your reading code