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 -triple nvptx64-nvidia-cuda -fcuda-is-device -std=c++11 \ 7 // RUN: -fno-threadsafe-statics -emit-llvm -o - %s | FileCheck %s 8 // RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fcuda-is-device -std=c++11 \ 9 // RUN: -emit-llvm -DERROR_CASE -verify -o /dev/null %s 10 11 #ifdef __clang__ 12 #include "Inputs/cuda.h" 13 #endif 14 15 // Base classes with different initializer variants. 16 17 // trivial constructor -- allowed 18 struct T { 19 int t; 20 }; 21 22 // empty constructor 23 struct EC { 24 int ec; 25 __device__ EC() {} // -- allowed 26 __device__ EC(int) {} // -- not allowed 27 }; 28 29 // empty templated constructor -- allowed with no arguments 30 struct ETC { 31 template <typename... T> __device__ ETC(T...) {} 32 }; 33 34 // undefined constructor -- not allowed 35 struct UC { 36 int uc; 37 __device__ UC(); 38 }; 39 40 // empty constructor w/ initializer list -- not allowed 41 struct ECI { 42 int eci; 43 __device__ ECI() : eci(1) {} 44 }; 45 46 // non-empty constructor -- not allowed 47 struct NEC { 48 int nec; 49 __device__ NEC() { nec = 1; } 50 }; 51 52 // no-constructor, virtual method -- not allowed 53 struct NCV { 54 int ncv; 55 __device__ virtual void vm() {} 56 }; 57 58 // dynamic in-class field initializer -- not allowed 59 __device__ int f(); 60 struct NCF { 61 int ncf = f(); 62 }; 63 64 // static in-class field initializer. NVCC does not allow it, but 65 // clang generates static initializer for this, so we'll accept it. 66 struct NCFS { 67 int ncfs = 3; 68 }; 69 70 // undefined templated constructor -- not allowed 71 struct UTC { 72 template <typename... T> __device__ UTC(T...); 73 }; 74 75 // non-empty templated constructor -- not allowed 76 struct NETC { 77 int netc; 78 template <typename... T> __device__ NETC(T...) { netc = 1; } 79 }; 80 81 __device__ int d_v; 82 // CHECK: @d_v = addrspace(1) externally_initialized global i32 0, 83 __shared__ int s_v; 84 // CHECK: @s_v = addrspace(3) global i32 undef, 85 __constant__ int c_v; 86 // CHECK: addrspace(4) externally_initialized global i32 0, 87 88 __device__ int d_v_i = 1; 89 // CHECK: @d_v_i = addrspace(1) externally_initialized global i32 1, 90 #ifdef ERROR_CASE 91 __shared__ int s_v_i = 1; 92 // expected-error@-1 {{initialization is not supported for __shared__ variables.}} 93 #endif 94 __constant__ int c_v_i = 1; 95 // CHECK: @c_v_i = addrspace(4) externally_initialized global i32 1, 96 97 #ifdef ERROR_CASE 98 __device__ int d_v_f = f(); 99 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}} 100 __shared__ int s_v_f = f(); 101 // expected-error@-1 {{initialization is not supported for __shared__ variables.}} 102 __constant__ int c_v_f = f(); 103 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}} 104 #endif 105 106 __device__ T d_t; 107 // CHECK: @d_t = addrspace(1) externally_initialized global %struct.T zeroinitializer 108 __shared__ T s_t; 109 // CHECK: @s_t = addrspace(3) global %struct.T undef, 110 __constant__ T c_t; 111 // CHECK: @c_t = addrspace(4) externally_initialized global %struct.T zeroinitializer, 112 113 __device__ T d_t_i = {2}; 114 // CHECKL @d_t_i = addrspace(1) externally_initialized global %struct.T { i32 2 }, 115 #ifdef ERROR_CASE 116 __shared__ T s_t_i = {2}; 117 // expected-error@-1 {{initialization is not supported for __shared__ variables.}} 118 #endif 119 __constant__ T c_t_i = {2}; 120 // CHECK: @c_t_i = addrspace(4) externally_initialized global %struct.T { i32 2 }, 121 122 __device__ EC d_ec; 123 // CHECK: @d_ec = addrspace(1) externally_initialized global %struct.EC zeroinitializer, 124 __shared__ EC s_ec; 125 // CHECK: @s_ec = addrspace(3) global %struct.EC undef, 126 __constant__ EC c_ec; 127 // CHECK: @c_ec = addrspace(4) externally_initialized global %struct.EC zeroinitializer, 128 129 #if ERROR_CASE 130 __device__ EC d_ec_i(3); 131 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}} 132 __shared__ EC s_ec_i(3); 133 // expected-error@-1 {{initialization is not supported for __shared__ variables.}} 134 __constant__ EC c_ec_i(3); 135 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}} 136 137 __device__ EC d_ec_i2 = {3}; 138 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}} 139 __shared__ EC s_ec_i2 = {3}; 140 // expected-error@-1 {{initialization is not supported for __shared__ variables.}} 141 __constant__ EC c_ec_i2 = {3}; 142 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}} 143 #endif 144 145 __device__ ETC d_etc; 146 // CHETCK: @d_etc = addrspace(1) externally_initialized global %struct.ETC zeroinitializer, 147 __shared__ ETC s_etc; 148 // CHETCK: @s_etc = addrspace(3) global %struct.ETC undef, 149 __constant__ ETC c_etc; 150 // CHETCK: @c_etc = addrspace(4) externally_initialized global %struct.ETC zeroinitializer, 151 152 #if ERROR_CASE 153 __device__ ETC d_etc_i(3); 154 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}} 155 __shared__ ETC s_etc_i(3); 156 // expected-error@-1 {{initialization is not supported for __shared__ variables.}} 157 __constant__ ETC c_etc_i(3); 158 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}} 159 160 __device__ ETC d_etc_i2 = {3}; 161 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}} 162 __shared__ ETC s_etc_i2 = {3}; 163 // expected-error@-1 {{initialization is not supported for __shared__ variables.}} 164 __constant__ ETC c_etc_i2 = {3}; 165 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}} 166 167 __device__ UC d_uc; 168 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}} 169 __shared__ UC s_uc; 170 // expected-error@-1 {{initialization is not supported for __shared__ variables.}} 171 __constant__ UC c_uc; 172 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}} 173 174 __device__ ECI d_eci; 175 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}} 176 __shared__ ECI s_eci; 177 // expected-error@-1 {{initialization is not supported for __shared__ variables.}} 178 __constant__ ECI c_eci; 179 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}} 180 181 __device__ NEC d_nec; 182 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}} 183 __shared__ NEC s_nec; 184 // expected-error@-1 {{initialization is not supported for __shared__ variables.}} 185 __constant__ NEC c_nec; 186 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}} 187 188 __device__ NCV d_ncv; 189 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}} 190 __shared__ NCV s_ncv; 191 // expected-error@-1 {{initialization is not supported for __shared__ variables.}} 192 __constant__ NCV c_ncv; 193 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}} 194 195 __device__ NCF d_ncf; 196 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}} 197 __shared__ NCF s_ncf; 198 // expected-error@-1 {{initialization is not supported for __shared__ variables.}} 199 __constant__ NCF c_ncf; 200 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}} 201 #endif 202 203 __device__ NCFS d_ncfs; 204 // CHECK: @d_ncfs = addrspace(1) externally_initialized global %struct.NCFS { i32 3 } 205 __constant__ NCFS c_ncfs; 206 // CHECK: @c_ncfs = addrspace(4) externally_initialized global %struct.NCFS { i32 3 } 207 208 #if ERROR_CASE 209 __shared__ NCFS s_ncfs; 210 // expected-error@-1 {{initialization is not supported for __shared__ variables.}} 211 212 __device__ UTC d_utc; 213 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}} 214 __shared__ UTC s_utc; 215 // expected-error@-1 {{initialization is not supported for __shared__ variables.}} 216 __constant__ UTC c_utc; 217 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}} 218 219 __device__ UTC d_utc_i(3); 220 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}} 221 __shared__ UTC s_utc_i(3); 222 // expected-error@-1 {{initialization is not supported for __shared__ variables.}} 223 __constant__ UTC c_utc_i(3); 224 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}} 225 226 __device__ NETC d_netc; 227 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}} 228 __shared__ NETC s_netc; 229 // expected-error@-1 {{initialization is not supported for __shared__ variables.}} 230 __constant__ NETC c_netc; 231 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}} 232 233 __device__ NETC d_netc_i(3); 234 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}} 235 __shared__ NETC s_netc_i(3); 236 // expected-error@-1 {{initialization is not supported for __shared__ variables.}} 237 __constant__ NETC c_netc_i(3); 238 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}} 239 #endif 240 241 // Regular base class -- allowed 242 struct T_B_T : T {}; 243 __device__ T_B_T d_t_b_t; 244 // CHECK: @d_t_b_t = addrspace(1) externally_initialized global %struct.T_B_T zeroinitializer, 245 __shared__ T_B_T s_t_b_t; 246 // CHECK: @s_t_b_t = addrspace(3) global %struct.T_B_T undef, 247 __constant__ T_B_T c_t_b_t; 248 // CHECK: @c_t_b_t = addrspace(4) externally_initialized global %struct.T_B_T zeroinitializer, 249 250 // Incapsulated object of allowed class -- allowed 251 struct T_F_T { 252 T t; 253 }; 254 __device__ T_F_T d_t_f_t; 255 // CHECK: @d_t_f_t = addrspace(1) externally_initialized global %struct.T_F_T zeroinitializer, 256 __shared__ T_F_T s_t_f_t; 257 // CHECK: @s_t_f_t = addrspace(3) global %struct.T_F_T undef, 258 __constant__ T_F_T c_t_f_t; 259 // CHECK: @c_t_f_t = addrspace(4) externally_initialized global %struct.T_F_T zeroinitializer, 260 261 // array of allowed objects -- allowed 262 struct T_FA_T { 263 T t[2]; 264 }; 265 __device__ T_FA_T d_t_fa_t; 266 // CHECK: @d_t_fa_t = addrspace(1) externally_initialized global %struct.T_FA_T zeroinitializer, 267 __shared__ T_FA_T s_t_fa_t; 268 // CHECK: @s_t_fa_t = addrspace(3) global %struct.T_FA_T undef, 269 __constant__ T_FA_T c_t_fa_t; 270 // CHECK: @c_t_fa_t = addrspace(4) externally_initialized global %struct.T_FA_T zeroinitializer, 271 272 273 // Calling empty base class initializer is OK 274 struct EC_I_EC : EC { 275 __device__ EC_I_EC() : EC() {} 276 }; 277 __device__ EC_I_EC d_ec_i_ec; 278 // CHECK: @d_ec_i_ec = addrspace(1) externally_initialized global %struct.EC_I_EC zeroinitializer, 279 __shared__ EC_I_EC s_ec_i_ec; 280 // CHECK: @s_ec_i_ec = addrspace(3) global %struct.EC_I_EC undef, 281 __constant__ EC_I_EC c_ec_i_ec; 282 // CHECK: @c_ec_i_ec = addrspace(4) externally_initialized global %struct.EC_I_EC zeroinitializer, 283 284 // .. though passing arguments is not allowed. 285 struct EC_I_EC1 : EC { 286 __device__ EC_I_EC1() : EC(1) {} 287 }; 288 #if ERROR_CASE 289 __device__ EC_I_EC1 d_ec_i_ec1; 290 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}} 291 __shared__ EC_I_EC1 s_ec_i_ec1; 292 // expected-error@-1 {{initialization is not supported for __shared__ variables.}} 293 __constant__ EC_I_EC1 c_ec_i_ec1; 294 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}} 295 #endif 296 297 // Virtual base class -- not allowed 298 struct T_V_T : virtual T {}; 299 #if ERROR_CASE 300 __device__ T_V_T d_t_v_t; 301 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}} 302 __shared__ T_V_T s_t_v_t; 303 // expected-error@-1 {{initialization is not supported for __shared__ variables.}} 304 __constant__ T_V_T c_t_v_t; 305 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}} 306 #endif 307 308 // Make sure that we don't allow if we inherit or incapsulate 309 // something with disallowed initializer. 310 311 // Inherited from or incapsulated class with non-empty constructor -- 312 // not allowed 313 struct T_B_NEC : NEC {}; 314 struct T_F_NEC { 315 NEC nec; 316 }; 317 struct T_FA_NEC { 318 NEC nec[2]; 319 }; 320 321 #if ERROR_CASE 322 __device__ T_B_NEC d_t_b_nec; 323 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}} 324 __shared__ T_B_NEC s_t_b_nec; 325 // expected-error@-1 {{initialization is not supported for __shared__ variables.}} 326 __constant__ T_B_NEC c_t_b_nec; 327 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}} 328 329 __device__ T_F_NEC d_t_f_nec; 330 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}} 331 __shared__ T_F_NEC s_t_f_nec; 332 // expected-error@-1 {{initialization is not supported for __shared__ variables.}} 333 __constant__ T_F_NEC c_t_f_nec; 334 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}} 335 336 __device__ T_FA_NEC d_t_fa_nec; 337 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}} 338 __shared__ T_FA_NEC s_t_fa_nec; 339 // expected-error@-1 {{initialization is not supported for __shared__ variables.}} 340 __constant__ T_FA_NEC c_t_fa_nec; 341 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}} 342 #endif 343 344 // We should not emit global initializers for device-side variables. 345 // CHECK-NOT: @__cxx_global_var_init 346 347 // Make sure that initialization restrictions do not apply to local 348 // variables. 349 __device__ void df() { 350 T t; 351 EC ec; 352 ETC etc; 353 UC uc; 354 ECI eci; 355 NEC nec; 356 NCV ncv; 357 NCF ncf; 358 NCFS ncfs; 359 UTC utc; 360 NETC netc; 361 T_B_T t_b_t; 362 T_F_T t_f_t; 363 T_FA_T t_fa_t; 364 EC_I_EC ec_i_ec; 365 EC_I_EC1 ec_i_ec1; 366 T_V_T t_v_t; 367 T_B_NEC t_b_nec; 368 T_F_NEC t_f_nec; 369 T_FA_NEC t_fa_nec; 370 static __shared__ UC s_uc; 371 } 372 373 // CHECK: call void @_ZN2ECC1Ev(%struct.EC* %ec) 374 // CHECK: call void @_ZN3ETCC1IJEEEDpT_(%struct.ETC* %etc) 375 // CHECK: call void @_ZN2UCC1Ev(%struct.UC* %uc) 376 // CHECK: call void @_ZN3ECIC1Ev(%struct.ECI* %eci) 377 // CHECK: call void @_ZN3NECC1Ev(%struct.NEC* %nec) 378 // CHECK: call void @_ZN3NCVC1Ev(%struct.NCV* %ncv) 379 // CHECK: call void @_ZN3NCFC1Ev(%struct.NCF* %ncf) 380 // CHECK: call void @_ZN4NCFSC1Ev(%struct.NCFS* %ncfs) 381 // CHECK: call void @_ZN3UTCC1IJEEEDpT_(%struct.UTC* %utc) 382 // CHECK: call void @_ZN4NETCC1IJEEEDpT_(%struct.NETC* %netc) 383 // CHECK: call void @_ZN7EC_I_ECC1Ev(%struct.EC_I_EC* %ec_i_ec) 384 // CHECK: call void @_ZN8EC_I_EC1C1Ev(%struct.EC_I_EC1* %ec_i_ec1) 385 // CHECK: call void @_ZN5T_V_TC1Ev(%struct.T_V_T* %t_v_t) 386 // CHECK: call void @_ZN7T_B_NECC1Ev(%struct.T_B_NEC* %t_b_nec) 387 // CHECK: call void @_ZN7T_F_NECC1Ev(%struct.T_F_NEC* %t_f_nec) 388 // CHECK: call void @_ZN8T_FA_NECC1Ev(%struct.T_FA_NEC* %t_fa_nec) 389 // CHECK: call void @_ZN2UCC1Ev(%struct.UC* addrspacecast (%struct.UC addrspace(3)* @_ZZ2dfvE4s_uc to %struct.UC*)) 390 // CHECK: ret void 391 392 // We should not emit global init function. 393 // CHECK-NOT: @_GLOBAL__sub_I 394