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.
hh()35 __host__ void hh() {} // expected-note {{previous definition is here}}
hh()36 void hh() {} // expected-error {{redefinition of 'hh'}}
37
38 // H/D overloading is OK.
dh()39 __host__ HostReturnTy dh() { return HostReturnTy(); }
dh()40 __device__ DeviceReturnTy dh() { return DeviceReturnTy(); }
41
42 // H/HD and D/HD are not allowed.
hdh()43 __host__ __device__ int hdh() { return 0; } // expected-note {{previous definition is here}}
hdh()44 __host__ int hdh() { return 0; } // expected-error {{redefinition of 'hdh'}}
45
hhd()46 __host__ int hhd() { return 0; } // expected-note {{previous definition is here}}
hhd()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
hdd()51 __host__ __device__ int hdd() { return 0; } // expected-note {{previous definition is here}}
hdd()52 __device__ int hdd() { return 0; } // expected-error {{redefinition of 'hdd'}}
53
dhd()54 __device__ int dhd() { return 0; } // expected-note {{previous definition is here}}
dhd()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.
chh()60 extern "C" __host__ int chh() { return 0; } // expected-note {{previous definition is here}}
chh()61 extern "C" int chh() { return 0; } // expected-error {{redefinition of 'chh'}}
62
63 // H/D overloading is OK.
cdh()64 extern "C" __device__ DeviceReturnTy cdh() { return DeviceReturnTy(); }
cdh()65 extern "C" __host__ HostReturnTy cdh() { return HostReturnTy(); }
66
67 // H/HD and D/HD overloading is not allowed.
chhd1()68 extern "C" __host__ __device__ int chhd1() { return 0; } // expected-note {{previous definition is here}}
chhd1()69 extern "C" __host__ int chhd1() { return 0; } // expected-error {{redefinition of 'chhd1'}}
70
chhd2()71 extern "C" __host__ int chhd2() { return 0; } // expected-note {{previous definition is here}}
chhd2()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.
d()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
h()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
g()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
cd()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
ch()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
hostf()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
devicef()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
globalf()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
hostdevicef()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 GlobalFnPtr fp_g = g;
185 #if defined(__CUDA_ARCH__)
186 // expected-error@-2 {{reference to __global__ function 'g' in __host__ __device__ function}}
187 #endif
188 g();
189 g<<<0,0>>>();
190 #if !defined(__CUDA_ARCH__)
191 // expected-error@-3 {{call to global function g not configured}}
192 #else
193 // expected-error@-5 {{no matching function for call to 'g'}}
194 // expected-error@-5 {{reference to __global__ function 'g' in __host__ __device__ function}}
195 #endif // __CUDA_ARCH__
196 }
197
198 // Test for address of overloaded function resolution in the global context.
199 HostFnPtr fp_h = h;
200 HostFnPtr fp_ch = ch;
201 CurrentFnPtr fp_dh = dh;
202 CurrentFnPtr fp_cdh = cdh;
203 GlobalFnPtr fp_g = g;
204
205
206 // Test overloading of destructors
207 // Can't mix H and unattributed destructors
208 struct d_h {
~d_hd_h209 ~d_h() {} // expected-note {{previous declaration is here}}
~d_hd_h210 __host__ ~d_h() {} // expected-error {{destructor cannot be redeclared}}
211 };
212
213 // H/D overloading is OK
214 struct d_dh {
~d_dhd_dh215 __device__ ~d_dh() {}
~d_dhd_dh216 __host__ ~d_dh() {}
217 };
218
219 // HD is OK
220 struct d_hd {
~d_hdd_hd221 __host__ __device__ ~d_hd() {}
222 };
223
224 // Mixing H/D and HD is not allowed.
225 struct d_dhhd {
~d_dhhdd_dhhd226 __device__ ~d_dhhd() {}
~d_dhhdd_dhhd227 __host__ ~d_dhhd() {} // expected-note {{previous declaration is here}}
~d_dhhdd_dhhd228 __host__ __device__ ~d_dhhd() {} // expected-error {{destructor cannot be redeclared}}
229 };
230
231 struct d_hhd {
~d_hhdd_hhd232 __host__ ~d_hhd() {} // expected-note {{previous declaration is here}}
~d_hhdd_hhd233 __host__ __device__ ~d_hhd() {} // expected-error {{destructor cannot be redeclared}}
234 };
235
236 struct d_hdh {
~d_hdhd_hdh237 __host__ __device__ ~d_hdh() {} // expected-note {{previous declaration is here}}
~d_hdhd_hdh238 __host__ ~d_hdh() {} // expected-error {{destructor cannot be redeclared}}
239 };
240
241 struct d_dhd {
~d_dhdd_dhd242 __device__ ~d_dhd() {} // expected-note {{previous declaration is here}}
~d_dhdd_dhd243 __host__ __device__ ~d_dhd() {} // expected-error {{destructor cannot be redeclared}}
244 };
245
246 struct d_hdd {
~d_hddd_hdd247 __host__ __device__ ~d_hdd() {} // expected-note {{previous declaration is here}}
~d_hddd_hdd248 __device__ ~d_hdd() {} // expected-error {{destructor cannot be redeclared}}
249 };
250
251 // Test overloading of member functions
252 struct m_h {
253 void operator delete(void *ptr); // expected-note {{previous declaration is here}}
254 __host__ void operator delete(void *ptr); // expected-error {{class member cannot be redeclared}}
255 };
256
257 // D/H overloading is OK
258 struct m_dh {
259 __device__ void operator delete(void *ptr);
260 __host__ void operator delete(void *ptr);
261 };
262
263 // HD by itself is OK
264 struct m_hd {
265 __device__ __host__ void operator delete(void *ptr);
266 };
267
268 struct m_hhd {
operator deletem_hhd269 __host__ void operator delete(void *ptr) {} // expected-note {{previous declaration is here}}
operator deletem_hhd270 __host__ __device__ void operator delete(void *ptr) {} // expected-error {{class member cannot be redeclared}}
271 };
272
273 struct m_hdh {
operator deletem_hdh274 __host__ __device__ void operator delete(void *ptr) {} // expected-note {{previous declaration is here}}
operator deletem_hdh275 __host__ void operator delete(void *ptr) {} // expected-error {{class member cannot be redeclared}}
276 };
277
278 struct m_dhd {
operator deletem_dhd279 __device__ void operator delete(void *ptr) {} // expected-note {{previous declaration is here}}
operator deletem_dhd280 __host__ __device__ void operator delete(void *ptr) {} // expected-error {{class member cannot be redeclared}}
281 };
282
283 struct m_hdd {
operator deletem_hdd284 __host__ __device__ void operator delete(void *ptr) {} // expected-note {{previous declaration is here}}
operator deletem_hdd285 __device__ void operator delete(void *ptr) {} // expected-error {{class member cannot be redeclared}}
286 };
287
288 // __global__ functions can't be overloaded based on attribute
289 // difference.
290 struct G {
291 friend void friend_of_g(G &arg);
292 private:
293 int x;
294 };
friend_of_g(G & arg)295 __global__ void friend_of_g(G &arg) { int x = arg.x; } // expected-note {{previous definition is here}}
friend_of_g(G & arg)296 void friend_of_g(G &arg) { int x = arg.x; } // expected-error {{redefinition of 'friend_of_g'}}
297
298 // HD functions are sometimes allowed to call H or D functions -- this
299 // is an artifact of the source-to-source splitting performed by nvcc
300 // that we need to mimic. During device mode compilation in nvcc, host
301 // functions aren't present at all, so don't participate in
302 // overloading. But in clang, H and D functions are present in both
303 // compilation modes. Clang normally uses the target attribute as a
304 // tiebreaker between overloads with otherwise identical priority, but
305 // in order to match nvcc's behavior, we sometimes need to wholly
306 // discard overloads that would not be present during compilation
307 // under nvcc.
308
template_vs_function(T arg)309 template <typename T> TemplateReturnTy template_vs_function(T arg) {
310 return TemplateReturnTy();
311 }
template_vs_function(float arg)312 __device__ DeviceReturnTy template_vs_function(float arg) {
313 return DeviceReturnTy();
314 }
315
316 // Here we expect to call the templated function during host compilation, even
317 // if -fcuda-disable-target-call-checks is passed, and even though C++ overload
318 // rules prefer the non-templated function.
test_host_device_calls_template(void)319 __host__ __device__ void test_host_device_calls_template(void) {
320 #ifdef __CUDA_ARCH__
321 typedef DeviceReturnTy ExpectedReturnTy;
322 #else
323 typedef TemplateReturnTy ExpectedReturnTy;
324 #endif
325
326 ExpectedReturnTy ret1 = template_vs_function(1.0f);
327 ExpectedReturnTy ret2 = template_vs_function(2.0);
328 }
329
330 // Calls from __host__ and __device__ functions should always call the
331 // overloaded function that matches their mode.
test_host_calls_template_fn()332 __host__ void test_host_calls_template_fn() {
333 TemplateReturnTy ret1 = template_vs_function(1.0f);
334 TemplateReturnTy ret2 = template_vs_function(2.0);
335 }
336
test_device_calls_template_fn()337 __device__ void test_device_calls_template_fn() {
338 DeviceReturnTy ret1 = template_vs_function(1.0f);
339 DeviceReturnTy ret2 = template_vs_function(2.0);
340 }
341
342 // If we have a mix of HD and H-only or D-only candidates in the overload set,
343 // normal C++ overload resolution rules apply first.
template_vs_hd_function(T arg)344 template <typename T> TemplateReturnTy template_vs_hd_function(T arg) {
345 return TemplateReturnTy();
346 }
template_vs_hd_function(float arg)347 __host__ __device__ HostDeviceReturnTy template_vs_hd_function(float arg) {
348 return HostDeviceReturnTy();
349 }
350
test_host_device_calls_hd_template()351 __host__ __device__ void test_host_device_calls_hd_template() {
352 HostDeviceReturnTy ret1 = template_vs_hd_function(1.0f);
353 TemplateReturnTy ret2 = template_vs_hd_function(1);
354 }
355
test_host_calls_hd_template()356 __host__ void test_host_calls_hd_template() {
357 HostDeviceReturnTy ret1 = template_vs_hd_function(1.0f);
358 TemplateReturnTy ret2 = template_vs_hd_function(1);
359 }
360
test_device_calls_hd_template()361 __device__ void test_device_calls_hd_template() {
362 HostDeviceReturnTy ret1 = template_vs_hd_function(1.0f);
363 // Host-only function template is not callable with strict call checks,
364 // so for device side HD function will be the only choice.
365 HostDeviceReturnTy ret2 = template_vs_hd_function(1);
366 }
367
368 // Check that overloads still work the same way on both host and
369 // device side when the overload set contains only functions from one
370 // side of compilation.
device_only_function(int arg)371 __device__ DeviceReturnTy device_only_function(int arg) { return DeviceReturnTy(); }
device_only_function(float arg)372 __device__ DeviceReturnTy2 device_only_function(float arg) { return DeviceReturnTy2(); }
host_only_function(int arg)373 __host__ HostReturnTy host_only_function(int arg) { return HostReturnTy(); }
host_only_function(float arg)374 __host__ HostReturnTy2 host_only_function(float arg) { return HostReturnTy2(); }
375
test_host_device_single_side_overloading()376 __host__ __device__ void test_host_device_single_side_overloading() {
377 DeviceReturnTy ret1 = device_only_function(1);
378 DeviceReturnTy2 ret2 = device_only_function(1.0f);
379 HostReturnTy ret3 = host_only_function(1);
380 HostReturnTy2 ret4 = host_only_function(1.0f);
381 }
382