497 lines · plain
1// REQUIRES: nvptx-registered-target2 3// Make sure we don't allow dynamic initialization for device4// 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 %s7 8#ifdef __clang__9#include "Inputs/cuda.h"10#endif11 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 device214// side and that they conform to the initialization constraints.215// __shared__ can't be initialized at all and others don't support dynamic initialization.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 local239 __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 434__host__ __device__ void hd_sema() {435 static int x = 42;436}437 438inline __host__ __device__ void hd_emitted_host_only() {439 static int x = 42; // no error on device because this is never codegen'ed there.440}441void call_hd_emitted_host_only() { hd_emitted_host_only(); }442 443// Verify that we also check field initializers in instantiated structs.444struct NontrivialInitializer {445 __host__ __device__ NontrivialInitializer() : x(43) {}446 int x;447};448 449template <typename T>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.467template <>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 484void instantiate() {485 bar<NontrivialInitializer><<<1, 1>>>();486// expected-note@-1 {{in instantiation of function template specialization 'bar<NontrivialInitializer>' requested here}}487}488 489__device__ void *ptr1 = nullptr;490__device__ void *ptr2 = ptr1;491// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables}}492 493__device__ [[gnu::constructor(101)]] void ctor() {}494// expected-error@-1 {{CUDA does not support global constructors for __device__ functions}}495__device__ [[gnu::destructor(101)]] void dtor() {}496// expected-error@-1 {{CUDA does not support global destructors for __device__ functions}}497