1 // REQUIRES: x86-registered-target
2 // REQUIRES: nvptx-registered-target
3
4 // RUN: %clang_cc1 -std=c++14 -triple x86_64-unknown-linux-gnu -fsyntax-only \
5 // RUN: -verify=host,hostdefer,devdefer,expected %s
6 // RUN: %clang_cc1 -std=c++14 -triple nvptx64-nvidia-cuda -fsyntax-only \
7 // RUN: -fcuda-is-device -verify=dev,devnodeferonly,hostdefer,devdefer,expected %s
8 // RUN: %clang_cc1 -fgpu-exclude-wrong-side-overloads -fgpu-defer-diag -DDEFER=1 \
9 // RUN: -std=c++14 -triple x86_64-unknown-linux-gnu -fsyntax-only \
10 // RUN: -verify=host,hostdefer,expected %s
11 // RUN: %clang_cc1 -fgpu-exclude-wrong-side-overloads -fgpu-defer-diag -DDEFER=1 \
12 // RUN: -std=c++14 -triple nvptx64-nvidia-cuda -fsyntax-only -fcuda-is-device \
13 // RUN: -verify=dev,devdeferonly,devdefer,expected %s
14
15 #include "Inputs/cuda.h"
16
17 // Opaque return types used to check that we pick the right overloads.
18 struct HostReturnTy {};
19 struct HostReturnTy2 {};
20 struct DeviceReturnTy {};
21 struct DeviceReturnTy2 {};
22 struct HostDeviceReturnTy {};
23 struct TemplateReturnTy {};
24
25 typedef HostReturnTy (*HostFnPtr)();
26 typedef DeviceReturnTy (*DeviceFnPtr)();
27 typedef HostDeviceReturnTy (*HostDeviceFnPtr)();
28 typedef void (*GlobalFnPtr)(); // __global__ functions must return void.
29
30 // CurrentReturnTy is {HostReturnTy,DeviceReturnTy} during {host,device}
31 // compilation.
32 #ifdef __CUDA_ARCH__
33 typedef DeviceReturnTy CurrentReturnTy;
34 #else
35 typedef HostReturnTy CurrentReturnTy;
36 #endif
37
38 // CurrentFnPtr is a function pointer to a {host,device} function during
39 // {host,device} compilation.
40 typedef CurrentReturnTy (*CurrentFnPtr)();
41
42 // Host and unattributed functions can't be overloaded.
hh()43 __host__ void hh() {} // expected-note {{previous definition is here}}
hh()44 void hh() {} // expected-error {{redefinition of 'hh'}}
45
46 // H/D overloading is OK.
dh()47 __host__ HostReturnTy dh() { return HostReturnTy(); }
dh()48 __device__ DeviceReturnTy dh() { return DeviceReturnTy(); }
49
50 // H/HD and D/HD are not allowed.
hdh()51 __host__ __device__ int hdh() { return 0; } // expected-note {{previous declaration is here}}
hdh()52 __host__ int hdh() { return 0; }
53 // expected-error@-1 {{__host__ function 'hdh' cannot overload __host__ __device__ function 'hdh'}}
54
hhd()55 __host__ int hhd() { return 0; } // expected-note {{previous declaration is here}}
hhd()56 __host__ __device__ int hhd() { return 0; }
57 // expected-error@-1 {{__host__ __device__ function 'hhd' cannot overload __host__ function 'hhd'}}
58
hdd()59 __host__ __device__ int hdd() { return 0; } // expected-note {{previous declaration is here}}
hdd()60 __device__ int hdd() { return 0; }
61 // expected-error@-1 {{__device__ function 'hdd' cannot overload __host__ __device__ function 'hdd'}}
62
dhd()63 __device__ int dhd() { return 0; } // expected-note {{previous declaration is here}}
dhd()64 __host__ __device__ int dhd() { return 0; }
65 // expected-error@-1 {{__host__ __device__ function 'dhd' cannot overload __device__ function 'dhd'}}
66
67 // Same tests for extern "C" functions.
chh()68 extern "C" __host__ int chh() { return 0; } // expected-note {{previous definition is here}}
chh()69 extern "C" int chh() { return 0; } // expected-error {{redefinition of 'chh'}}
70
71 // H/D overloading is OK.
cdh()72 extern "C" __device__ DeviceReturnTy cdh() { return DeviceReturnTy(); }
cdh()73 extern "C" __host__ HostReturnTy cdh() { return HostReturnTy(); }
74
75 // H/HD and D/HD overloading is not allowed.
chhd1()76 extern "C" __host__ __device__ int chhd1() { return 0; } // expected-note {{previous declaration is here}}
chhd1()77 extern "C" __host__ int chhd1() { return 0; }
78 // expected-error@-1 {{__host__ function 'chhd1' cannot overload __host__ __device__ function 'chhd1'}}
79
chhd2()80 extern "C" __host__ int chhd2() { return 0; } // expected-note {{previous declaration is here}}
chhd2()81 extern "C" __host__ __device__ int chhd2() { return 0; }
82 // expected-error@-1 {{__host__ __device__ function 'chhd2' cannot overload __host__ function 'chhd2'}}
83
84 // Helper functions to verify calling restrictions.
d()85 __device__ DeviceReturnTy d() { return DeviceReturnTy(); }
86 // host-note@-1 1+ {{'d' declared here}}
87 // hostdefer-note@-2 1+ {{candidate function not viable: call to __device__ function from __host__ function}}
88 // expected-note@-3 0+ {{candidate function not viable: call to __device__ function from __host__ __device__ function}}
89
h()90 __host__ HostReturnTy h() { return HostReturnTy(); }
91 // dev-note@-1 1+ {{'h' declared here}}
92 // devdefer-note@-2 1+ {{candidate function not viable: call to __host__ function from __device__ function}}
93 // expected-note@-3 0+ {{candidate function not viable: call to __host__ function from __host__ __device__ function}}
94 // devdefer-note@-4 1+ {{candidate function not viable: call to __host__ function from __global__ function}}
95
g()96 __global__ void g() {}
97 // dev-note@-1 1+ {{'g' declared here}}
98 // devdefer-note@-2 1+ {{candidate function not viable: call to __global__ function from __device__ function}}
99 // expected-note@-3 0+ {{candidate function not viable: call to __global__ function from __host__ __device__ function}}
100 // devdefer-note@-4 1+ {{candidate function not viable: call to __global__ function from __global__ function}}
101
cd()102 extern "C" __device__ DeviceReturnTy cd() { return DeviceReturnTy(); }
103 // host-note@-1 1+ {{'cd' declared here}}
104 // hostdefer-note@-2 1+ {{candidate function not viable: call to __device__ function from __host__ function}}
105 // expected-note@-3 0+ {{candidate function not viable: call to __device__ function from __host__ __device__ function}}
106
ch()107 extern "C" __host__ HostReturnTy ch() { return HostReturnTy(); }
108 // dev-note@-1 1+ {{'ch' declared here}}
109 // devdefer-note@-2 1+ {{candidate function not viable: call to __host__ function from __device__ function}}
110 // expected-note@-3 0+ {{candidate function not viable: call to __host__ function from __host__ __device__ function}}
111 // devdefer-note@-4 1+ {{candidate function not viable: call to __host__ function from __global__ function}}
112
hostf()113 __host__ void hostf() {
114 DeviceFnPtr fp_d = d; // host-error {{reference to __device__ function 'd' in __host__ function}}
115 DeviceReturnTy ret_d = d(); // hostdefer-error {{no matching function for call to 'd'}}
116 DeviceFnPtr fp_cd = cd; // host-error {{reference to __device__ function 'cd' in __host__ function}}
117 DeviceReturnTy ret_cd = cd(); // hostdefer-error {{no matching function for call to 'cd'}}
118
119 HostFnPtr fp_h = h;
120 HostReturnTy ret_h = h();
121 HostFnPtr fp_ch = ch;
122 HostReturnTy ret_ch = ch();
123
124 HostFnPtr fp_dh = dh;
125 HostReturnTy ret_dh = dh();
126 HostFnPtr fp_cdh = cdh;
127 HostReturnTy ret_cdh = cdh();
128
129 GlobalFnPtr fp_g = g;
130 g(); // expected-error {{call to global function 'g' not configured}}
131 g<<<0, 0>>>();
132 }
133
devicef()134 __device__ void devicef() {
135 DeviceFnPtr fp_d = d;
136 DeviceReturnTy ret_d = d();
137 DeviceFnPtr fp_cd = cd;
138 DeviceReturnTy ret_cd = cd();
139
140 HostFnPtr fp_h = h; // dev-error {{reference to __host__ function 'h' in __device__ function}}
141 HostReturnTy ret_h = h(); // devdefer-error {{no matching function for call to 'h'}}
142 HostFnPtr fp_ch = ch; // dev-error {{reference to __host__ function 'ch' in __device__ function}}
143 HostReturnTy ret_ch = ch(); // devdefer-error {{no matching function for call to 'ch'}}
144
145 DeviceFnPtr fp_dh = dh;
146 DeviceReturnTy ret_dh = dh();
147 DeviceFnPtr fp_cdh = cdh;
148 DeviceReturnTy ret_cdh = cdh();
149
150 GlobalFnPtr fp_g = g; // dev-error {{reference to __global__ function 'g' in __device__ function}}
151 g(); // devdefer-error {{no matching function for call to 'g'}}
152 g<<<0,0>>>(); // dev-error {{reference to __global__ function 'g' in __device__ function}}
153 }
154
globalf()155 __global__ void globalf() {
156 DeviceFnPtr fp_d = d;
157 DeviceReturnTy ret_d = d();
158 DeviceFnPtr fp_cd = cd;
159 DeviceReturnTy ret_cd = cd();
160
161 HostFnPtr fp_h = h; // dev-error {{reference to __host__ function 'h' in __global__ function}}
162 HostReturnTy ret_h = h(); // devdefer-error {{no matching function for call to 'h'}}
163 HostFnPtr fp_ch = ch; // dev-error {{reference to __host__ function 'ch' in __global__ function}}
164 HostReturnTy ret_ch = ch(); // devdefer-error {{no matching function for call to 'ch'}}
165
166 DeviceFnPtr fp_dh = dh;
167 DeviceReturnTy ret_dh = dh();
168 DeviceFnPtr fp_cdh = cdh;
169 DeviceReturnTy ret_cdh = cdh();
170
171 GlobalFnPtr fp_g = g; // dev-error {{reference to __global__ function 'g' in __global__ function}}
172 g(); // devdefer-error {{no matching function for call to 'g'}}
173 g<<<0,0>>>(); // dev-error {{reference to __global__ function 'g' in __global__ function}}
174 }
175
hostdevicef()176 __host__ __device__ void hostdevicef() {
177 DeviceFnPtr fp_d = d;
178 DeviceReturnTy ret_d = d();
179 DeviceFnPtr fp_cd = cd;
180 DeviceReturnTy ret_cd = cd();
181 #if !defined(__CUDA_ARCH__)
182 // expected-error@-5 {{reference to __device__ function 'd' in __host__ __device__ function}}
183 // expected-error@-5 {{reference to __device__ function 'd' in __host__ __device__ function}}
184 // expected-error@-5 {{reference to __device__ function 'cd' in __host__ __device__ function}}
185 // expected-error@-5 {{reference to __device__ function 'cd' in __host__ __device__ function}}
186 #endif
187
188 HostFnPtr fp_h = h;
189 HostReturnTy ret_h = h();
190 HostFnPtr fp_ch = ch;
191 HostReturnTy ret_ch = ch();
192 #if defined(__CUDA_ARCH__)
193 // expected-error@-5 {{reference to __host__ function 'h' in __host__ __device__ function}}
194 // expected-error@-5 {{reference to __host__ function 'h' in __host__ __device__ function}}
195 // devdefer-error@-5 {{reference to __host__ function 'ch' in __host__ __device__ function}}
196 // expected-error@-5 {{reference to __host__ function 'ch' in __host__ __device__ function}}
197 #endif
198
199 CurrentFnPtr fp_dh = dh;
200 CurrentReturnTy ret_dh = dh();
201 CurrentFnPtr fp_cdh = cdh;
202 CurrentReturnTy ret_cdh = cdh();
203
204 GlobalFnPtr fp_g = g;
205 #if defined(__CUDA_ARCH__)
206 // expected-error@-2 {{reference to __global__ function 'g' in __host__ __device__ function}}
207 #endif
208
209 g();
210 #if defined (__CUDA_ARCH__)
211 // expected-error@-2 {{reference to __global__ function 'g' in __host__ __device__ function}}
212 #else
213 // expected-error@-4 {{call to global function 'g' not configured}}
214 #endif
215
216 g<<<0,0>>>();
217 #if defined(__CUDA_ARCH__)
218 // expected-error@-2 {{reference to __global__ function 'g' in __host__ __device__ function}}
219 #endif
220 }
221
222 // Test for address of overloaded function resolution in the global context.
223 HostFnPtr fp_h = h;
224 HostFnPtr fp_ch = ch;
225 #if defined (__CUDA_ARCH__)
226 __device__
227 #endif
228 CurrentFnPtr fp_dh = dh;
229 #if defined (__CUDA_ARCH__)
230 __device__
231 #endif
232 CurrentFnPtr fp_cdh = cdh;
233 GlobalFnPtr fp_g = g;
234
235
236 // Test overloading of destructors
237 // Can't mix H and unattributed destructors
238 struct d_h {
~d_hd_h239 ~d_h() {} // expected-note {{previous definition is here}}
~d_hd_h240 __host__ ~d_h() {} // expected-error {{destructor cannot be redeclared}}
241 };
242
243 // HD is OK
244 struct d_hd {
~d_hdd_hd245 __host__ __device__ ~d_hd() {}
246 };
247
248 // Test overloading of member functions
249 struct m_h {
250 void operator delete(void *ptr); // expected-note {{previous declaration is here}}
251 __host__ void operator delete(void *ptr); // expected-error {{class member cannot be redeclared}}
252 };
253
254 // D/H overloading is OK
255 struct m_dh {
256 __device__ void operator delete(void *ptr);
257 __host__ void operator delete(void *ptr);
258 };
259
260 // HD by itself is OK
261 struct m_hd {
262 __device__ __host__ void operator delete(void *ptr);
263 };
264
265 struct m_hhd {
operator deletem_hhd266 __host__ void operator delete(void *ptr) {} // expected-note {{previous declaration is here}}
operator deletem_hhd267 __host__ __device__ void operator delete(void *ptr) {}
268 // expected-error@-1 {{__host__ __device__ function 'operator delete' cannot overload __host__ function 'operator delete'}}
269 };
270
271 struct m_hdh {
operator deletem_hdh272 __host__ __device__ void operator delete(void *ptr) {} // expected-note {{previous declaration is here}}
operator deletem_hdh273 __host__ void operator delete(void *ptr) {}
274 // expected-error@-1 {{__host__ function 'operator delete' cannot overload __host__ __device__ function 'operator delete'}}
275 };
276
277 struct m_dhd {
operator deletem_dhd278 __device__ void operator delete(void *ptr) {} // expected-note {{previous declaration is here}}
operator deletem_dhd279 __host__ __device__ void operator delete(void *ptr) {}
280 // expected-error@-1 {{__host__ __device__ function 'operator delete' cannot overload __device__ function 'operator delete'}}
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) {}
286 // expected-error@-1 {{__device__ function 'operator delete' cannot overload __host__ __device__ function 'operator delete'}}
287 };
288
289 // __global__ functions can't be overloaded based on attribute
290 // difference.
291 struct G {
292 friend void friend_of_g(G &arg); // expected-note {{previous declaration is here}}
293 private:
294 int x; // expected-note {{declared private here}}
295 };
friend_of_g(G & arg)296 __global__ void friend_of_g(G &arg) { int x = arg.x; }
297 // expected-error@-1 {{__global__ function 'friend_of_g' cannot overload __host__ function 'friend_of_g'}}
298 // expected-error@-2 {{'x' is a private member of 'G'}}
friend_of_g(G & arg)299 void friend_of_g(G &arg) { int x = arg.x; }
300
301 // HD functions are sometimes allowed to call H or D functions -- this
302 // is an artifact of the source-to-source splitting performed by nvcc
303 // that we need to mimic. During device mode compilation in nvcc, host
304 // functions aren't present at all, so don't participate in
305 // overloading. But in clang, H and D functions are present in both
306 // compilation modes. Clang normally uses the target attribute as a
307 // tiebreaker between overloads with otherwise identical priority, but
308 // in order to match nvcc's behavior, we sometimes need to wholly
309 // discard overloads that would not be present during compilation
310 // under nvcc.
311
template_vs_function(T arg)312 template <typename T> TemplateReturnTy template_vs_function(T arg) {
313 return TemplateReturnTy();
314 }
template_vs_function(float arg)315 __device__ DeviceReturnTy template_vs_function(float arg) {
316 return DeviceReturnTy();
317 }
318
319 // Here we expect to call the templated function during host compilation, even
320 // if -fcuda-disable-target-call-checks is passed, and even though C++ overload
321 // rules prefer the non-templated function.
test_host_device_calls_template(void)322 __host__ __device__ void test_host_device_calls_template(void) {
323 #ifdef __CUDA_ARCH__
324 typedef DeviceReturnTy ExpectedReturnTy;
325 #else
326 typedef TemplateReturnTy ExpectedReturnTy;
327 #endif
328
329 ExpectedReturnTy ret1 = template_vs_function(1.0f);
330 ExpectedReturnTy ret2 = template_vs_function(2.0);
331 }
332
333 // Calls from __host__ and __device__ functions should always call the
334 // overloaded function that matches their mode.
test_host_calls_template_fn()335 __host__ void test_host_calls_template_fn() {
336 TemplateReturnTy ret1 = template_vs_function(1.0f);
337 TemplateReturnTy ret2 = template_vs_function(2.0);
338 }
339
test_device_calls_template_fn()340 __device__ void test_device_calls_template_fn() {
341 DeviceReturnTy ret1 = template_vs_function(1.0f);
342 DeviceReturnTy ret2 = template_vs_function(2.0);
343 }
344
345 // If we have a mix of HD and H-only or D-only candidates in the overload set,
346 // normal C++ overload resolution rules apply first.
template_vs_hd_function(T arg)347 template <typename T> TemplateReturnTy template_vs_hd_function(T arg)
348 // devnodeferonly-note@-1{{'template_vs_hd_function<int>' declared here}}
349 {
350 return TemplateReturnTy();
351 }
template_vs_hd_function(float arg)352 __host__ __device__ HostDeviceReturnTy template_vs_hd_function(float arg) {
353 return HostDeviceReturnTy();
354 }
355
test_host_device_calls_hd_template()356 __host__ __device__ void test_host_device_calls_hd_template() {
357 #if __CUDA_ARCH__ && DEFER
358 typedef HostDeviceReturnTy ExpectedReturnTy;
359 #else
360 typedef TemplateReturnTy ExpectedReturnTy;
361 #endif
362 HostDeviceReturnTy ret1 = template_vs_hd_function(1.0f);
363 ExpectedReturnTy ret2 = template_vs_hd_function(1);
364 // devnodeferonly-error@-1{{reference to __host__ function 'template_vs_hd_function<int>' in __host__ __device__ function}}
365 }
366
test_host_calls_hd_template()367 __host__ void test_host_calls_hd_template() {
368 HostDeviceReturnTy ret1 = template_vs_hd_function(1.0f);
369 TemplateReturnTy ret2 = template_vs_hd_function(1);
370 }
371
test_device_calls_hd_template()372 __device__ void test_device_calls_hd_template() {
373 HostDeviceReturnTy ret1 = template_vs_hd_function(1.0f);
374 // Host-only function template is not callable with strict call checks,
375 // so for device side HD function will be the only choice.
376 HostDeviceReturnTy ret2 = template_vs_hd_function(1);
377 }
378
379 // Check that overloads still work the same way on both host and
380 // device side when the overload set contains only functions from one
381 // side of compilation.
device_only_function(int arg)382 __device__ DeviceReturnTy device_only_function(int arg) { return DeviceReturnTy(); }
device_only_function(float arg)383 __device__ DeviceReturnTy2 device_only_function(float arg) { return DeviceReturnTy2(); }
384 #ifndef __CUDA_ARCH__
385 // expected-note@-3 2{{'device_only_function' declared here}}
386 // expected-note@-3 2{{'device_only_function' declared here}}
387 #endif
host_only_function(int arg)388 __host__ HostReturnTy host_only_function(int arg) { return HostReturnTy(); }
host_only_function(float arg)389 __host__ HostReturnTy2 host_only_function(float arg) { return HostReturnTy2(); }
390 #ifdef __CUDA_ARCH__
391 // expected-note@-3 2{{'host_only_function' declared here}}
392 // expected-note@-3 2{{'host_only_function' declared here}}
393 #endif
394
test_host_device_single_side_overloading()395 __host__ __device__ void test_host_device_single_side_overloading() {
396 DeviceReturnTy ret1 = device_only_function(1);
397 DeviceReturnTy2 ret2 = device_only_function(1.0f);
398 #ifndef __CUDA_ARCH__
399 // expected-error@-3 {{reference to __device__ function 'device_only_function' in __host__ __device__ function}}
400 // expected-error@-3 {{reference to __device__ function 'device_only_function' in __host__ __device__ function}}
401 #endif
402 HostReturnTy ret3 = host_only_function(1);
403 HostReturnTy2 ret4 = host_only_function(1.0f);
404 #ifdef __CUDA_ARCH__
405 // expected-error@-3 {{reference to __host__ function 'host_only_function' in __host__ __device__ function}}
406 // expected-error@-3 {{reference to __host__ function 'host_only_function' in __host__ __device__ function}}
407 #endif
408 }
409
410 // wrong-sided overloading should not cause diagnostic unless it is emitted.
411 // This inline function is not emitted.
test_host_device_wrong_side_overloading_inline_no_diag()412 inline __host__ __device__ void test_host_device_wrong_side_overloading_inline_no_diag() {
413 DeviceReturnTy ret1 = device_only_function(1);
414 DeviceReturnTy2 ret2 = device_only_function(1.0f);
415 HostReturnTy ret3 = host_only_function(1);
416 HostReturnTy2 ret4 = host_only_function(1.0f);
417 }
418
419 // wrong-sided overloading should cause diagnostic if it is emitted.
420 // This inline function is emitted since it is called by an emitted function.
test_host_device_wrong_side_overloading_inline_diag()421 inline __host__ __device__ void test_host_device_wrong_side_overloading_inline_diag() {
422 DeviceReturnTy ret1 = device_only_function(1);
423 DeviceReturnTy2 ret2 = device_only_function(1.0f);
424 #ifndef __CUDA_ARCH__
425 // expected-error@-3 {{reference to __device__ function 'device_only_function' in __host__ __device__ function}}
426 // expected-error@-3 {{reference to __device__ function 'device_only_function' in __host__ __device__ function}}
427 #endif
428 HostReturnTy ret3 = host_only_function(1);
429 HostReturnTy2 ret4 = host_only_function(1.0f);
430 #ifdef __CUDA_ARCH__
431 // expected-error@-3 {{reference to __host__ function 'host_only_function' in __host__ __device__ function}}
432 // expected-error@-3 {{reference to __host__ function 'host_only_function' in __host__ __device__ function}}
433 #endif
434 }
435
test_host_device_wrong_side_overloading_inline_diag_caller()436 __host__ __device__ void test_host_device_wrong_side_overloading_inline_diag_caller() {
437 test_host_device_wrong_side_overloading_inline_diag();
438 // expected-note@-1 {{called by 'test_host_device_wrong_side_overloading_inline_diag_caller'}}
439 }
440
441 // Verify that we allow overloading function templates.
template_overload(const T & a)442 template <typename T> __host__ T template_overload(const T &a) { return a; };
template_overload(const T & a)443 template <typename T> __device__ T template_overload(const T &a) { return a; };
444
test_host_template_overload()445 __host__ void test_host_template_overload() {
446 template_overload(1); // OK. Attribute-based overloading picks __host__ variant.
447 }
test_device_template_overload()448 __device__ void test_device_template_overload() {
449 template_overload(1); // OK. Attribute-based overloading picks __device__ variant.
450 }
451
452 // Two classes with `operator-` defined. One of them is device only.
453 struct C1;
454 struct C2;
455 __device__
456 int operator-(const C1 &x, const C1 &y);
457 int operator-(const C2 &x, const C2 &y);
458
459 template <typename T>
constexpr_overload(const T & x,const T & y)460 __host__ __device__ int constexpr_overload(const T &x, const T &y) {
461 return x - y;
462 }
463
464 // Verify that function overloading doesn't prune candidate wrongly.
test_constexpr_overload(C2 & x,C2 & y)465 int test_constexpr_overload(C2 &x, C2 &y) {
466 return constexpr_overload(x, y);
467 }
468
469 // Verify no ambiguity for new operator.
470 void *a = new int;
471 __device__ void *b = new int;
472 // expected-error@-1{{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables}}
473
474 // Verify no ambiguity for new operator.
475 template<typename _Tp> _Tp&& f();
476 template<typename _Tp, typename = decltype(new _Tp(f<_Tp>()))>
477 void __test();
478
foo()479 void foo() {
480 __test<int>();
481 }
482
483 // Test resolving implicit host device candidate vs wrong-sided candidate.
484 // In device compilation, implicit host device caller choose implicit host
485 // device candidate and wrong-sided candidate with equal preference.
486 // Resolution result should not change with/without pragma.
487 namespace ImplicitHostDeviceVsWrongSided {
488 HostReturnTy callee(double x);
489 #pragma clang force_cuda_host_device begin
490 HostDeviceReturnTy callee(int x);
implicit_hd_caller()491 inline HostReturnTy implicit_hd_caller() {
492 return callee(1.0);
493 }
494 #pragma clang force_cuda_host_device end
495 }
496
497 // Test resolving implicit host device candidate vs same-sided candidate.
498 // In host compilation, implicit host device caller choose implicit host
499 // device candidate and same-sided candidate with equal preference.
500 // Resolution result should not change with/without pragma.
501 namespace ImplicitHostDeviceVsSameSide {
502 HostReturnTy callee(int x);
503 #pragma clang force_cuda_host_device begin
504 HostDeviceReturnTy callee(double x);
implicit_hd_caller()505 inline HostDeviceReturnTy implicit_hd_caller() {
506 return callee(1.0);
507 }
508 #pragma clang force_cuda_host_device end
509 }
510
511 // Test resolving explicit host device candidate vs. wrong-sided candidate.
512 // When -fgpu-defer-diag is off, wrong-sided candidate is not excluded, therefore
513 // the first callee is chosen.
514 // When -fgpu-defer-diag is on, wrong-sided candidate is excluded, therefore
515 // the second callee is chosen.
516 namespace ExplicitHostDeviceVsWrongSided {
517 HostReturnTy callee(double x);
518 __host__ __device__ HostDeviceReturnTy callee(int x);
519 #if __CUDA_ARCH__ && DEFER
520 typedef HostDeviceReturnTy ExpectedRetTy;
521 #else
522 typedef HostReturnTy ExpectedRetTy;
523 #endif
explicit_hd_caller()524 inline __host__ __device__ ExpectedRetTy explicit_hd_caller() {
525 return callee(1.0);
526 }
527 }
528
529 // In the implicit host device function 'caller', the second 'callee' should be
530 // chosen since it has better match, even though it is an implicit host device
531 // function whereas the first 'callee' is a host function. A diagnostic will be
532 // emitted if the first 'callee' is chosen since deduced return type cannot be
533 // used before it is defined.
534 namespace ImplicitHostDeviceByConstExpr {
535 template <class a> a b;
536 auto callee(...);
537 template <class d> constexpr auto callee(d) -> decltype(0);
538 struct e {
gImplicitHostDeviceByConstExpr::e539 template <class ad, class... f> static auto g(ad, f...) {
540 return h<e, decltype(b<f>)...>;
541 }
542 struct i {
callerImplicitHostDeviceByConstExpr::e::i543 template <class, class... f> static constexpr auto caller(f... k) {
544 return callee(k...);
545 }
546 };
hImplicitHostDeviceByConstExpr::e547 template <class, class... f> static auto h() {
548 return i::caller<int, f...>;
549 }
550 };
551 class l {
l()552 l() {
553 e::g([] {}, this);
554 }
555 };
556 }
557
558 // Implicit HD candidate competes with device candidate.
559 // a and b have implicit HD copy ctor. In copy ctor of b, ctor of a is resolved.
560 // copy ctor of a should win over a(short), otherwise there will be ambiguity
561 // due to conversion operator.
562 namespace TestImplicitHDWithD {
563 struct a {
564 __device__ a(short);
565 __device__ operator unsigned() const;
566 __device__ operator int() const;
567 };
568 struct b {
569 a d;
570 };
f(b g)571 void f(b g) { b e = g; }
572 }
573
574 // Implicit HD candidate competes with host candidate.
575 // a and b have implicit HD copy ctor. In copy ctor of b, ctor of a is resolved.
576 // copy ctor of a should win over a(short), otherwise there will be ambiguity
577 // due to conversion operator.
578 namespace TestImplicitHDWithH {
579 struct a {
580 a(short);
581 __device__ operator unsigned() const;
582 __device__ operator int() const;
583 };
584 struct b {
585 a d;
586 };
f(b g)587 void f(b g) { b e = g; }
588 }
589
590 // Implicit HD candidate competes with HD candidate.
591 // a and b have implicit HD copy ctor. In copy ctor of b, ctor of a is resolved.
592 // copy ctor of a should win over a(short), otherwise there will be ambiguity
593 // due to conversion operator.
594 namespace TestImplicitHDWithHD {
595 struct a {
596 __host__ __device__ a(short);
597 __device__ operator unsigned() const;
598 __device__ operator int() const;
599 };
600 struct b {
601 a d;
602 };
f(b g)603 void f(b g) { b e = g; }
604 }
605
606 // HD candidate competes with H candidate.
607 // HD has type mismatch whereas H has type match.
608 // In device compilation, H wins when -fgpu-defer-diag is off and HD wins
609 // when -fgpu-defer-diags is on. In both cases the diagnostic should be
610 // deferred.
611 namespace TestDeferNoMatchingFuncNotEmitted {
612 template <typename> struct a {};
613 namespace b {
614 struct c : a<int> {};
615 template <typename d> void ag(d);
616 } // namespace b
617 template <typename ae>
ag(a<ae>)618 __host__ __device__ void ag(a<ae>) {
619 ae e;
620 ag(e);
621 }
f()622 void f() { (void)ag<b::c>; }
623 }
624
625 namespace TestDeferNoMatchingFuncEmitted {
626 template <typename> struct a {};
627 namespace b {
628 struct c : a<int> {};
629 template <typename d> void ag(d);
630 // devnodeferonly-note@-1{{'ag<TestDeferNoMatchingFuncEmitted::b::c>' declared here}}
631 } // namespace b
632 template <typename ae>
ag(a<ae>)633 __host__ __device__ void ag(a<ae>) {
634 ae e;
635 ag(e);
636 // devnodeferonly-error@-1{{reference to __host__ function 'ag<TestDeferNoMatchingFuncEmitted::b::c>' in __host__ __device__ function}}
637 // devdeferonly-error@-2{{no matching function for call to 'ag'}}
638 // devdeferonly-note@-3{{called by 'ag<TestDeferNoMatchingFuncEmitted::b::c>'}}
639 }
f()640 __host__ __device__ void f() { (void)ag<b::c>; }
641 // devnodeferonly-note@-1{{called by 'f'}}
642 // devdeferonly-note@-2{{called by 'f'}}
643 }
644
645 // Two HD candidates compete with H candidate.
646 // HDs have type mismatch whereas H has type match.
647 // In device compilation, H wins when -fgpu-defer-diag is off and two HD win
648 // when -fgpu-defer-diags is on. In both cases the diagnostic should be
649 // deferred.
650 namespace TestDeferAmbiguityNotEmitted {
651 template <typename> struct a {};
652 namespace b {
653 struct c : a<int> {};
654 template <typename d> void ag(d, int);
655 } // namespace b
656 template <typename ae>
ag(a<ae>,float)657 __host__ __device__ void ag(a<ae>, float) {
658 ae e;
659 ag(e, 1);
660 }
661 template <typename ae>
ag(a<ae>,double)662 __host__ __device__ void ag(a<ae>, double) {
663 }
f()664 void f() {
665 b::c x;
666 ag(x, 1);
667 }
668 }
669
670 namespace TestDeferAmbiguityEmitted {
671 template <typename> struct a {};
672 namespace b {
673 struct c : a<int> {};
674 template <typename d> void ag(d, int);
675 // devnodeferonly-note@-1{{'ag<TestDeferAmbiguityEmitted::b::c>' declared here}}
676 } // namespace b
677 template <typename ae>
ag(a<ae>,float)678 __host__ __device__ void ag(a<ae>, float) {
679 // devdeferonly-note@-1{{candidate function [with ae = int]}}
680 ae e;
681 ag(e, 1);
682 }
683 template <typename ae>
ag(a<ae>,double)684 __host__ __device__ void ag(a<ae>, double) {
685 // devdeferonly-note@-1{{candidate function [with ae = int]}}
686 }
f()687 __host__ __device__ void f() {
688 b::c x;
689 ag(x, 1);
690 // devnodeferonly-error@-1{{reference to __host__ function 'ag<TestDeferAmbiguityEmitted::b::c>' in __host__ __device__ function}}
691 // devdeferonly-error@-2{{call to 'ag' is ambiguous}}
692 }
693 }
694
695 // Implicit HD functions compute with H function and D function.
696 // In host compilation, foo(0.0, 2) should resolve to X::foo<double, int>.
697 // In device compilation, foo(0.0, 2) should resolve to foo(double, int).
698 // In either case there should be no ambiguity.
699 namespace TestImplicitHDWithHAndD {
700 namespace X {
foo(double,double)701 inline double foo(double, double) { return 0;}
foo(float,float)702 inline constexpr float foo(float, float) { return 1;}
foo(long double,long double)703 inline constexpr long double foo(long double, long double) { return 2;}
foo(_Tp,_Up)704 template<typename _Tp, typename _Up> inline constexpr double foo(_Tp, _Up) { return 3;}
705 };
706 using X::foo;
foo(double,double)707 inline __device__ double foo(double, double) { return 4;}
foo(float,int)708 inline __device__ float foo(float, int) { return 5;}
foo(int,int)709 inline __device__ float foo(int, int) { return 6;}
foo(double,int)710 inline __device__ double foo(double, int) { return 7;}
foo(float,float)711 inline __device__ float foo(float, float) { return 9;}
foo(_Tp,_Up)712 template<typename _Tp, typename _Up> inline __device__ double foo(_Tp, _Up) { return 10;}
713
g()714 int g() {
715 return [](){
716 return foo(0.0, 2);
717 }();
718 }
719 }
720