1 // REQUIRES: x86-registered-target 2 // REQUIRES: nvptx-registered-target 3 4 // Make sure we handle target overloads correctly. 5 // RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu \ 6 // RUN: -fsyntax-only -fcuda-target-overloads -verify %s 7 // RUN: %clang_cc1 -triple nvptx64-nvidia-cuda \ 8 // RUN: -fsyntax-only -fcuda-target-overloads -fcuda-is-device -verify %s 9 10 // Check target overloads handling with disabled call target checks. 11 // RUN: %clang_cc1 -DNOCHECKS -triple x86_64-unknown-linux-gnu -fsyntax-only \ 12 // RUN: -fcuda-disable-target-call-checks -fcuda-target-overloads -verify %s 13 // RUN: %clang_cc1 -DNOCHECKS -triple nvptx64-nvidia-cuda -fsyntax-only \ 14 // RUN: -fcuda-disable-target-call-checks -fcuda-target-overloads \ 15 // RUN: -fcuda-is-device -verify %s 16 17 #include "Inputs/cuda.h" 18 19 typedef int (*fp_t)(void); 20 typedef void (*gp_t)(void); 21 22 // Host and unattributed functions can't be overloaded 23 __host__ int hh(void) { return 1; } // expected-note {{previous definition is here}} 24 int hh(void) { return 1; } // expected-error {{redefinition of 'hh'}} 25 26 // H/D overloading is OK 27 __host__ int dh(void) { return 2; } 28 __device__ int dh(void) { return 2; } 29 30 // H/HD and D/HD are not allowed 31 __host__ __device__ int hdh(void) { return 5; } // expected-note {{previous definition is here}} 32 __host__ int hdh(void) { return 4; } // expected-error {{redefinition of 'hdh'}} 33 34 __host__ int hhd(void) { return 4; } // expected-note {{previous definition is here}} 35 __host__ __device__ int hhd(void) { return 5; } // expected-error {{redefinition of 'hhd'}} 36 // expected-warning@-1 {{attribute declaration must precede definition}} 37 // expected-note@-3 {{previous definition is here}} 38 39 __host__ __device__ int hdd(void) { return 7; } // expected-note {{previous definition is here}} 40 __device__ int hdd(void) { return 6; } // expected-error {{redefinition of 'hdd'}} 41 42 __device__ int dhd(void) { return 6; } // expected-note {{previous definition is here}} 43 __host__ __device__ int dhd(void) { return 7; } // expected-error {{redefinition of 'dhd'}} 44 // expected-warning@-1 {{attribute declaration must precede definition}} 45 // expected-note@-3 {{previous definition is here}} 46 47 // Same tests for extern "C" functions 48 extern "C" __host__ int chh(void) {return 11;} // expected-note {{previous definition is here}} 49 extern "C" int chh(void) {return 11;} // expected-error {{redefinition of 'chh'}} 50 51 // H/D overloading is OK 52 extern "C" __device__ int cdh(void) {return 10;} 53 extern "C" __host__ int cdh(void) {return 11;} 54 55 // H/HD and D/HD overloading is not allowed. 56 extern "C" __host__ __device__ int chhd1(void) {return 12;} // expected-note {{previous definition is here}} 57 extern "C" __host__ int chhd1(void) {return 13;} // expected-error {{redefinition of 'chhd1'}} 58 59 extern "C" __host__ int chhd2(void) {return 13;} // expected-note {{previous definition is here}} 60 extern "C" __host__ __device__ int chhd2(void) {return 12;} // expected-error {{redefinition of 'chhd2'}} 61 // expected-warning@-1 {{attribute declaration must precede definition}} 62 // expected-note@-3 {{previous definition is here}} 63 64 // Helper functions to verify calling restrictions. 65 __device__ int d(void) { return 8; } 66 __host__ int h(void) { return 9; } 67 __global__ void g(void) {} 68 extern "C" __device__ int cd(void) {return 10;} 69 extern "C" __host__ int ch(void) {return 11;} 70 71 __host__ void hostf(void) { 72 fp_t dp = d; 73 fp_t cdp = cd; 74 #if !defined(NOCHECKS) 75 // expected-error@-3 {{reference to __device__ function 'd' in __host__ function}} 76 // expected-note@65 {{'d' declared here}} 77 // expected-error@-4 {{reference to __device__ function 'cd' in __host__ function}} 78 // expected-note@68 {{'cd' declared here}} 79 #endif 80 fp_t hp = h; 81 fp_t chp = ch; 82 fp_t dhp = dh; 83 fp_t cdhp = cdh; 84 gp_t gp = g; 85 86 d(); 87 cd(); 88 #if !defined(NOCHECKS) 89 // expected-error@-3 {{no matching function for call to 'd'}} 90 // expected-note@65 {{candidate function not viable: call to __device__ function from __host__ function}} 91 // expected-error@-4 {{no matching function for call to 'cd'}} 92 // expected-note@68 {{candidate function not viable: call to __device__ function from __host__ function}} 93 #endif 94 h(); 95 ch(); 96 dh(); 97 cdh(); 98 g(); // expected-error {{call to global function g not configured}} 99 g<<<0,0>>>(); 100 } 101 102 103 __device__ void devicef(void) { 104 fp_t dp = d; 105 fp_t cdp = cd; 106 fp_t hp = h; 107 fp_t chp = ch; 108 #if !defined(NOCHECKS) 109 // expected-error@-3 {{reference to __host__ function 'h' in __device__ function}} 110 // expected-note@66 {{'h' declared here}} 111 // expected-error@-4 {{reference to __host__ function 'ch' in __device__ function}} 112 // expected-note@69 {{'ch' declared here}} 113 #endif 114 fp_t dhp = dh; 115 fp_t cdhp = cdh; 116 gp_t gp = g; // expected-error {{reference to __global__ function 'g' in __device__ function}} 117 // expected-note@67 {{'g' declared here}} 118 119 d(); 120 cd(); 121 h(); 122 ch(); 123 #if !defined(NOCHECKS) 124 // expected-error@-3 {{no matching function for call to 'h'}} 125 // expected-note@66 {{candidate function not viable: call to __host__ function from __device__ function}} 126 // expected-error@-4 {{no matching function for call to 'ch'}} 127 // expected-note@69 {{candidate function not viable: call to __host__ function from __device__ function}} 128 #endif 129 dh(); 130 cdh(); 131 g(); // expected-error {{no matching function for call to 'g'}} 132 // expected-note@67 {{candidate function not viable: call to __global__ function from __device__ function}} 133 g<<<0,0>>>(); // expected-error {{reference to __global__ function 'g' in __device__ function}} 134 // expected-note@67 {{'g' declared here}} 135 } 136 137 __global__ void globalf(void) { 138 fp_t dp = d; 139 fp_t cdp = cd; 140 fp_t hp = h; 141 fp_t chp = ch; 142 #if !defined(NOCHECKS) 143 // expected-error@-3 {{reference to __host__ function 'h' in __global__ function}} 144 // expected-note@66 {{'h' declared here}} 145 // expected-error@-4 {{reference to __host__ function 'ch' in __global__ function}} 146 // expected-note@69 {{'ch' declared here}} 147 #endif 148 fp_t dhp = dh; 149 fp_t cdhp = cdh; 150 gp_t gp = g; // expected-error {{reference to __global__ function 'g' in __global__ function}} 151 // expected-note@67 {{'g' declared here}} 152 153 d(); 154 cd(); 155 h(); 156 ch(); 157 #if !defined(NOCHECKS) 158 // expected-error@-3 {{no matching function for call to 'h'}} 159 // expected-note@66 {{candidate function not viable: call to __host__ function from __global__ function}} 160 // expected-error@-4 {{no matching function for call to 'ch'}} 161 // expected-note@69 {{candidate function not viable: call to __host__ function from __global__ function}} 162 #endif 163 dh(); 164 cdh(); 165 g(); // expected-error {{no matching function for call to 'g'}} 166 // expected-note@67 {{candidate function not viable: call to __global__ function from __global__ function}} 167 g<<<0,0>>>(); // expected-error {{reference to __global__ function 'g' in __global__ function}} 168 // expected-note@67 {{'g' declared here}} 169 } 170 171 __host__ __device__ void hostdevicef(void) { 172 fp_t dp = d; 173 fp_t cdp = cd; 174 fp_t hp = h; 175 fp_t chp = ch; 176 #if !defined(NOCHECKS) 177 #if !defined(__CUDA_ARCH__) 178 // expected-error@-6 {{reference to __device__ function 'd' in __host__ __device__ function}} 179 // expected-note@65 {{'d' declared here}} 180 // expected-error@-7 {{reference to __device__ function 'cd' in __host__ __device__ function}} 181 // expected-note@68 {{'cd' declared here}} 182 #else 183 // expected-error@-9 {{reference to __host__ function 'h' in __host__ __device__ function}} 184 // expected-note@66 {{'h' declared here}} 185 // expected-error@-10 {{reference to __host__ function 'ch' in __host__ __device__ function}} 186 // expected-note@69 {{'ch' declared here}} 187 #endif 188 #endif 189 fp_t dhp = dh; 190 fp_t cdhp = cdh; 191 gp_t gp = g; 192 #if defined(__CUDA_ARCH__) 193 // expected-error@-2 {{reference to __global__ function 'g' in __host__ __device__ function}} 194 // expected-note@67 {{'g' declared here}} 195 #endif 196 197 d(); 198 cd(); 199 h(); 200 ch(); 201 #if !defined(NOCHECKS) 202 #if !defined(__CUDA_ARCH__) 203 // expected-error@-6 {{no matching function for call to 'd'}} 204 // expected-note@65 {{candidate function not viable: call to __device__ function from __host__ __device__ function}} 205 // expected-error@-7 {{no matching function for call to 'cd'}} 206 // expected-note@68 {{candidate function not viable: call to __device__ function from __host__ __device__ function}} 207 #else 208 // expected-error@-9 {{no matching function for call to 'h'}} 209 // expected-note@66 {{candidate function not viable: call to __host__ function from __host__ __device__ function}} 210 // expected-error@-10 {{no matching function for call to 'ch'}} 211 // expected-note@69 {{candidate function not viable: call to __host__ function from __host__ __device__ function}} 212 #endif 213 #endif 214 215 dh(); 216 cdh(); 217 g(); 218 g<<<0,0>>>(); 219 #if !defined(__CUDA_ARCH__) 220 // expected-error@-3 {{call to global function g not configured}} 221 #else 222 // expected-error@-5 {{no matching function for call to 'g'}} 223 // expected-note@67 {{candidate function not viable: call to __global__ function from __host__ __device__ function}} 224 // expected-error@-6 {{reference to __global__ function 'g' in __host__ __device__ function}} 225 // expected-note@67 {{'g' declared here}} 226 #endif // __CUDA_ARCH__ 227 } 228 229 // Test for address of overloaded function resolution in the global context. 230 fp_t hp = h; 231 fp_t chp = ch; 232 fp_t dhp = dh; 233 fp_t cdhp = cdh; 234 gp_t gp = g; 235 236 237 // Test overloading of destructors 238 // Can't mix H and unattributed destructors 239 struct d_h { 240 ~d_h() {} // expected-note {{previous declaration is here}} 241 __host__ ~d_h() {} // expected-error {{destructor cannot be redeclared}} 242 }; 243 244 // H/D overloading is OK 245 struct d_dh { 246 __device__ ~d_dh() {} 247 __host__ ~d_dh() {} 248 }; 249 250 // HD is OK 251 struct d_hd { 252 __host__ __device__ ~d_hd() {} 253 }; 254 255 // Mixing H/D and HD is not allowed. 256 struct d_dhhd { 257 __device__ ~d_dhhd() {} 258 __host__ ~d_dhhd() {} // expected-note {{previous declaration is here}} 259 __host__ __device__ ~d_dhhd() {} // expected-error {{destructor cannot be redeclared}} 260 }; 261 262 struct d_hhd { 263 __host__ ~d_hhd() {} // expected-note {{previous declaration is here}} 264 __host__ __device__ ~d_hhd() {} // expected-error {{destructor cannot be redeclared}} 265 }; 266 267 struct d_hdh { 268 __host__ __device__ ~d_hdh() {} // expected-note {{previous declaration is here}} 269 __host__ ~d_hdh() {} // expected-error {{destructor cannot be redeclared}} 270 }; 271 272 struct d_dhd { 273 __device__ ~d_dhd() {} // expected-note {{previous declaration is here}} 274 __host__ __device__ ~d_dhd() {} // expected-error {{destructor cannot be redeclared}} 275 }; 276 277 struct d_hdd { 278 __host__ __device__ ~d_hdd() {} // expected-note {{previous declaration is here}} 279 __device__ ~d_hdd() {} // expected-error {{destructor cannot be redeclared}} 280 }; 281 282 // Test overloading of member functions 283 struct m_h { 284 void operator delete(void *ptr); // expected-note {{previous declaration is here}} 285 __host__ void operator delete(void *ptr); // expected-error {{class member cannot be redeclared}} 286 }; 287 288 // D/H overloading is OK 289 struct m_dh { 290 __device__ void operator delete(void *ptr); 291 __host__ void operator delete(void *ptr); 292 }; 293 294 // HD by itself is OK 295 struct m_hd { 296 __device__ __host__ void operator delete(void *ptr); 297 }; 298 299 struct m_hhd { 300 __host__ void operator delete(void *ptr) {} // expected-note {{previous declaration is here}} 301 __host__ __device__ void operator delete(void *ptr) {} // expected-error {{class member cannot be redeclared}} 302 }; 303 304 struct m_hdh { 305 __host__ __device__ void operator delete(void *ptr) {} // expected-note {{previous declaration is here}} 306 __host__ void operator delete(void *ptr) {} // expected-error {{class member cannot be redeclared}} 307 }; 308 309 struct m_dhd { 310 __device__ void operator delete(void *ptr) {} // expected-note {{previous declaration is here}} 311 __host__ __device__ void operator delete(void *ptr) {} // expected-error {{class member cannot be redeclared}} 312 }; 313 314 struct m_hdd { 315 __host__ __device__ void operator delete(void *ptr) {} // expected-note {{previous declaration is here}} 316 __device__ void operator delete(void *ptr) {} // expected-error {{class member cannot be redeclared}} 317 }; 318
