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
