clang/test/SemaCUDA/function-overload.cu

   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.
  43 __host__ void hh() {} // expected-note {{previous definition is here}}
  44 void hh() {} // expected-error {{redefinition of 'hh'}}
  45
  46 // H/D overloading is OK.
  47 __host__ HostReturnTy dh() { return HostReturnTy(); }
  48 __device__ DeviceReturnTy dh() { return DeviceReturnTy(); }
  49
  50 // H/HD and D/HD are not allowed.
  51 __host__ __device__ int hdh() { return 0; } // expected-note {{previous declaration is here}}
  52 __host__ int hdh() { return 0; }
  53 // expected-error@-1 {{__host__ function 'hdh' cannot overload __host__ __device__ function 'hdh'}}
  54
  55 __host__ int hhd() { return 0; }            // expected-note {{previous declaration is here}}
  56 __host__ __device__ int hhd() { return 0; }
  57 // expected-error@-1 {{__host__ __device__ function 'hhd' cannot overload __host__ function 'hhd'}}
  58
  59 __host__ __device__ int hdd() { return 0; } // expected-note {{previous declaration is here}}
  60 __device__ int hdd() { return 0; }
  61 // expected-error@-1 {{__device__ function 'hdd' cannot overload __host__ __device__ function 'hdd'}}
  62
  63 __device__ int dhd() { return 0; }          // expected-note {{previous declaration is here}}
  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.
  68 extern "C" __host__ int chh() { return 0; } // expected-note {{previous definition is here}}
  69 extern "C" int chh() { return 0; }          // expected-error {{redefinition of 'chh'}}
  70
  71 // H/D overloading is OK.
  72 extern "C" __device__ DeviceReturnTy cdh() { return DeviceReturnTy(); }
  73 extern "C" __host__ HostReturnTy cdh() { return HostReturnTy(); }
  74
  75 // H/HD and D/HD overloading is not allowed.
  76 extern "C" __host__ __device__ int chhd1() { return 0; } // expected-note {{previous declaration is here}}
  77 extern "C" __host__ int chhd1() { return 0; }
  78 // expected-error@-1 {{__host__ function 'chhd1' cannot overload __host__ __device__ function 'chhd1'}}
  79
  80 extern "C" __host__ int chhd2() { return 0; } // expected-note {{previous declaration is here}}
  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.
  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
  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
  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
 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
 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
 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
 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
 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
 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 {
 239   ~d_h() {} // expected-note {{previous definition is here}}
 240   __host__ ~d_h() {} // expected-error {{destructor cannot be redeclared}}
 241 };
 242
 243 // HD is OK
 244 struct d_hd {
 245   __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 {
 266   __host__ void operator delete(void *ptr) {} // expected-note {{previous declaration is here}}
 267   __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 {
 272   __host__ __device__ void operator delete(void *ptr) {} // expected-note {{previous declaration is here}}
 273   __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 {
 278   __device__ void operator delete(void *ptr) {} // expected-note {{previous declaration is here}}
 279   __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 {
 284   __host__ __device__ void operator delete(void *ptr) {} // expected-note {{previous declaration is here}}
 285   __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 };
 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'}}
 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
 312 template <typename T> TemplateReturnTy template_vs_function(T arg) {
 313   return TemplateReturnTy();
 314 }
 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.
 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.
 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
 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.
 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 }
 352 __host__ __device__ HostDeviceReturnTy template_vs_hd_function(float arg) {
 353   return HostDeviceReturnTy();
 354 }
 355
 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
 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
 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.
 382 __device__ DeviceReturnTy device_only_function(int arg) { return DeviceReturnTy(); }
 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
 388 __host__ HostReturnTy host_only_function(int arg) { return HostReturnTy(); }
 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
 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.
 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.
 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
 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.
 442 template <typename T> __host__ T template_overload(const T &a) { return a; };
 443 template <typename T> __device__ T template_overload(const T &a) { return a; };
 444
 445 __host__ void test_host_template_overload() {
 446   template_overload(1); // OK. Attribute-based overloading picks __host__ variant.
 447 }
 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>
 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.
 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
 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);
 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);
 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
 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 {
 539   template <class ad, class... f> static auto g(ad, f...) {
 540     return h<e, decltype(b<f>)...>;
 541   }
 542   struct i {
 543     template <class, class... f> static constexpr auto caller(f... k) {
 544       return callee(k...);
 545     }
 546   };
 547   template <class, class... f> static auto h() {
 548     return i::caller<int, f...>;
 549   }
 550 };
 551 class 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   };
 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   };
 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   };
 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>
 618   __host__ __device__ void ag(a<ae>) {
 619     ae e;
 620     ag(e);
 621   }
 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>
 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   }
 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>
 657   __host__ __device__ void ag(a<ae>, float) {
 658     ae e;
 659     ag(e, 1);
 660   }
 661   template <typename ae>
 662   __host__ __device__ void ag(a<ae>, double) {
 663   }
 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>
 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>
 684   __host__ __device__ void ag(a<ae>, double) {
 685     // devdeferonly-note@-1{{candidate function [with ae = int]}}
 686   }
 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 {
 701     inline double foo(double, double) { return 0;}
 702     inline constexpr float foo(float, float) { return 1;}
 703     inline constexpr long double foo(long double, long double) { return 2;}
 704     template<typename _Tp, typename _Up> inline constexpr double foo(_Tp, _Up) { return 3;}
 705   };
 706   using X::foo;
 707   inline __device__ double foo(double, double) { return 4;}
 708   inline __device__ float foo(float, int) { return 5;}
 709   inline __device__ float foo(int, int) { return 6;}
 710   inline __device__ double foo(double, int) { return 7;}
 711   inline __device__ float foo(float, float) { return 9;}
 712   template<typename _Tp, typename _Up> inline __device__ double foo(_Tp, _Up) { return 10;}
 713
 714   int g() {
 715     return [](){
 716     return foo(0.0, 2);
 717     }();
 718   }
 719 }