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 CurrentFnPtr fp_dh = dh;
226 CurrentFnPtr fp_cdh = cdh;
227 GlobalFnPtr fp_g = g;
228
229
230 // Test overloading of destructors
231 // Can't mix H and unattributed destructors
232 struct d_h {
~d_hd_h233 ~d_h() {} // expected-note {{previous definition is here}}
~d_hd_h234 __host__ ~d_h() {} // expected-error {{destructor cannot be redeclared}}
235 };
236
237 // HD is OK
238 struct d_hd {
~d_hdd_hd239 __host__ __device__ ~d_hd() {}
240 };
241
242 // Test overloading of member functions
243 struct m_h {
244 void operator delete(void *ptr); // expected-note {{previous declaration is here}}
245 __host__ void operator delete(void *ptr); // expected-error {{class member cannot be redeclared}}
246 };
247
248 // D/H overloading is OK
249 struct m_dh {
250 __device__ void operator delete(void *ptr);
251 __host__ void operator delete(void *ptr);
252 };
253
254 // HD by itself is OK
255 struct m_hd {
256 __device__ __host__ void operator delete(void *ptr);
257 };
258
259 struct m_hhd {
operator deletem_hhd260 __host__ void operator delete(void *ptr) {} // expected-note {{previous declaration is here}}
operator deletem_hhd261 __host__ __device__ void operator delete(void *ptr) {}
262 // expected-error@-1 {{__host__ __device__ function 'operator delete' cannot overload __host__ function 'operator delete'}}
263 };
264
265 struct m_hdh {
operator deletem_hdh266 __host__ __device__ void operator delete(void *ptr) {} // expected-note {{previous declaration is here}}
operator deletem_hdh267 __host__ void operator delete(void *ptr) {}
268 // expected-error@-1 {{__host__ function 'operator delete' cannot overload __host__ __device__ function 'operator delete'}}
269 };
270
271 struct m_dhd {
operator deletem_dhd272 __device__ void operator delete(void *ptr) {} // expected-note {{previous declaration is here}}
operator deletem_dhd273 __host__ __device__ void operator delete(void *ptr) {}
274 // expected-error@-1 {{__host__ __device__ function 'operator delete' cannot overload __device__ function 'operator delete'}}
275 };
276
277 struct m_hdd {
operator deletem_hdd278 __host__ __device__ void operator delete(void *ptr) {} // expected-note {{previous declaration is here}}
operator deletem_hdd279 __device__ void operator delete(void *ptr) {}
280 // expected-error@-1 {{__device__ function 'operator delete' cannot overload __host__ __device__ function 'operator delete'}}
281 };
282
283 // __global__ functions can't be overloaded based on attribute
284 // difference.
285 struct G {
286 friend void friend_of_g(G &arg); // expected-note {{previous declaration is here}}
287 private:
288 int x; // expected-note {{declared private here}}
289 };
friend_of_g(G & arg)290 __global__ void friend_of_g(G &arg) { int x = arg.x; }
291 // expected-error@-1 {{__global__ function 'friend_of_g' cannot overload __host__ function 'friend_of_g'}}
292 // expected-error@-2 {{'x' is a private member of 'G'}}
friend_of_g(G & arg)293 void friend_of_g(G &arg) { int x = arg.x; }
294
295 // HD functions are sometimes allowed to call H or D functions -- this
296 // is an artifact of the source-to-source splitting performed by nvcc
297 // that we need to mimic. During device mode compilation in nvcc, host
298 // functions aren't present at all, so don't participate in
299 // overloading. But in clang, H and D functions are present in both
300 // compilation modes. Clang normally uses the target attribute as a
301 // tiebreaker between overloads with otherwise identical priority, but
302 // in order to match nvcc's behavior, we sometimes need to wholly
303 // discard overloads that would not be present during compilation
304 // under nvcc.
305
template_vs_function(T arg)306 template <typename T> TemplateReturnTy template_vs_function(T arg) {
307 return TemplateReturnTy();
308 }
template_vs_function(float arg)309 __device__ DeviceReturnTy template_vs_function(float arg) {
310 return DeviceReturnTy();
311 }
312
313 // Here we expect to call the templated function during host compilation, even
314 // if -fcuda-disable-target-call-checks is passed, and even though C++ overload
315 // rules prefer the non-templated function.
test_host_device_calls_template(void)316 __host__ __device__ void test_host_device_calls_template(void) {
317 #ifdef __CUDA_ARCH__
318 typedef DeviceReturnTy ExpectedReturnTy;
319 #else
320 typedef TemplateReturnTy ExpectedReturnTy;
321 #endif
322
323 ExpectedReturnTy ret1 = template_vs_function(1.0f);
324 ExpectedReturnTy ret2 = template_vs_function(2.0);
325 }
326
327 // Calls from __host__ and __device__ functions should always call the
328 // overloaded function that matches their mode.
test_host_calls_template_fn()329 __host__ void test_host_calls_template_fn() {
330 TemplateReturnTy ret1 = template_vs_function(1.0f);
331 TemplateReturnTy ret2 = template_vs_function(2.0);
332 }
333
test_device_calls_template_fn()334 __device__ void test_device_calls_template_fn() {
335 DeviceReturnTy ret1 = template_vs_function(1.0f);
336 DeviceReturnTy ret2 = template_vs_function(2.0);
337 }
338
339 // If we have a mix of HD and H-only or D-only candidates in the overload set,
340 // normal C++ overload resolution rules apply first.
template_vs_hd_function(T arg)341 template <typename T> TemplateReturnTy template_vs_hd_function(T arg)
342 // devnodeferonly-note@-1{{'template_vs_hd_function<int>' declared here}}
343 {
344 return TemplateReturnTy();
345 }
template_vs_hd_function(float arg)346 __host__ __device__ HostDeviceReturnTy template_vs_hd_function(float arg) {
347 return HostDeviceReturnTy();
348 }
349
test_host_device_calls_hd_template()350 __host__ __device__ void test_host_device_calls_hd_template() {
351 #if __CUDA_ARCH__ && DEFER
352 typedef HostDeviceReturnTy ExpectedReturnTy;
353 #else
354 typedef TemplateReturnTy ExpectedReturnTy;
355 #endif
356 HostDeviceReturnTy ret1 = template_vs_hd_function(1.0f);
357 ExpectedReturnTy ret2 = template_vs_hd_function(1);
358 // devnodeferonly-error@-1{{reference to __host__ function 'template_vs_hd_function<int>' in __host__ __device__ function}}
359 }
360
test_host_calls_hd_template()361 __host__ void test_host_calls_hd_template() {
362 HostDeviceReturnTy ret1 = template_vs_hd_function(1.0f);
363 TemplateReturnTy ret2 = template_vs_hd_function(1);
364 }
365
test_device_calls_hd_template()366 __device__ void test_device_calls_hd_template() {
367 HostDeviceReturnTy ret1 = template_vs_hd_function(1.0f);
368 // Host-only function template is not callable with strict call checks,
369 // so for device side HD function will be the only choice.
370 HostDeviceReturnTy ret2 = template_vs_hd_function(1);
371 }
372
373 // Check that overloads still work the same way on both host and
374 // device side when the overload set contains only functions from one
375 // side of compilation.
device_only_function(int arg)376 __device__ DeviceReturnTy device_only_function(int arg) { return DeviceReturnTy(); }
device_only_function(float arg)377 __device__ DeviceReturnTy2 device_only_function(float arg) { return DeviceReturnTy2(); }
378 #ifndef __CUDA_ARCH__
379 // expected-note@-3 2{{'device_only_function' declared here}}
380 // expected-note@-3 2{{'device_only_function' declared here}}
381 #endif
host_only_function(int arg)382 __host__ HostReturnTy host_only_function(int arg) { return HostReturnTy(); }
host_only_function(float arg)383 __host__ HostReturnTy2 host_only_function(float arg) { return HostReturnTy2(); }
384 #ifdef __CUDA_ARCH__
385 // expected-note@-3 2{{'host_only_function' declared here}}
386 // expected-note@-3 2{{'host_only_function' declared here}}
387 #endif
388
test_host_device_single_side_overloading()389 __host__ __device__ void test_host_device_single_side_overloading() {
390 DeviceReturnTy ret1 = device_only_function(1);
391 DeviceReturnTy2 ret2 = device_only_function(1.0f);
392 #ifndef __CUDA_ARCH__
393 // expected-error@-3 {{reference to __device__ function 'device_only_function' in __host__ __device__ function}}
394 // expected-error@-3 {{reference to __device__ function 'device_only_function' in __host__ __device__ function}}
395 #endif
396 HostReturnTy ret3 = host_only_function(1);
397 HostReturnTy2 ret4 = host_only_function(1.0f);
398 #ifdef __CUDA_ARCH__
399 // expected-error@-3 {{reference to __host__ function 'host_only_function' in __host__ __device__ function}}
400 // expected-error@-3 {{reference to __host__ function 'host_only_function' in __host__ __device__ function}}
401 #endif
402 }
403
404 // wrong-sided overloading should not cause diagnostic unless it is emitted.
405 // This inline function is not emitted.
test_host_device_wrong_side_overloading_inline_no_diag()406 inline __host__ __device__ void test_host_device_wrong_side_overloading_inline_no_diag() {
407 DeviceReturnTy ret1 = device_only_function(1);
408 DeviceReturnTy2 ret2 = device_only_function(1.0f);
409 HostReturnTy ret3 = host_only_function(1);
410 HostReturnTy2 ret4 = host_only_function(1.0f);
411 }
412
413 // wrong-sided overloading should cause diagnostic if it is emitted.
414 // This inline function is emitted since it is called by an emitted function.
test_host_device_wrong_side_overloading_inline_diag()415 inline __host__ __device__ void test_host_device_wrong_side_overloading_inline_diag() {
416 DeviceReturnTy ret1 = device_only_function(1);
417 DeviceReturnTy2 ret2 = device_only_function(1.0f);
418 #ifndef __CUDA_ARCH__
419 // expected-error@-3 {{reference to __device__ function 'device_only_function' in __host__ __device__ function}}
420 // expected-error@-3 {{reference to __device__ function 'device_only_function' in __host__ __device__ function}}
421 #endif
422 HostReturnTy ret3 = host_only_function(1);
423 HostReturnTy2 ret4 = host_only_function(1.0f);
424 #ifdef __CUDA_ARCH__
425 // expected-error@-3 {{reference to __host__ function 'host_only_function' in __host__ __device__ function}}
426 // expected-error@-3 {{reference to __host__ function 'host_only_function' in __host__ __device__ function}}
427 #endif
428 }
429
test_host_device_wrong_side_overloading_inline_diag_caller()430 __host__ __device__ void test_host_device_wrong_side_overloading_inline_diag_caller() {
431 test_host_device_wrong_side_overloading_inline_diag();
432 // expected-note@-1 {{called by 'test_host_device_wrong_side_overloading_inline_diag_caller'}}
433 }
434
435 // Verify that we allow overloading function templates.
template_overload(const T & a)436 template <typename T> __host__ T template_overload(const T &a) { return a; };
template_overload(const T & a)437 template <typename T> __device__ T template_overload(const T &a) { return a; };
438
test_host_template_overload()439 __host__ void test_host_template_overload() {
440 template_overload(1); // OK. Attribute-based overloading picks __host__ variant.
441 }
test_device_template_overload()442 __device__ void test_device_template_overload() {
443 template_overload(1); // OK. Attribute-based overloading picks __device__ variant.
444 }
445
446 // Two classes with `operator-` defined. One of them is device only.
447 struct C1;
448 struct C2;
449 __device__
450 int operator-(const C1 &x, const C1 &y);
451 int operator-(const C2 &x, const C2 &y);
452
453 template <typename T>
constexpr_overload(const T & x,const T & y)454 __host__ __device__ int constexpr_overload(const T &x, const T &y) {
455 return x - y;
456 }
457
458 // Verify that function overloading doesn't prune candidate wrongly.
test_constexpr_overload(C2 & x,C2 & y)459 int test_constexpr_overload(C2 &x, C2 &y) {
460 return constexpr_overload(x, y);
461 }
462
463 // Verify no ambiguity for new operator.
464 void *a = new int;
465 __device__ void *b = new int;
466 // expected-error@-1{{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
467
468 // Verify no ambiguity for new operator.
469 template<typename _Tp> _Tp&& f();
470 template<typename _Tp, typename = decltype(new _Tp(f<_Tp>()))>
471 void __test();
472
foo()473 void foo() {
474 __test<int>();
475 }
476
477 // Test resolving implicit host device candidate vs wrong-sided candidate.
478 // In device compilation, implicit host device caller choose implicit host
479 // device candidate and wrong-sided candidate with equal preference.
480 // Resolution result should not change with/without pragma.
481 namespace ImplicitHostDeviceVsWrongSided {
482 HostReturnTy callee(double x);
483 #pragma clang force_cuda_host_device begin
484 HostDeviceReturnTy callee(int x);
implicit_hd_caller()485 inline HostReturnTy implicit_hd_caller() {
486 return callee(1.0);
487 }
488 #pragma clang force_cuda_host_device end
489 }
490
491 // Test resolving implicit host device candidate vs same-sided candidate.
492 // In host compilation, implicit host device caller choose implicit host
493 // device candidate and same-sided candidate with equal preference.
494 // Resolution result should not change with/without pragma.
495 namespace ImplicitHostDeviceVsSameSide {
496 HostReturnTy callee(int x);
497 #pragma clang force_cuda_host_device begin
498 HostDeviceReturnTy callee(double x);
implicit_hd_caller()499 inline HostDeviceReturnTy implicit_hd_caller() {
500 return callee(1.0);
501 }
502 #pragma clang force_cuda_host_device end
503 }
504
505 // Test resolving explicit host device candidate vs. wrong-sided candidate.
506 // When -fgpu-defer-diag is off, wrong-sided candidate is not excluded, therefore
507 // the first callee is chosen.
508 // When -fgpu-defer-diag is on, wrong-sided candidate is excluded, therefore
509 // the second callee is chosen.
510 namespace ExplicitHostDeviceVsWrongSided {
511 HostReturnTy callee(double x);
512 __host__ __device__ HostDeviceReturnTy callee(int x);
513 #if __CUDA_ARCH__ && DEFER
514 typedef HostDeviceReturnTy ExpectedRetTy;
515 #else
516 typedef HostReturnTy ExpectedRetTy;
517 #endif
explicit_hd_caller()518 inline __host__ __device__ ExpectedRetTy explicit_hd_caller() {
519 return callee(1.0);
520 }
521 }
522
523 // In the implicit host device function 'caller', the second 'callee' should be
524 // chosen since it has better match, even though it is an implicit host device
525 // function whereas the first 'callee' is a host function. A diagnostic will be
526 // emitted if the first 'callee' is chosen since deduced return type cannot be
527 // used before it is defined.
528 namespace ImplicitHostDeviceByConstExpr {
529 template <class a> a b;
530 auto callee(...);
531 template <class d> constexpr auto callee(d) -> decltype(0);
532 struct e {
gImplicitHostDeviceByConstExpr::e533 template <class ad, class... f> static auto g(ad, f...) {
534 return h<e, decltype(b<f>)...>;
535 }
536 struct i {
callerImplicitHostDeviceByConstExpr::e::i537 template <class, class... f> static constexpr auto caller(f... k) {
538 return callee(k...);
539 }
540 };
hImplicitHostDeviceByConstExpr::e541 template <class, class... f> static auto h() {
542 return i::caller<int, f...>;
543 }
544 };
545 class l {
l()546 l() {
547 e::g([] {}, this);
548 }
549 };
550 }
551
552 // Implicit HD candidate competes with device candidate.
553 // a and b have implicit HD copy ctor. In copy ctor of b, ctor of a is resolved.
554 // copy ctor of a should win over a(short), otherwise there will be ambiguity
555 // due to conversion operator.
556 namespace TestImplicitHDWithD {
557 struct a {
558 __device__ a(short);
559 __device__ operator unsigned() const;
560 __device__ operator int() const;
561 };
562 struct b {
563 a d;
564 };
f(b g)565 void f(b g) { b e = g; }
566 }
567
568 // Implicit HD candidate competes with host candidate.
569 // a and b have implicit HD copy ctor. In copy ctor of b, ctor of a is resolved.
570 // copy ctor of a should win over a(short), otherwise there will be ambiguity
571 // due to conversion operator.
572 namespace TestImplicitHDWithH {
573 struct a {
574 a(short);
575 __device__ operator unsigned() const;
576 __device__ operator int() const;
577 };
578 struct b {
579 a d;
580 };
f(b g)581 void f(b g) { b e = g; }
582 }
583
584 // Implicit HD candidate competes with HD candidate.
585 // a and b have implicit HD copy ctor. In copy ctor of b, ctor of a is resolved.
586 // copy ctor of a should win over a(short), otherwise there will be ambiguity
587 // due to conversion operator.
588 namespace TestImplicitHDWithHD {
589 struct a {
590 __host__ __device__ a(short);
591 __device__ operator unsigned() const;
592 __device__ operator int() const;
593 };
594 struct b {
595 a d;
596 };
f(b g)597 void f(b g) { b e = g; }
598 }
599
600 // HD candidate competes with H candidate.
601 // HD has type mismatch whereas H has type match.
602 // In device compilation, H wins when -fgpu-defer-diag is off and HD wins
603 // when -fgpu-defer-diags is on. In both cases the diagnostic should be
604 // deferred.
605 namespace TestDeferNoMatchingFuncNotEmitted {
606 template <typename> struct a {};
607 namespace b {
608 struct c : a<int> {};
609 template <typename d> void ag(d);
610 } // namespace b
611 template <typename ae>
ag(a<ae>)612 __host__ __device__ void ag(a<ae>) {
613 ae e;
614 ag(e);
615 }
f()616 void f() { (void)ag<b::c>; }
617 }
618
619 namespace TestDeferNoMatchingFuncEmitted {
620 template <typename> struct a {};
621 namespace b {
622 struct c : a<int> {};
623 template <typename d> void ag(d);
624 // devnodeferonly-note@-1{{'ag<TestDeferNoMatchingFuncEmitted::b::c>' declared here}}
625 } // namespace b
626 template <typename ae>
ag(a<ae>)627 __host__ __device__ void ag(a<ae>) {
628 ae e;
629 ag(e);
630 // devnodeferonly-error@-1{{reference to __host__ function 'ag<TestDeferNoMatchingFuncEmitted::b::c>' in __host__ __device__ function}}
631 // devdeferonly-error@-2{{no matching function for call to 'ag'}}
632 // devdeferonly-note@-3{{called by 'ag<TestDeferNoMatchingFuncEmitted::b::c>'}}
633 }
f()634 __host__ __device__ void f() { (void)ag<b::c>; }
635 // devnodeferonly-note@-1{{called by 'f'}}
636 // devdeferonly-note@-2{{called by 'f'}}
637 }
638
639 // Two HD candidates compete with H candidate.
640 // HDs have type mismatch whereas H has type match.
641 // In device compilation, H wins when -fgpu-defer-diag is off and two HD win
642 // when -fgpu-defer-diags is on. In both cases the diagnostic should be
643 // deferred.
644 namespace TestDeferAmbiguityNotEmitted {
645 template <typename> struct a {};
646 namespace b {
647 struct c : a<int> {};
648 template <typename d> void ag(d, int);
649 } // namespace b
650 template <typename ae>
ag(a<ae>,float)651 __host__ __device__ void ag(a<ae>, float) {
652 ae e;
653 ag(e, 1);
654 }
655 template <typename ae>
ag(a<ae>,double)656 __host__ __device__ void ag(a<ae>, double) {
657 }
f()658 void f() {
659 b::c x;
660 ag(x, 1);
661 }
662 }
663
664 namespace TestDeferAmbiguityEmitted {
665 template <typename> struct a {};
666 namespace b {
667 struct c : a<int> {};
668 template <typename d> void ag(d, int);
669 // devnodeferonly-note@-1{{'ag<TestDeferAmbiguityEmitted::b::c>' declared here}}
670 } // namespace b
671 template <typename ae>
ag(a<ae>,float)672 __host__ __device__ void ag(a<ae>, float) {
673 // devdeferonly-note@-1{{candidate function [with ae = int]}}
674 ae e;
675 ag(e, 1);
676 }
677 template <typename ae>
ag(a<ae>,double)678 __host__ __device__ void ag(a<ae>, double) {
679 // devdeferonly-note@-1{{candidate function [with ae = int]}}
680 }
f()681 __host__ __device__ void f() {
682 b::c x;
683 ag(x, 1);
684 // devnodeferonly-error@-1{{reference to __host__ function 'ag<TestDeferAmbiguityEmitted::b::c>' in __host__ __device__ function}}
685 // devdeferonly-error@-2{{call to 'ag' is ambiguous}}
686 }
687 }
688
689 // Implicit HD functions compute with H function and D function.
690 // In host compilation, foo(0.0, 2) should resolve to X::foo<double, int>.
691 // In device compilation, foo(0.0, 2) should resolve to foo(double, int).
692 // In either case there should be no ambiguity.
693 namespace TestImplicitHDWithHAndD {
694 namespace X {
foo(double,double)695 inline double foo(double, double) { return 0;}
foo(float,float)696 inline constexpr float foo(float, float) { return 1;}
foo(long double,long double)697 inline constexpr long double foo(long double, long double) { return 2;}
foo(_Tp,_Up)698 template<typename _Tp, typename _Up> inline constexpr double foo(_Tp, _Up) { return 3;}
699 };
700 using X::foo;
foo(double,double)701 inline __device__ double foo(double, double) { return 4;}
foo(float,int)702 inline __device__ float foo(float, int) { return 5;}
foo(int,int)703 inline __device__ float foo(int, int) { return 6;}
foo(double,int)704 inline __device__ double foo(double, int) { return 7;}
foo(float,float)705 inline __device__ float foo(float, float) { return 9;}
foo(_Tp,_Up)706 template<typename _Tp, typename _Up> inline __device__ double foo(_Tp, _Up) { return 10;}
707
g()708 int g() {
709 return [](){
710 return foo(0.0, 2);
711 }();
712 }
713 }
714