granicus.if.org Git - clang/blob - test/SemaCUDA/function-overload.cu

   1 // REQUIRES: x86-registered-target
   2 // REQUIRES: nvptx-registered-target
   3
   4 // RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fsyntax-only -verify %s
   5 // RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fsyntax-only -fcuda-is-device -verify %s
   6
   7 #include "Inputs/cuda.h"
   8
   9 // Opaque return types used to check that we pick the right overloads.
  10 struct HostReturnTy {};
  11 struct HostReturnTy2 {};
  12 struct DeviceReturnTy {};
  13 struct DeviceReturnTy2 {};
  14 struct HostDeviceReturnTy {};
  15 struct TemplateReturnTy {};
  16
  17 typedef HostReturnTy (*HostFnPtr)();
  18 typedef DeviceReturnTy (*DeviceFnPtr)();
  19 typedef HostDeviceReturnTy (*HostDeviceFnPtr)();
  20 typedef void (*GlobalFnPtr)();  // __global__ functions must return void.
  21
  22 // CurrentReturnTy is {HostReturnTy,DeviceReturnTy} during {host,device}
  23 // compilation.
  24 #ifdef __CUDA_ARCH__
  25 typedef DeviceReturnTy CurrentReturnTy;
  26 #else
  27 typedef HostReturnTy CurrentReturnTy;
  28 #endif
  29
  30 // CurrentFnPtr is a function pointer to a {host,device} function during
  31 // {host,device} compilation.
  32 typedef CurrentReturnTy (*CurrentFnPtr)();
  33
  34 // Host and unattributed functions can't be overloaded.
  35 __host__ void hh() {} // expected-note {{previous definition is here}}
  36 void hh() {} // expected-error {{redefinition of 'hh'}}
  37
  38 // H/D overloading is OK.
  39 __host__ HostReturnTy dh() { return HostReturnTy(); }
  40 __device__ DeviceReturnTy dh() { return DeviceReturnTy(); }
  41
  42 // H/HD and D/HD are not allowed.
  43 __host__ __device__ int hdh() { return 0; } // expected-note {{previous declaration is here}}
  44 __host__ int hdh() { return 0; }
  45 // expected-error@-1 {{__host__ function 'hdh' cannot overload __host__ __device__ function 'hdh'}}
  46
  47 __host__ int hhd() { return 0; }            // expected-note {{previous declaration is here}}
  48 __host__ __device__ int hhd() { return 0; }
  49 // expected-error@-1 {{__host__ __device__ function 'hhd' cannot overload __host__ function 'hhd'}}
  50
  51 __host__ __device__ int hdd() { return 0; } // expected-note {{previous declaration is here}}
  52 __device__ int hdd() { return 0; }
  53 // expected-error@-1 {{__device__ function 'hdd' cannot overload __host__ __device__ function 'hdd'}}
  54
  55 __device__ int dhd() { return 0; }          // expected-note {{previous declaration is here}}
  56 __host__ __device__ int dhd() { return 0; }
  57 // expected-error@-1 {{__host__ __device__ function 'dhd' cannot overload __device__ function 'dhd'}}
  58
  59 // Same tests for extern "C" functions.
  60 extern "C" __host__ int chh() { return 0; } // expected-note {{previous definition is here}}
  61 extern "C" int chh() { return 0; }          // expected-error {{redefinition of 'chh'}}
  62
  63 // H/D overloading is OK.
  64 extern "C" __device__ DeviceReturnTy cdh() { return DeviceReturnTy(); }
  65 extern "C" __host__ HostReturnTy cdh() { return HostReturnTy(); }
  66
  67 // H/HD and D/HD overloading is not allowed.
  68 extern "C" __host__ __device__ int chhd1() { return 0; } // expected-note {{previous declaration is here}}
  69 extern "C" __host__ int chhd1() { return 0; }
  70 // expected-error@-1 {{__host__ function 'chhd1' cannot overload __host__ __device__ function 'chhd1'}}
  71
  72 extern "C" __host__ int chhd2() { return 0; } // expected-note {{previous declaration is here}}
  73 extern "C" __host__ __device__ int chhd2() { return 0; }
  74 // expected-error@-1 {{__host__ __device__ function 'chhd2' cannot overload __host__ function 'chhd2'}}
  75
  76 // Helper functions to verify calling restrictions.
  77 __device__ DeviceReturnTy d() { return DeviceReturnTy(); }
  78 // expected-note@-1 1+ {{'d' declared here}}
  79 // expected-note@-2 1+ {{candidate function not viable: call to __device__ function from __host__ function}}
  80 // expected-note@-3 0+ {{candidate function not viable: call to __device__ function from __host__ __device__ function}}
  81
  82 __host__ HostReturnTy h() { return HostReturnTy(); }
  83 // expected-note@-1 1+ {{'h' declared here}}
  84 // expected-note@-2 1+ {{candidate function not viable: call to __host__ function from __device__ function}}
  85 // expected-note@-3 0+ {{candidate function not viable: call to __host__ function from __host__ __device__ function}}
  86 // expected-note@-4 1+ {{candidate function not viable: call to __host__ function from __global__ function}}
  87
  88 __global__ void g() {}
  89 // expected-note@-1 1+ {{'g' declared here}}
  90 // expected-note@-2 1+ {{candidate function not viable: call to __global__ function from __device__ function}}
  91 // expected-note@-3 0+ {{candidate function not viable: call to __global__ function from __host__ __device__ function}}
  92 // expected-note@-4 1+ {{candidate function not viable: call to __global__ function from __global__ function}}
  93
  94 extern "C" __device__ DeviceReturnTy cd() { return DeviceReturnTy(); }
  95 // expected-note@-1 1+ {{'cd' declared here}}
  96 // expected-note@-2 1+ {{candidate function not viable: call to __device__ function from __host__ function}}
  97 // expected-note@-3 0+ {{candidate function not viable: call to __device__ function from __host__ __device__ function}}
  98
  99 extern "C" __host__ HostReturnTy ch() { return HostReturnTy(); }
 100 // expected-note@-1 1+ {{'ch' declared here}}
 101 // expected-note@-2 1+ {{candidate function not viable: call to __host__ function from __device__ function}}
 102 // expected-note@-3 0+ {{candidate function not viable: call to __host__ function from __host__ __device__ function}}
 103 // expected-note@-4 1+ {{candidate function not viable: call to __host__ function from __global__ function}}
 104
 105 __host__ void hostf() {
 106   DeviceFnPtr fp_d = d;         // expected-error {{reference to __device__ function 'd' in __host__ function}}
 107   DeviceReturnTy ret_d = d();   // expected-error {{no matching function for call to 'd'}}
 108   DeviceFnPtr fp_cd = cd;       // expected-error {{reference to __device__ function 'cd' in __host__ function}}
 109   DeviceReturnTy ret_cd = cd(); // expected-error {{no matching function for call to 'cd'}}
 110
 111   HostFnPtr fp_h = h;
 112   HostReturnTy ret_h = h();
 113   HostFnPtr fp_ch = ch;
 114   HostReturnTy ret_ch = ch();
 115
 116   HostFnPtr fp_dh = dh;
 117   HostReturnTy ret_dh = dh();
 118   HostFnPtr fp_cdh = cdh;
 119   HostReturnTy ret_cdh = cdh();
 120
 121   GlobalFnPtr fp_g = g;
 122   g(); // expected-error {{call to global function 'g' not configured}}
 123   g<<<0, 0>>>();
 124 }
 125
 126 __device__ void devicef() {
 127   DeviceFnPtr fp_d = d;
 128   DeviceReturnTy ret_d = d();
 129   DeviceFnPtr fp_cd = cd;
 130   DeviceReturnTy ret_cd = cd();
 131
 132   HostFnPtr fp_h = h;         // expected-error {{reference to __host__ function 'h' in __device__ function}}
 133   HostReturnTy ret_h = h();   // expected-error {{no matching function for call to 'h'}}
 134   HostFnPtr fp_ch = ch;       // expected-error {{reference to __host__ function 'ch' in __device__ function}}
 135   HostReturnTy ret_ch = ch(); // expected-error {{no matching function for call to 'ch'}}
 136
 137   DeviceFnPtr fp_dh = dh;
 138   DeviceReturnTy ret_dh = dh();
 139   DeviceFnPtr fp_cdh = cdh;
 140   DeviceReturnTy ret_cdh = cdh();
 141
 142   GlobalFnPtr fp_g = g; // expected-error {{reference to __global__ function 'g' in __device__ function}}
 143   g(); // expected-error {{no matching function for call to 'g'}}
 144   g<<<0,0>>>(); // expected-error {{reference to __global__ function 'g' in __device__ function}}
 145 }
 146
 147 __global__ void globalf() {
 148   DeviceFnPtr fp_d = d;
 149   DeviceReturnTy ret_d = d();
 150   DeviceFnPtr fp_cd = cd;
 151   DeviceReturnTy ret_cd = cd();
 152
 153   HostFnPtr fp_h = h;         // expected-error {{reference to __host__ function 'h' in __global__ function}}
 154   HostReturnTy ret_h = h();   // expected-error {{no matching function for call to 'h'}}
 155   HostFnPtr fp_ch = ch;       // expected-error {{reference to __host__ function 'ch' in __global__ function}}
 156   HostReturnTy ret_ch = ch(); // expected-error {{no matching function for call to 'ch'}}
 157
 158   DeviceFnPtr fp_dh = dh;
 159   DeviceReturnTy ret_dh = dh();
 160   DeviceFnPtr fp_cdh = cdh;
 161   DeviceReturnTy ret_cdh = cdh();
 162
 163   GlobalFnPtr fp_g = g; // expected-error {{reference to __global__ function 'g' in __global__ function}}
 164   g(); // expected-error {{no matching function for call to 'g'}}
 165   g<<<0,0>>>(); // expected-error {{reference to __global__ function 'g' in __global__ function}}
 166 }
 167
 168 __host__ __device__ void hostdevicef() {
 169   DeviceFnPtr fp_d = d;
 170   DeviceReturnTy ret_d = d();
 171   DeviceFnPtr fp_cd = cd;
 172   DeviceReturnTy ret_cd = cd();
 173 #if !defined(__CUDA_ARCH__)
 174   // expected-error@-5 {{reference to __device__ function 'd' in __host__ __device__ function}}
 175   // expected-error@-5 {{reference to __device__ function 'd' in __host__ __device__ function}}
 176   // expected-error@-5 {{reference to __device__ function 'cd' in __host__ __device__ function}}
 177   // expected-error@-5 {{reference to __device__ function 'cd' in __host__ __device__ function}}
 178 #endif
 179
 180   HostFnPtr fp_h = h;
 181   HostReturnTy ret_h = h();
 182   HostFnPtr fp_ch = ch;
 183   HostReturnTy ret_ch = ch();
 184 #if defined(__CUDA_ARCH__)
 185   // expected-error@-5 {{reference to __host__ function 'h' in __host__ __device__ function}}
 186   // expected-error@-5 {{reference to __host__ function 'h' in __host__ __device__ function}}
 187   // expected-error@-5 {{reference to __host__ function 'ch' in __host__ __device__ function}}
 188   // expected-error@-5 {{reference to __host__ function 'ch' in __host__ __device__ function}}
 189 #endif
 190
 191   CurrentFnPtr fp_dh = dh;
 192   CurrentReturnTy ret_dh = dh();
 193   CurrentFnPtr fp_cdh = cdh;
 194   CurrentReturnTy ret_cdh = cdh();
 195
 196   GlobalFnPtr fp_g = g;
 197 #if defined(__CUDA_ARCH__)
 198   // expected-error@-2 {{reference to __global__ function 'g' in __host__ __device__ function}}
 199 #endif
 200
 201   g();
 202 #if defined (__CUDA_ARCH__)
 203   // expected-error@-2 {{reference to __global__ function 'g' in __host__ __device__ function}}
 204 #else
 205   // expected-error@-4 {{call to global function 'g' not configured}}
 206 #endif
 207
 208   g<<<0,0>>>();
 209 #if defined(__CUDA_ARCH__)
 210   // expected-error@-2 {{reference to __global__ function 'g' in __host__ __device__ function}}
 211 #endif
 212 }
 213
 214 // Test for address of overloaded function resolution in the global context.
 215 HostFnPtr fp_h = h;
 216 HostFnPtr fp_ch = ch;
 217 CurrentFnPtr fp_dh = dh;
 218 CurrentFnPtr fp_cdh = cdh;
 219 GlobalFnPtr fp_g = g;
 220
 221
 222 // Test overloading of destructors
 223 // Can't mix H and unattributed destructors
 224 struct d_h {
 225   ~d_h() {} // expected-note {{previous definition is here}}
 226   __host__ ~d_h() {} // expected-error {{destructor cannot be redeclared}}
 227 };
 228
 229 // HD is OK
 230 struct d_hd {
 231   __host__ __device__ ~d_hd() {}
 232 };
 233
 234 // Test overloading of member functions
 235 struct m_h {
 236   void operator delete(void *ptr); // expected-note {{previous declaration is here}}
 237   __host__ void operator delete(void *ptr); // expected-error {{class member cannot be redeclared}}
 238 };
 239
 240 // D/H overloading is OK
 241 struct m_dh {
 242   __device__ void operator delete(void *ptr);
 243   __host__ void operator delete(void *ptr);
 244 };
 245
 246 // HD by itself is OK
 247 struct m_hd {
 248   __device__ __host__ void operator delete(void *ptr);
 249 };
 250
 251 struct m_hhd {
 252   __host__ void operator delete(void *ptr) {} // expected-note {{previous declaration is here}}
 253   __host__ __device__ void operator delete(void *ptr) {}
 254   // expected-error@-1 {{__host__ __device__ function 'operator delete' cannot overload __host__ function 'operator delete'}}
 255 };
 256
 257 struct m_hdh {
 258   __host__ __device__ void operator delete(void *ptr) {} // expected-note {{previous declaration is here}}
 259   __host__ void operator delete(void *ptr) {}
 260   // expected-error@-1 {{__host__ function 'operator delete' cannot overload __host__ __device__ function 'operator delete'}}
 261 };
 262
 263 struct m_dhd {
 264   __device__ void operator delete(void *ptr) {} // expected-note {{previous declaration is here}}
 265   __host__ __device__ void operator delete(void *ptr) {}
 266   // expected-error@-1 {{__host__ __device__ function 'operator delete' cannot overload __device__ function 'operator delete'}}
 267 };
 268
 269 struct m_hdd {
 270   __host__ __device__ void operator delete(void *ptr) {} // expected-note {{previous declaration is here}}
 271   __device__ void operator delete(void *ptr) {}
 272   // expected-error@-1 {{__device__ function 'operator delete' cannot overload __host__ __device__ function 'operator delete'}}
 273 };
 274
 275 // __global__ functions can't be overloaded based on attribute
 276 // difference.
 277 struct G {
 278   friend void friend_of_g(G &arg); // expected-note {{previous declaration is here}}
 279 private:
 280   int x; // expected-note {{declared private here}}
 281 };
 282 __global__ void friend_of_g(G &arg) { int x = arg.x; }
 283 // expected-error@-1 {{__global__ function 'friend_of_g' cannot overload __host__ function 'friend_of_g'}}
 284 // expected-error@-2 {{'x' is a private member of 'G'}}
 285 void friend_of_g(G &arg) { int x = arg.x; }
 286
 287 // HD functions are sometimes allowed to call H or D functions -- this
 288 // is an artifact of the source-to-source splitting performed by nvcc
 289 // that we need to mimic. During device mode compilation in nvcc, host
 290 // functions aren't present at all, so don't participate in
 291 // overloading. But in clang, H and D functions are present in both
 292 // compilation modes. Clang normally uses the target attribute as a
 293 // tiebreaker between overloads with otherwise identical priority, but
 294 // in order to match nvcc's behavior, we sometimes need to wholly
 295 // discard overloads that would not be present during compilation
 296 // under nvcc.
 297
 298 template <typename T> TemplateReturnTy template_vs_function(T arg) {
 299   return TemplateReturnTy();
 300 }
 301 __device__ DeviceReturnTy template_vs_function(float arg) {
 302   return DeviceReturnTy();
 303 }
 304
 305 // Here we expect to call the templated function during host compilation, even
 306 // if -fcuda-disable-target-call-checks is passed, and even though C++ overload
 307 // rules prefer the non-templated function.
 308 __host__ __device__ void test_host_device_calls_template(void) {
 309 #ifdef __CUDA_ARCH__
 310   typedef DeviceReturnTy ExpectedReturnTy;
 311 #else
 312   typedef TemplateReturnTy ExpectedReturnTy;
 313 #endif
 314
 315   ExpectedReturnTy ret1 = template_vs_function(1.0f);
 316   ExpectedReturnTy ret2 = template_vs_function(2.0);
 317 }
 318
 319 // Calls from __host__ and __device__ functions should always call the
 320 // overloaded function that matches their mode.
 321 __host__ void test_host_calls_template_fn() {
 322   TemplateReturnTy ret1 = template_vs_function(1.0f);
 323   TemplateReturnTy ret2 = template_vs_function(2.0);
 324 }
 325
 326 __device__ void test_device_calls_template_fn() {
 327   DeviceReturnTy ret1 = template_vs_function(1.0f);
 328   DeviceReturnTy ret2 = template_vs_function(2.0);
 329 }
 330
 331 // If we have a mix of HD and H-only or D-only candidates in the overload set,
 332 // normal C++ overload resolution rules apply first.
 333 template <typename T> TemplateReturnTy template_vs_hd_function(T arg)
 334 #ifdef __CUDA_ARCH__
 335 //expected-note@-2 {{declared here}}
 336 #endif
 337 {
 338   return TemplateReturnTy();
 339 }
 340 __host__ __device__ HostDeviceReturnTy template_vs_hd_function(float arg) {
 341   return HostDeviceReturnTy();
 342 }
 343
 344 __host__ __device__ void test_host_device_calls_hd_template() {
 345   HostDeviceReturnTy ret1 = template_vs_hd_function(1.0f);
 346   TemplateReturnTy ret2 = template_vs_hd_function(1);
 347 #ifdef __CUDA_ARCH__
 348   // expected-error@-2 {{reference to __host__ function 'template_vs_hd_function<int>' in __host__ __device__ function}}
 349 #endif
 350 }
 351
 352 __host__ void test_host_calls_hd_template() {
 353   HostDeviceReturnTy ret1 = template_vs_hd_function(1.0f);
 354   TemplateReturnTy ret2 = template_vs_hd_function(1);
 355 }
 356
 357 __device__ void test_device_calls_hd_template() {
 358   HostDeviceReturnTy ret1 = template_vs_hd_function(1.0f);
 359   // Host-only function template is not callable with strict call checks,
 360   // so for device side HD function will be the only choice.
 361   HostDeviceReturnTy ret2 = template_vs_hd_function(1);
 362 }
 363
 364 // Check that overloads still work the same way on both host and
 365 // device side when the overload set contains only functions from one
 366 // side of compilation.
 367 __device__ DeviceReturnTy device_only_function(int arg) { return DeviceReturnTy(); }
 368 __device__ DeviceReturnTy2 device_only_function(float arg) { return DeviceReturnTy2(); }
 369 #ifndef __CUDA_ARCH__
 370   // expected-note@-3 {{'device_only_function' declared here}}
 371   // expected-note@-3 {{'device_only_function' declared here}}
 372 #endif
 373 __host__ HostReturnTy host_only_function(int arg) { return HostReturnTy(); }
 374 __host__ HostReturnTy2 host_only_function(float arg) { return HostReturnTy2(); }
 375 #ifdef __CUDA_ARCH__
 376   // expected-note@-3 {{'host_only_function' declared here}}
 377   // expected-note@-3 {{'host_only_function' declared here}}
 378 #endif
 379
 380 __host__ __device__ void test_host_device_single_side_overloading() {
 381   DeviceReturnTy ret1 = device_only_function(1);
 382   DeviceReturnTy2 ret2 = device_only_function(1.0f);
 383 #ifndef __CUDA_ARCH__
 384   // expected-error@-3 {{reference to __device__ function 'device_only_function' in __host__ __device__ function}}
 385   // expected-error@-3 {{reference to __device__ function 'device_only_function' in __host__ __device__ function}}
 386 #endif
 387   HostReturnTy ret3 = host_only_function(1);
 388   HostReturnTy2 ret4 = host_only_function(1.0f);
 389 #ifdef __CUDA_ARCH__
 390   // expected-error@-3 {{reference to __host__ function 'host_only_function' in __host__ __device__ function}}
 391   // expected-error@-3 {{reference to __host__ function 'host_only_function' in __host__ __device__ function}}
 392 #endif
 393 }
 394
 395 // Verify that we allow overloading function templates.
 396 template <typename T> __host__ T template_overload(const T &a) { return a; };
 397 template <typename T> __device__ T template_overload(const T &a) { return a; };
 398
 399 __host__ void test_host_template_overload() {
 400   template_overload(1); // OK. Attribute-based overloading picks __host__ variant.
 401 }
 402 __device__ void test_device_template_overload() {
 403   template_overload(1); // OK. Attribute-based overloading picks __device__ variant.
 404 }
 405
 406 // Two classes with `operator-` defined. One of them is device only.
 407 struct C1;
 408 struct C2;
 409 __device__
 410 int operator-(const C1 &x, const C1 &y);
 411 int operator-(const C2 &x, const C2 &y);
 412
 413 template <typename T>
 414 __host__ __device__ int constexpr_overload(const T &x, const T &y) {
 415   return x - y;
 416 }
 417
 418 // Verify that function overloading doesn't prune candidate wrongly.
 419 int test_constexpr_overload(C2 &x, C2 &y) {
 420   return constexpr_overload(x, y);
 421 }