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