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
hh(void)23 __host__ int hh(void) { return 1; } // expected-note {{previous definition is here}}
hh(void)24 int hh(void) { return 1; } // expected-error {{redefinition of 'hh'}}
25
26 // H/D overloading is OK
dh(void)27 __host__ int dh(void) { return 2; }
dh(void)28 __device__ int dh(void) { return 2; }
29
30 // H/HD and D/HD are not allowed
hdh(void)31 __host__ __device__ int hdh(void) { return 5; } // expected-note {{previous definition is here}}
hdh(void)32 __host__ int hdh(void) { return 4; } // expected-error {{redefinition of 'hdh'}}
33
hhd(void)34 __host__ int hhd(void) { return 4; } // expected-note {{previous definition is here}}
hhd(void)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
hdd(void)39 __host__ __device__ int hdd(void) { return 7; } // expected-note {{previous definition is here}}
hdd(void)40 __device__ int hdd(void) { return 6; } // expected-error {{redefinition of 'hdd'}}
41
dhd(void)42 __device__ int dhd(void) { return 6; } // expected-note {{previous definition is here}}
dhd(void)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
chh(void)48 extern "C" __host__ int chh(void) {return 11;} // expected-note {{previous definition is here}}
chh(void)49 extern "C" int chh(void) {return 11;} // expected-error {{redefinition of 'chh'}}
50
51 // H/D overloading is OK
cdh(void)52 extern "C" __device__ int cdh(void) {return 10;}
cdh(void)53 extern "C" __host__ int cdh(void) {return 11;}
54
55 // H/HD and D/HD overloading is not allowed.
chhd1(void)56 extern "C" __host__ __device__ int chhd1(void) {return 12;} // expected-note {{previous definition is here}}
chhd1(void)57 extern "C" __host__ int chhd1(void) {return 13;} // expected-error {{redefinition of 'chhd1'}}
58
chhd2(void)59 extern "C" __host__ int chhd2(void) {return 13;} // expected-note {{previous definition is here}}
chhd2(void)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.
d(void)65 __device__ int d(void) { return 8; }
h(void)66 __host__ int h(void) { return 9; }
g(void)67 __global__ void g(void) {}
cd(void)68 extern "C" __device__ int cd(void) {return 10;}
ch(void)69 extern "C" __host__ int ch(void) {return 11;}
70
hostf(void)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
devicef(void)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
globalf(void)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
hostdevicef(void)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 {
~d_hd_h240 ~d_h() {} // expected-note {{previous declaration is here}}
~d_hd_h241 __host__ ~d_h() {} // expected-error {{destructor cannot be redeclared}}
242 };
243
244 // H/D overloading is OK
245 struct d_dh {
~d_dhd_dh246 __device__ ~d_dh() {}
~d_dhd_dh247 __host__ ~d_dh() {}
248 };
249
250 // HD is OK
251 struct d_hd {
~d_hdd_hd252 __host__ __device__ ~d_hd() {}
253 };
254
255 // Mixing H/D and HD is not allowed.
256 struct d_dhhd {
~d_dhhdd_dhhd257 __device__ ~d_dhhd() {}
~d_dhhdd_dhhd258 __host__ ~d_dhhd() {} // expected-note {{previous declaration is here}}
~d_dhhdd_dhhd259 __host__ __device__ ~d_dhhd() {} // expected-error {{destructor cannot be redeclared}}
260 };
261
262 struct d_hhd {
~d_hhdd_hhd263 __host__ ~d_hhd() {} // expected-note {{previous declaration is here}}
~d_hhdd_hhd264 __host__ __device__ ~d_hhd() {} // expected-error {{destructor cannot be redeclared}}
265 };
266
267 struct d_hdh {
~d_hdhd_hdh268 __host__ __device__ ~d_hdh() {} // expected-note {{previous declaration is here}}
~d_hdhd_hdh269 __host__ ~d_hdh() {} // expected-error {{destructor cannot be redeclared}}
270 };
271
272 struct d_dhd {
~d_dhdd_dhd273 __device__ ~d_dhd() {} // expected-note {{previous declaration is here}}
~d_dhdd_dhd274 __host__ __device__ ~d_dhd() {} // expected-error {{destructor cannot be redeclared}}
275 };
276
277 struct d_hdd {
~d_hddd_hdd278 __host__ __device__ ~d_hdd() {} // expected-note {{previous declaration is here}}
~d_hddd_hdd279 __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 {
operator deletem_hhd300 __host__ void operator delete(void *ptr) {} // expected-note {{previous declaration is here}}
operator deletem_hhd301 __host__ __device__ void operator delete(void *ptr) {} // expected-error {{class member cannot be redeclared}}
302 };
303
304 struct m_hdh {
operator deletem_hdh305 __host__ __device__ void operator delete(void *ptr) {} // expected-note {{previous declaration is here}}
operator deletem_hdh306 __host__ void operator delete(void *ptr) {} // expected-error {{class member cannot be redeclared}}
307 };
308
309 struct m_dhd {
operator deletem_dhd310 __device__ void operator delete(void *ptr) {} // expected-note {{previous declaration is here}}
operator deletem_dhd311 __host__ __device__ void operator delete(void *ptr) {} // expected-error {{class member cannot be redeclared}}
312 };
313
314 struct m_hdd {
operator deletem_hdd315 __host__ __device__ void operator delete(void *ptr) {} // expected-note {{previous declaration is here}}
operator deletem_hdd316 __device__ void operator delete(void *ptr) {} // expected-error {{class member cannot be redeclared}}
317 };
318