1 // REQUIRES: x86-registered-target
2 // REQUIRES: nvptx-registered-target
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
7 #include "Inputs/cuda.h"
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 {};
17 typedef HostReturnTy (*HostFnPtr)();
18 typedef DeviceReturnTy (*DeviceFnPtr)();
19 typedef HostDeviceReturnTy (*HostDeviceFnPtr)();
20 typedef void (*GlobalFnPtr)(); // __global__ functions must return void.
22 // CurrentReturnTy is {HostReturnTy,DeviceReturnTy} during {host,device}
25 typedef DeviceReturnTy CurrentReturnTy;
27 typedef HostReturnTy CurrentReturnTy;
30 // CurrentFnPtr is a function pointer to a {host,device} function during
31 // {host,device} compilation.
32 typedef CurrentReturnTy (*CurrentFnPtr)();
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'}}
38 // H/D overloading is OK.
39 __host__ HostReturnTy dh() { return HostReturnTy(); }
40 __device__ DeviceReturnTy dh() { return DeviceReturnTy(); }
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'}}
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'}}
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'}}
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'}}
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'}}
63 // H/D overloading is OK.
64 extern "C" __device__ DeviceReturnTy cdh() { return DeviceReturnTy(); }
65 extern "C" __host__ HostReturnTy cdh() { return HostReturnTy(); }
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'}}
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'}}
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}}
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}}
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}}
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}}
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}}
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'}}
112 HostReturnTy ret_h = h();
113 HostFnPtr fp_ch = ch;
114 HostReturnTy ret_ch = ch();
116 HostFnPtr fp_dh = dh;
117 HostReturnTy ret_dh = dh();
118 HostFnPtr fp_cdh = cdh;
119 HostReturnTy ret_cdh = cdh();
121 GlobalFnPtr fp_g = g;
122 g(); // expected-error {{call to global function 'g' not configured}}
126 __device__ void devicef() {
127 DeviceFnPtr fp_d = d;
128 DeviceReturnTy ret_d = d();
129 DeviceFnPtr fp_cd = cd;
130 DeviceReturnTy ret_cd = cd();
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'}}
137 DeviceFnPtr fp_dh = dh;
138 DeviceReturnTy ret_dh = dh();
139 DeviceFnPtr fp_cdh = cdh;
140 DeviceReturnTy ret_cdh = cdh();
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}}
147 __global__ void globalf() {
148 DeviceFnPtr fp_d = d;
149 DeviceReturnTy ret_d = d();
150 DeviceFnPtr fp_cd = cd;
151 DeviceReturnTy ret_cd = cd();
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'}}
158 DeviceFnPtr fp_dh = dh;
159 DeviceReturnTy ret_dh = dh();
160 DeviceFnPtr fp_cdh = cdh;
161 DeviceReturnTy ret_cdh = cdh();
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}}
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}}
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}}
191 CurrentFnPtr fp_dh = dh;
192 CurrentReturnTy ret_dh = dh();
193 CurrentFnPtr fp_cdh = cdh;
194 CurrentReturnTy ret_cdh = cdh();
196 GlobalFnPtr fp_g = g;
197 #if defined(__CUDA_ARCH__)
198 // expected-error@-2 {{reference to __global__ function 'g' in __host__ __device__ function}}
202 #if defined (__CUDA_ARCH__)
203 // expected-error@-2 {{reference to __global__ function 'g' in __host__ __device__ function}}
205 // expected-error@-4 {{call to global function 'g' not configured}}
209 #if defined(__CUDA_ARCH__)
210 // expected-error@-2 {{reference to __global__ function 'g' in __host__ __device__ function}}
214 // Test for address of overloaded function resolution in the global context.
216 HostFnPtr fp_ch = ch;
217 CurrentFnPtr fp_dh = dh;
218 CurrentFnPtr fp_cdh = cdh;
219 GlobalFnPtr fp_g = g;
222 // Test overloading of destructors
223 // Can't mix H and unattributed destructors
225 ~d_h() {} // expected-note {{previous definition is here}}
226 __host__ ~d_h() {} // expected-error {{destructor cannot be redeclared}}
231 __host__ __device__ ~d_hd() {}
234 // Test overloading of member functions
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}}
240 // D/H overloading is OK
242 __device__ void operator delete(void *ptr);
243 __host__ void operator delete(void *ptr);
246 // HD by itself is OK
248 __device__ __host__ void operator delete(void *ptr);
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'}}
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'}}
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'}}
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'}}
275 // __global__ functions can't be overloaded based on attribute
278 friend void friend_of_g(G &arg); // expected-note {{previous declaration is here}}
280 int x; // expected-note {{declared private here}}
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; }
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
298 template <typename T> TemplateReturnTy template_vs_function(T arg) {
299 return TemplateReturnTy();
301 __device__ DeviceReturnTy template_vs_function(float arg) {
302 return DeviceReturnTy();
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) {
310 typedef DeviceReturnTy ExpectedReturnTy;
312 typedef TemplateReturnTy ExpectedReturnTy;
315 ExpectedReturnTy ret1 = template_vs_function(1.0f);
316 ExpectedReturnTy ret2 = template_vs_function(2.0);
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);
326 __device__ void test_device_calls_template_fn() {
327 DeviceReturnTy ret1 = template_vs_function(1.0f);
328 DeviceReturnTy ret2 = template_vs_function(2.0);
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)
335 //expected-note@-2 {{declared here}}
338 return TemplateReturnTy();
340 __host__ __device__ HostDeviceReturnTy template_vs_hd_function(float arg) {
341 return HostDeviceReturnTy();
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);
348 // expected-error@-2 {{reference to __host__ function 'template_vs_hd_function<int>' in __host__ __device__ function}}
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);
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);
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}}
373 __host__ HostReturnTy host_only_function(int arg) { return HostReturnTy(); }
374 __host__ HostReturnTy2 host_only_function(float arg) { return HostReturnTy2(); }
376 // expected-note@-3 {{'host_only_function' declared here}}
377 // expected-note@-3 {{'host_only_function' declared here}}
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}}
387 HostReturnTy ret3 = host_only_function(1);
388 HostReturnTy2 ret4 = host_only_function(1.0f);
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}}
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; };
399 __host__ void test_host_template_overload() {
400 template_overload(1); // OK. Attribute-based overloading picks __host__ variant.
402 __device__ void test_device_template_overload() {
403 template_overload(1); // OK. Attribute-based overloading picks __device__ variant.
406 // Two classes with `operator-` defined. One of them is device only.
410 int operator-(const C1 &x, const C1 &y);
411 int operator-(const C2 &x, const C2 &y);
413 template <typename T>
414 __host__ __device__ int constexpr_overload(const T &x, const T &y) {
418 // Verify that function overloading doesn't prune candidate wrongly.
419 int test_constexpr_overload(C2 &x, C2 &y) {
420 return constexpr_overload(x, y);