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 definition is here}}
  44 __host__ int hdh() { return 0; }            // expected-error {{redefinition of 'hdh'}}
  45
  46 __host__ int hhd() { return 0; }            // expected-note {{previous definition is here}}
  47 __host__ __device__ int hhd() { return 0; } // expected-error {{redefinition of 'hhd'}}
  48 // expected-warning@-1 {{attribute declaration must precede definition}}
  49 // expected-note@-3 {{previous definition is here}}
  50
  51 __host__ __device__ int hdd() { return 0; } // expected-note {{previous definition is here}}
  52 __device__ int hdd() { return 0; }          // expected-error {{redefinition of 'hdd'}}
  53
  54 __device__ int dhd() { return 0; }          // expected-note {{previous definition is here}}
  55 __host__ __device__ int dhd() { return 0; } // expected-error {{redefinition of 'dhd'}}
  56 // expected-warning@-1 {{attribute declaration must precede definition}}
  57 // expected-note@-3 {{previous definition is here}}
  58
  59 // Same tests for extern "C" functions.
  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 definition is here}}
  69 extern "C" __host__ int chhd1() { return 0; }            // expected-error {{redefinition of 'chhd1'}}
  70
  71 extern "C" __host__ int chhd2() { return 0; }            // expected-note {{previous definition is here}}
  72 extern "C" __host__ __device__ int chhd2() { return 0; } // expected-error {{redefinition of 'chhd2'}}
  73 // expected-warning@-1 {{attribute declaration must precede definition}}
  74 // expected-note@-3 {{previous definition is here}}
  75
  76 // Helper functions to verify calling restrictions.
  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
 174   HostFnPtr fp_h = h;
 175   HostReturnTy ret_h = h();
 176   HostFnPtr fp_ch = ch;
 177   HostReturnTy ret_ch = ch();
 178
 179   CurrentFnPtr fp_dh = dh;
 180   CurrentReturnTy ret_dh = dh();
 181   CurrentFnPtr fp_cdh = cdh;
 182   CurrentReturnTy ret_cdh = cdh();
 183
 184   GlobalFnPtr fp_g = g;
 185 #if defined(__CUDA_ARCH__)
 186   // expected-error@-2 {{reference to __global__ function 'g' in __host__ __device__ function}}
 187 #endif
 188   g();
 189   g<<<0,0>>>();
 190 #if !defined(__CUDA_ARCH__)
 191   // expected-error@-3 {{call to global function g not configured}}
 192 #else
 193   // expected-error@-5 {{no matching function for call to 'g'}}
 194   // expected-error@-5 {{reference to __global__ function 'g' in __host__ __device__ function}}
 195 #endif  // __CUDA_ARCH__
 196 }
 197
 198 // Test for address of overloaded function resolution in the global context.
 199 HostFnPtr fp_h = h;
 200 HostFnPtr fp_ch = ch;
 201 CurrentFnPtr fp_dh = dh;
 202 CurrentFnPtr fp_cdh = cdh;
 203 GlobalFnPtr fp_g = g;
 204
 205
 206 // Test overloading of destructors
 207 // Can't mix H and unattributed destructors
 208 struct d_h {
 209   ~d_h() {} // expected-note {{previous declaration is here}}
 210   __host__ ~d_h() {} // expected-error {{destructor cannot be redeclared}}
 211 };
 212
 213 // H/D overloading is OK
 214 struct d_dh {
 215   __device__ ~d_dh() {}
 216   __host__ ~d_dh() {}
 217 };
 218
 219 // HD is OK
 220 struct d_hd {
 221   __host__ __device__ ~d_hd() {}
 222 };
 223
 224 // Mixing H/D and HD is not allowed.
 225 struct d_dhhd {
 226   __device__ ~d_dhhd() {}
 227   __host__ ~d_dhhd() {} // expected-note {{previous declaration is here}}
 228   __host__ __device__ ~d_dhhd() {} // expected-error {{destructor cannot be redeclared}}
 229 };
 230
 231 struct d_hhd {
 232   __host__ ~d_hhd() {} // expected-note {{previous declaration is here}}
 233   __host__ __device__ ~d_hhd() {} // expected-error {{destructor cannot be redeclared}}
 234 };
 235
 236 struct d_hdh {
 237   __host__ __device__ ~d_hdh() {} // expected-note {{previous declaration is here}}
 238   __host__ ~d_hdh() {} // expected-error {{destructor cannot be redeclared}}
 239 };
 240
 241 struct d_dhd {
 242   __device__ ~d_dhd() {} // expected-note {{previous declaration is here}}
 243   __host__ __device__ ~d_dhd() {} // expected-error {{destructor cannot be redeclared}}
 244 };
 245
 246 struct d_hdd {
 247   __host__ __device__ ~d_hdd() {} // expected-note {{previous declaration is here}}
 248   __device__ ~d_hdd() {} // expected-error {{destructor cannot be redeclared}}
 249 };
 250
 251 // Test overloading of member functions
 252 struct m_h {
 253   void operator delete(void *ptr); // expected-note {{previous declaration is here}}
 254   __host__ void operator delete(void *ptr); // expected-error {{class member cannot be redeclared}}
 255 };
 256
 257 // D/H overloading is OK
 258 struct m_dh {
 259   __device__ void operator delete(void *ptr);
 260   __host__ void operator delete(void *ptr);
 261 };
 262
 263 // HD by itself is OK
 264 struct m_hd {
 265   __device__ __host__ void operator delete(void *ptr);
 266 };
 267
 268 struct m_hhd {
 269   __host__ void operator delete(void *ptr) {} // expected-note {{previous declaration is here}}
 270   __host__ __device__ void operator delete(void *ptr) {} // expected-error {{class member cannot be redeclared}}
 271 };
 272
 273 struct m_hdh {
 274   __host__ __device__ void operator delete(void *ptr) {} // expected-note {{previous declaration is here}}
 275   __host__ void operator delete(void *ptr) {} // expected-error {{class member cannot be redeclared}}
 276 };
 277
 278 struct m_dhd {
 279   __device__ void operator delete(void *ptr) {} // expected-note {{previous declaration is here}}
 280   __host__ __device__ void operator delete(void *ptr) {} // expected-error {{class member cannot be redeclared}}
 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) {} // expected-error {{class member cannot be redeclared}}
 286 };
 287
 288 // __global__ functions can't be overloaded based on attribute
 289 // difference.
 290 struct G {
 291   friend void friend_of_g(G &arg);
 292 private:
 293   int x;
 294 };
 295 __global__ void friend_of_g(G &arg) { int x = arg.x; } // expected-note {{previous definition is here}}
 296 void friend_of_g(G &arg) { int x = arg.x; } // expected-error {{redefinition of 'friend_of_g'}}
 297
 298 // HD functions are sometimes allowed to call H or D functions -- this
 299 // is an artifact of the source-to-source splitting performed by nvcc
 300 // that we need to mimic. During device mode compilation in nvcc, host
 301 // functions aren't present at all, so don't participate in
 302 // overloading. But in clang, H and D functions are present in both
 303 // compilation modes. Clang normally uses the target attribute as a
 304 // tiebreaker between overloads with otherwise identical priority, but
 305 // in order to match nvcc's behavior, we sometimes need to wholly
 306 // discard overloads that would not be present during compilation
 307 // under nvcc.
 308
 309 template <typename T> TemplateReturnTy template_vs_function(T arg) {
 310   return TemplateReturnTy();
 311 }
 312 __device__ DeviceReturnTy template_vs_function(float arg) {
 313   return DeviceReturnTy();
 314 }
 315
 316 // Here we expect to call the templated function during host compilation, even
 317 // if -fcuda-disable-target-call-checks is passed, and even though C++ overload
 318 // rules prefer the non-templated function.
 319 __host__ __device__ void test_host_device_calls_template(void) {
 320 #ifdef __CUDA_ARCH__
 321   typedef DeviceReturnTy ExpectedReturnTy;
 322 #else
 323   typedef TemplateReturnTy ExpectedReturnTy;
 324 #endif
 325
 326   ExpectedReturnTy ret1 = template_vs_function(1.0f);
 327   ExpectedReturnTy ret2 = template_vs_function(2.0);
 328 }
 329
 330 // Calls from __host__ and __device__ functions should always call the
 331 // overloaded function that matches their mode.
 332 __host__ void test_host_calls_template_fn() {
 333   TemplateReturnTy ret1 = template_vs_function(1.0f);
 334   TemplateReturnTy ret2 = template_vs_function(2.0);
 335 }
 336
 337 __device__ void test_device_calls_template_fn() {
 338   DeviceReturnTy ret1 = template_vs_function(1.0f);
 339   DeviceReturnTy ret2 = template_vs_function(2.0);
 340 }
 341
 342 // If we have a mix of HD and H-only or D-only candidates in the overload set,
 343 // normal C++ overload resolution rules apply first.
 344 template <typename T> TemplateReturnTy template_vs_hd_function(T arg) {
 345   return TemplateReturnTy();
 346 }
 347 __host__ __device__ HostDeviceReturnTy template_vs_hd_function(float arg) {
 348   return HostDeviceReturnTy();
 349 }
 350
 351 __host__ __device__ void test_host_device_calls_hd_template() {
 352   HostDeviceReturnTy ret1 = template_vs_hd_function(1.0f);
 353   TemplateReturnTy ret2 = template_vs_hd_function(1);
 354 }
 355
 356 __host__ void test_host_calls_hd_template() {
 357   HostDeviceReturnTy ret1 = template_vs_hd_function(1.0f);
 358   TemplateReturnTy ret2 = template_vs_hd_function(1);
 359 }
 360
 361 __device__ void test_device_calls_hd_template() {
 362   HostDeviceReturnTy ret1 = template_vs_hd_function(1.0f);
 363   // Host-only function template is not callable with strict call checks,
 364   // so for device side HD function will be the only choice.
 365   HostDeviceReturnTy ret2 = template_vs_hd_function(1);
 366 }
 367
 368 // Check that overloads still work the same way on both host and
 369 // device side when the overload set contains only functions from one
 370 // side of compilation.
 371 __device__ DeviceReturnTy device_only_function(int arg) { return DeviceReturnTy(); }
 372 __device__ DeviceReturnTy2 device_only_function(float arg) { return DeviceReturnTy2(); }
 373 __host__ HostReturnTy host_only_function(int arg) { return HostReturnTy(); }
 374 __host__ HostReturnTy2 host_only_function(float arg) { return HostReturnTy2(); }
 375
 376 __host__ __device__ void test_host_device_single_side_overloading() {
 377   DeviceReturnTy ret1 = device_only_function(1);
 378   DeviceReturnTy2 ret2 = device_only_function(1.0f);
 379   HostReturnTy ret3 = host_only_function(1);
 380   HostReturnTy2 ret4 = host_only_function(1.0f);
 381 }
 382
 383 // Verify that we allow overloading function templates.
 384 template <typename T> __host__ T template_overload(const T &a) { return a; };
 385 template <typename T> __device__ T template_overload(const T &a) { return a; };
 386
 387 __host__ void test_host_template_overload() {
 388   template_overload(1); // OK. Attribute-based overloading picks __host__ variant.
 389 }
 390 __device__ void test_device_template_overload() {
 391   template_overload(1); // OK. Attribute-based overloading picks __device__ variant.
 392 }