1 // REQUIRES: nvptx-registered-target
2
3 // Make sure we don't allow dynamic initialization for device
4 // variables, but accept empty constructors allowed by CUDA.
5
6 // RUN: %clang_cc1 -verify %s -triple nvptx64-nvidia-cuda -fcuda-is-device -std=c++11 %s
7
8 #ifdef __clang__
9 #include "Inputs/cuda.h"
10 #endif
11
12 // Use the types we share with CodeGen tests.
13 #include "Inputs/cuda-initializers.h"
14
15 __shared__ int s_v_i = 1;
16 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
17
18 __device__ int d_v_f = f();
19 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
20 __shared__ int s_v_f = f();
21 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
22 __constant__ int c_v_f = f();
23 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
24
25 __shared__ T s_t_i = {2};
26 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
27 __device__ T d_t_i = {2};
28 __constant__ T c_t_i = {2};
29
30 __device__ ECD d_ecd_i{};
31 __shared__ ECD s_ecd_i{};
32 __constant__ ECD c_ecd_i{};
33
34 __device__ EC d_ec_i(3);
35 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
36 __shared__ EC s_ec_i(3);
37 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
38 __constant__ EC c_ec_i(3);
39 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
40
41 __device__ EC d_ec_i2 = {3};
42 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
43 __shared__ EC s_ec_i2 = {3};
44 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
45 __constant__ EC c_ec_i2 = {3};
46 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
47
48 __device__ ETC d_etc_i(3);
49 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
50 __shared__ ETC s_etc_i(3);
51 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
52 __constant__ ETC c_etc_i(3);
53 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
54
55 __device__ ETC d_etc_i2 = {3};
56 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
57 __shared__ ETC s_etc_i2 = {3};
58 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
59 __constant__ ETC c_etc_i2 = {3};
60 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
61
62 __device__ UC d_uc;
63 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
64 __shared__ UC s_uc;
65 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
66 __constant__ UC c_uc;
67 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
68
69 __device__ UD d_ud;
70 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
71 __shared__ UD s_ud;
72 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
73 __constant__ UD c_ud;
74 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
75
76 __device__ ECI d_eci;
77 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
78 __shared__ ECI s_eci;
79 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
80 __constant__ ECI c_eci;
81 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
82
83 __device__ NEC d_nec;
84 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
85 __shared__ NEC s_nec;
86 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
87 __constant__ NEC c_nec;
88 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
89
90 __device__ NED d_ned;
91 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
92 __shared__ NED s_ned;
93 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
94 __constant__ NED c_ned;
95 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
96
97 __device__ NCV d_ncv;
98 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
99 __shared__ NCV s_ncv;
100 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
101 __constant__ NCV c_ncv;
102 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
103
104 __device__ VD d_vd;
105 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
106 __shared__ VD s_vd;
107 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
108 __constant__ VD c_vd;
109 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
110
111 __device__ NCF d_ncf;
112 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
113 __shared__ NCF s_ncf;
114 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
115 __constant__ NCF c_ncf;
116 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
117
118 __shared__ NCFS s_ncfs;
119 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
120
121 __device__ UTC d_utc;
122 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
123 __shared__ UTC s_utc;
124 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
125 __constant__ UTC c_utc;
126 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
127
128 __device__ UTC d_utc_i(3);
129 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
130 __shared__ UTC s_utc_i(3);
131 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
132 __constant__ UTC c_utc_i(3);
133 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
134
135 __device__ NETC d_netc;
136 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
137 __shared__ NETC s_netc;
138 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
139 __constant__ NETC c_netc;
140 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
141
142 __device__ NETC d_netc_i(3);
143 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
144 __shared__ NETC s_netc_i(3);
145 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
146 __constant__ NETC c_netc_i(3);
147 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
148
149 __device__ EC_I_EC1 d_ec_i_ec1;
150 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
151 __shared__ EC_I_EC1 s_ec_i_ec1;
152 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
153 __constant__ EC_I_EC1 c_ec_i_ec1;
154 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
155
156 __device__ T_V_T d_t_v_t;
157 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
158 __shared__ T_V_T s_t_v_t;
159 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
160 __constant__ T_V_T c_t_v_t;
161 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
162
163 __device__ T_B_NEC d_t_b_nec;
164 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
165 __shared__ T_B_NEC s_t_b_nec;
166 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
167 __constant__ T_B_NEC c_t_b_nec;
168 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
169
170 __device__ T_F_NEC d_t_f_nec;
171 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
172 __shared__ T_F_NEC s_t_f_nec;
173 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
174 __constant__ T_F_NEC c_t_f_nec;
175 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
176
177 __device__ T_FA_NEC d_t_fa_nec;
178 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
179 __shared__ T_FA_NEC s_t_fa_nec;
180 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
181 __constant__ T_FA_NEC c_t_fa_nec;
182 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
183
184 __device__ T_B_NED d_t_b_ned;
185 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
186 __shared__ T_B_NED s_t_b_ned;
187 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
188 __constant__ T_B_NED c_t_b_ned;
189 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
190
191 __device__ T_F_NED d_t_f_ned;
192 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
193 __shared__ T_F_NED s_t_f_ned;
194 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
195 __constant__ T_F_NED c_t_f_ned;
196 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
197
198 __device__ T_FA_NED d_t_fa_ned;
199 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
200 __shared__ T_FA_NED s_t_fa_ned;
201 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
202 __constant__ T_FA_NED c_t_fa_ned;
203 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
204
205 // Verify that local variables may be static on device
206 // side and that they conform to the initialization constraints.
207 // __shared__ can't be initialized at all and others don't support dynamic initialization.
df_sema()208 __device__ void df_sema() {
209 static __device__ int ds;
210 static __constant__ int dc;
211 static int v;
212 static const int cv = 1;
213 static const __device__ int cds = 1;
214 static const __constant__ int cdc = 1;
215
216
217 // __shared__ does not need to be explicitly static.
218 __shared__ int lsi;
219 // __constant__ and __device__ can not be non-static local
220 __constant__ int lci;
221 // expected-error@-1 {{__constant__ and __device__ are not allowed on non-static local variables}}
222 __device__ int ldi;
223 // expected-error@-1 {{__constant__ and __device__ are not allowed on non-static local variables}}
224
225 // Same test cases as for the globals above.
226
227 static __device__ int d_v_f = f();
228 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
229 static __shared__ int s_v_f = f();
230 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
231 static __constant__ int c_v_f = f();
232 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
233
234 static __shared__ T s_t_i = {2};
235 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
236 static __device__ T d_t_i = {2};
237 static __constant__ T c_t_i = {2};
238
239 static __device__ ECD d_ecd_i;
240 static __shared__ ECD s_ecd_i;
241 static __constant__ ECD c_ecd_i;
242
243 static __device__ EC d_ec_i(3);
244 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
245 static __shared__ EC s_ec_i(3);
246 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
247 static __constant__ EC c_ec_i(3);
248 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
249
250 static __device__ EC d_ec_i2 = {3};
251 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
252 static __shared__ EC s_ec_i2 = {3};
253 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
254 static __constant__ EC c_ec_i2 = {3};
255 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
256
257 static __device__ ETC d_etc_i(3);
258 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
259 static __shared__ ETC s_etc_i(3);
260 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
261 static __constant__ ETC c_etc_i(3);
262 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
263
264 static __device__ ETC d_etc_i2 = {3};
265 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
266 static __shared__ ETC s_etc_i2 = {3};
267 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
268 static __constant__ ETC c_etc_i2 = {3};
269 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
270
271 static __device__ UC d_uc;
272 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
273 static __shared__ UC s_uc;
274 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
275 static __constant__ UC c_uc;
276 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
277
278 static __device__ UD d_ud;
279 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
280 static __shared__ UD s_ud;
281 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
282 static __constant__ UD c_ud;
283 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
284
285 static __device__ ECI d_eci;
286 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
287 static __shared__ ECI s_eci;
288 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
289 static __constant__ ECI c_eci;
290 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
291
292 static __device__ NEC d_nec;
293 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
294 static __shared__ NEC s_nec;
295 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
296 static __constant__ NEC c_nec;
297 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
298
299 static __device__ NED d_ned;
300 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
301 static __shared__ NED s_ned;
302 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
303 static __constant__ NED c_ned;
304 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
305
306 static __device__ NCV d_ncv;
307 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
308 static __shared__ NCV s_ncv;
309 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
310 static __constant__ NCV c_ncv;
311 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
312
313 static __device__ VD d_vd;
314 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
315 static __shared__ VD s_vd;
316 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
317 static __constant__ VD c_vd;
318 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
319
320 static __device__ NCF d_ncf;
321 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
322 static __shared__ NCF s_ncf;
323 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
324 static __constant__ NCF c_ncf;
325 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
326
327 static __shared__ NCFS s_ncfs;
328 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
329
330 static __device__ UTC d_utc;
331 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
332 static __shared__ UTC s_utc;
333 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
334 static __constant__ UTC c_utc;
335 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
336
337 static __device__ UTC d_utc_i(3);
338 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
339 static __shared__ UTC s_utc_i(3);
340 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
341 static __constant__ UTC c_utc_i(3);
342 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
343
344 static __device__ NETC d_netc;
345 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
346 static __shared__ NETC s_netc;
347 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
348 static __constant__ NETC c_netc;
349 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
350
351 static __device__ NETC d_netc_i(3);
352 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
353 static __shared__ NETC s_netc_i(3);
354 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
355 static __constant__ NETC c_netc_i(3);
356 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
357
358 static __device__ EC_I_EC1 d_ec_i_ec1;
359 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
360 static __shared__ EC_I_EC1 s_ec_i_ec1;
361 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
362 static __constant__ EC_I_EC1 c_ec_i_ec1;
363 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
364
365 static __device__ T_V_T d_t_v_t;
366 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
367 static __shared__ T_V_T s_t_v_t;
368 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
369 static __constant__ T_V_T c_t_v_t;
370 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
371
372 static __device__ T_B_NEC d_t_b_nec;
373 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
374 static __shared__ T_B_NEC s_t_b_nec;
375 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
376 static __constant__ T_B_NEC c_t_b_nec;
377 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
378
379 static __device__ T_F_NEC d_t_f_nec;
380 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
381 static __shared__ T_F_NEC s_t_f_nec;
382 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
383 static __constant__ T_F_NEC c_t_f_nec;
384 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
385
386 static __device__ T_FA_NEC d_t_fa_nec;
387 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
388 static __shared__ T_FA_NEC s_t_fa_nec;
389 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
390 static __constant__ T_FA_NEC c_t_fa_nec;
391 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
392
393 static __device__ T_B_NED d_t_b_ned;
394 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
395 static __shared__ T_B_NED s_t_b_ned;
396 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
397 static __constant__ T_B_NED c_t_b_ned;
398 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
399
400 static __device__ T_F_NED d_t_f_ned;
401 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
402 static __shared__ T_F_NED s_t_f_ned;
403 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
404 static __constant__ T_F_NED c_t_f_ned;
405 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
406
407 static __device__ T_FA_NED d_t_fa_ned;
408 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
409 static __shared__ T_FA_NED s_t_fa_ned;
410 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
411 static __constant__ T_FA_NED c_t_fa_ned;
412 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
413 }
414
hd_sema()415 __host__ __device__ void hd_sema() {
416 static int x = 42;
417 }
418
hd_emitted_host_only()419 inline __host__ __device__ void hd_emitted_host_only() {
420 static int x = 42; // no error on device because this is never codegen'ed there.
421 }
call_hd_emitted_host_only()422 void call_hd_emitted_host_only() { hd_emitted_host_only(); }
423
424 // Verify that we also check field initializers in instantiated structs.
425 struct NontrivialInitializer {
NontrivialInitializerNontrivialInitializer426 __host__ __device__ NontrivialInitializer() : x(43) {}
427 int x;
428 };
429
430 template <typename T>
bar()431 __global__ void bar() {
432 __shared__ T bad;
433 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
434 }
435
instantiate()436 void instantiate() {
437 bar<NontrivialInitializer><<<1, 1>>>();
438 // expected-note@-1 {{in instantiation of function template specialization 'bar<NontrivialInitializer>' requested here}}
439 }
440