clang/test/SemaCUDA/device-var-init.cu

   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.
 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
 434 __host__ __device__ void hd_sema() {
 435   static int x = 42;
 436 }
 437
 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 }
 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 {
 445   __host__ __device__ NontrivialInitializer() : x(43) {}
 446   int x;
 447 };
 448
 449 template <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.
 467 template <>
 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
 484 void instantiate() {
 485   bar<NontrivialInitializer><<<1, 1>>>();
 486 // expected-note@-1 {{in instantiation of function template specialization 'bar<NontrivialInitializer>' requested here}}
 487 }