1 // REQUIRES: nvptx-registered-target 2 // REQUIRES: amdgpu-registered-target 3 4 // Make sure we don't allow dynamic initialization for device 5 // variables, but accept empty constructors allowed by CUDA. 6 7 // RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fcuda-is-device -std=c++11 \ 8 // RUN: -fno-threadsafe-statics -emit-llvm -o - %s | FileCheck -check-prefixes=DEVICE,NVPTX %s 9 // RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -std=c++11 \ 10 // RUN: -fno-threadsafe-statics -emit-llvm -o - %s | FileCheck -check-prefixes=HOST %s 11 12 // RUN: %clang_cc1 -triple amdgcn -fcuda-is-device -std=c++11 \ 13 // RUN: -fno-threadsafe-statics -emit-llvm -o - %s | FileCheck -check-prefixes=DEVICE,AMDGCN %s 14 15 #ifdef __clang__ 16 #include "Inputs/cuda.h" 17 #endif 18 19 // Use the types we share with Sema tests. 20 #include "Inputs/cuda-initializers.h" 21 22 __device__ int d_v; 23 // DEVICE: @d_v = addrspace(1) externally_initialized global i32 0, 24 // HOST: @d_v = internal global i32 undef, 25 __shared__ int s_v; 26 // DEVICE: @s_v = addrspace(3) global i32 undef, 27 // HOST: @s_v = internal global i32 undef, 28 __constant__ int c_v; 29 // DEVICE: addrspace(4) externally_initialized global i32 0, 30 // HOST: @c_v = internal global i32 undef, 31 32 __device__ int d_v_i = 1; 33 // DEVICE: @d_v_i = addrspace(1) externally_initialized global i32 1, 34 // HOST: @d_v_i = internal global i32 undef, 35 36 // trivial constructor -- allowed 37 __device__ T d_t; 38 // DEVICE: @d_t = addrspace(1) externally_initialized global %struct.T zeroinitializer 39 // HOST: @d_t = internal global %struct.T undef, 40 __shared__ T s_t; 41 // DEVICE: @s_t = addrspace(3) global %struct.T undef, 42 // HOST: @s_t = internal global %struct.T undef, 43 __constant__ T c_t; 44 // DEVICE: @c_t = addrspace(4) externally_initialized global %struct.T zeroinitializer, 45 // HOST: @c_t = internal global %struct.T undef, 46 47 __device__ T d_t_i = {2}; 48 // DEVICE: @d_t_i = addrspace(1) externally_initialized global %struct.T { i32 2 }, 49 // HOST: @d_t_i = internal global %struct.T undef, 50 __constant__ T c_t_i = {2}; 51 // DEVICE: @c_t_i = addrspace(4) externally_initialized global %struct.T { i32 2 }, 52 // HOST: @c_t_i = internal global %struct.T undef, 53 54 // empty constructor 55 __device__ EC d_ec; 56 // DEVICE: @d_ec = addrspace(1) externally_initialized global %struct.EC zeroinitializer, 57 // HOST: @d_ec = internal global %struct.EC undef, 58 __shared__ EC s_ec; 59 // DEVICE: @s_ec = addrspace(3) global %struct.EC undef, 60 // HOST: @s_ec = internal global %struct.EC undef, 61 __constant__ EC c_ec; 62 // DEVICE: @c_ec = addrspace(4) externally_initialized global %struct.EC zeroinitializer, 63 // HOST: @c_ec = internal global %struct.EC undef 64 65 // empty destructor 66 __device__ ED d_ed; 67 // DEVICE: @d_ed = addrspace(1) externally_initialized global %struct.ED zeroinitializer, 68 // HOST: @d_ed = internal global %struct.ED undef, 69 __shared__ ED s_ed; 70 // DEVICE: @s_ed = addrspace(3) global %struct.ED undef, 71 // HOST: @s_ed = internal global %struct.ED undef, 72 __constant__ ED c_ed; 73 // DEVICE: @c_ed = addrspace(4) externally_initialized global %struct.ED zeroinitializer, 74 // HOST: @c_ed = internal global %struct.ED undef, 75 76 __device__ ECD d_ecd; 77 // DEVICE: @d_ecd = addrspace(1) externally_initialized global %struct.ECD zeroinitializer, 78 // HOST: @d_ecd = internal global %struct.ECD undef, 79 __shared__ ECD s_ecd; 80 // DEVICE: @s_ecd = addrspace(3) global %struct.ECD undef, 81 // HOST: @s_ecd = internal global %struct.ECD undef, 82 __constant__ ECD c_ecd; 83 // DEVICE: @c_ecd = addrspace(4) externally_initialized global %struct.ECD zeroinitializer, 84 // HOST: @c_ecd = internal global %struct.ECD undef, 85 86 // empty templated constructor -- allowed with no arguments 87 __device__ ETC d_etc; 88 // DEVICE: @d_etc = addrspace(1) externally_initialized global %struct.ETC zeroinitializer, 89 // HOST: @d_etc = internal global %struct.ETC undef, 90 __shared__ ETC s_etc; 91 // DEVICE: @s_etc = addrspace(3) global %struct.ETC undef, 92 // HOST: @s_etc = internal global %struct.ETC undef, 93 __constant__ ETC c_etc; 94 // DEVICE: @c_etc = addrspace(4) externally_initialized global %struct.ETC zeroinitializer, 95 // HOST: @c_etc = internal global %struct.ETC undef, 96 97 __device__ NCFS d_ncfs; 98 // DEVICE: @d_ncfs = addrspace(1) externally_initialized global %struct.NCFS { i32 3 } 99 // HOST: @d_ncfs = internal global %struct.NCFS undef, 100 __constant__ NCFS c_ncfs; 101 // DEVICE: @c_ncfs = addrspace(4) externally_initialized global %struct.NCFS { i32 3 } 102 // HOST: @c_ncfs = internal global %struct.NCFS undef, 103 104 // Regular base class -- allowed 105 __device__ T_B_T d_t_b_t; 106 // DEVICE: @d_t_b_t = addrspace(1) externally_initialized global %struct.T_B_T zeroinitializer, 107 // HOST: @d_t_b_t = internal global %struct.T_B_T undef, 108 __shared__ T_B_T s_t_b_t; 109 // DEVICE: @s_t_b_t = addrspace(3) global %struct.T_B_T undef, 110 // HOST: @s_t_b_t = internal global %struct.T_B_T undef, 111 __constant__ T_B_T c_t_b_t; 112 // DEVICE: @c_t_b_t = addrspace(4) externally_initialized global %struct.T_B_T zeroinitializer, 113 // HOST: @c_t_b_t = internal global %struct.T_B_T undef, 114 115 // Incapsulated object of allowed class -- allowed 116 __device__ T_F_T d_t_f_t; 117 // DEVICE: @d_t_f_t = addrspace(1) externally_initialized global %struct.T_F_T zeroinitializer, 118 // HOST: @d_t_f_t = internal global %struct.T_F_T undef, 119 __shared__ T_F_T s_t_f_t; 120 // DEVICE: @s_t_f_t = addrspace(3) global %struct.T_F_T undef, 121 // HOST: @s_t_f_t = internal global %struct.T_F_T undef, 122 __constant__ T_F_T c_t_f_t; 123 // DEVICE: @c_t_f_t = addrspace(4) externally_initialized global %struct.T_F_T zeroinitializer, 124 // HOST: @c_t_f_t = internal global %struct.T_F_T undef, 125 126 // array of allowed objects -- allowed 127 __device__ T_FA_T d_t_fa_t; 128 // DEVICE: @d_t_fa_t = addrspace(1) externally_initialized global %struct.T_FA_T zeroinitializer, 129 // HOST: @d_t_fa_t = internal global %struct.T_FA_T undef, 130 __shared__ T_FA_T s_t_fa_t; 131 // DEVICE: @s_t_fa_t = addrspace(3) global %struct.T_FA_T undef, 132 // HOST: @s_t_fa_t = internal global %struct.T_FA_T undef, 133 __constant__ T_FA_T c_t_fa_t; 134 // DEVICE: @c_t_fa_t = addrspace(4) externally_initialized global %struct.T_FA_T zeroinitializer, 135 // HOST: @c_t_fa_t = internal global %struct.T_FA_T undef, 136 137 138 // Calling empty base class initializer is OK 139 __device__ EC_I_EC d_ec_i_ec; 140 // DEVICE: @d_ec_i_ec = addrspace(1) externally_initialized global %struct.EC_I_EC zeroinitializer, 141 // HOST: @d_ec_i_ec = internal global %struct.EC_I_EC undef, 142 __shared__ EC_I_EC s_ec_i_ec; 143 // DEVICE: @s_ec_i_ec = addrspace(3) global %struct.EC_I_EC undef, 144 // HOST: @s_ec_i_ec = internal global %struct.EC_I_EC undef, 145 __constant__ EC_I_EC c_ec_i_ec; 146 // DEVICE: @c_ec_i_ec = addrspace(4) externally_initialized global %struct.EC_I_EC zeroinitializer, 147 // HOST: @c_ec_i_ec = internal global %struct.EC_I_EC undef, 148 149 // DEVICE: @_ZZ2dfvE4s_ec = internal addrspace(3) global %struct.EC undef 150 // DEVICE: @_ZZ2dfvE5s_etc = internal addrspace(3) global %struct.ETC undef 151 152 // DEVICE: @_ZZ2dfvE11const_array = internal addrspace(4) constant [5 x i32] [i32 1, i32 2, i32 3, i32 4, i32 5] 153 // DEVICE: @_ZZ2dfvE9const_int = internal addrspace(4) constant i32 123 154 155 // We should not emit global initializers for device-side variables. 156 // DEVICE-NOT: @__cxx_global_var_init 157 158 // Make sure that initialization restrictions do not apply to local 159 // variables. 160 __device__ void df() { 161 // NVPTX: %[[ec:.*]] = alloca %struct.EC 162 // NVPTX: %[[ed:.*]] = alloca %struct.ED 163 // NVPTX: %[[ecd:.*]] = alloca %struct.ECD 164 // NVPTX: %[[etc:.*]] = alloca %struct.ETC 165 // NVPTX: %[[uc:.*]] = alloca %struct.UC 166 // NVPTX: %[[ud:.*]] = alloca %struct.UD 167 // NVPTX: %[[eci:.*]] = alloca %struct.ECI 168 // NVPTX: %[[nec:.*]] = alloca %struct.NEC 169 // NVPTX: %[[ned:.*]] = alloca %struct.NED 170 // NVPTX: %[[ncv:.*]] = alloca %struct.NCV 171 // NVPTX: %[[vd:.*]] = alloca %struct.VD 172 // NVPTX: %[[ncf:.*]] = alloca %struct.NCF 173 // NVPTX: %[[ncfs:.*]] = alloca %struct.NCFS 174 // NVPTX: %[[utc:.*]] = alloca %struct.UTC 175 // NVPTX: %[[netc:.*]] = alloca %struct.NETC 176 // NVPTX: %[[ec_i_ec:.*]] = alloca %struct.EC_I_EC 177 // NVPTX: %[[ec_i_ec1:.*]] = alloca %struct.EC_I_EC1 178 // NVPTX: %[[t_v_t:.*]] = alloca %struct.T_V_T 179 // NVPTX: %[[t_b_nec:.*]] = alloca %struct.T_B_NEC 180 // NVPTX: %[[t_f_nec:.*]] = alloca %struct.T_F_NEC 181 // NVPTX: %[[t_fa_nec:.*]] = alloca %struct.T_FA_NEC 182 // NVPTX: %[[t_b_ned:.*]] = alloca %struct.T_B_NED 183 // NVPTX: %[[t_f_ned:.*]] = alloca %struct.T_F_NED 184 // NVPTX: %[[t_fa_ned:.*]] = alloca %struct.T_FA_NED 185 // AMDGCN: %[[ec:.*]] = addrspacecast %struct.EC addrspace(5)* %ec to %struct.EC* 186 // AMDGCN: %[[ed:.*]] = addrspacecast %struct.ED addrspace(5)* %ed to %struct.ED* 187 // AMDGCN: %[[ecd:.*]] = addrspacecast %struct.ECD addrspace(5)* %ecd to %struct.ECD* 188 // AMDGCN: %[[etc:.*]] = addrspacecast %struct.ETC addrspace(5)* %etc to %struct.ETC* 189 // AMDGCN: %[[uc:.*]] = addrspacecast %struct.UC addrspace(5)* %uc to %struct.UC* 190 // AMDGCN: %[[ud:.*]] = addrspacecast %struct.UD addrspace(5)* %ud to %struct.UD* 191 // AMDGCN: %[[eci:.*]] = addrspacecast %struct.ECI addrspace(5)* %eci to %struct.ECI* 192 // AMDGCN: %[[nec:.*]] = addrspacecast %struct.NEC addrspace(5)* %nec to %struct.NEC* 193 // AMDGCN: %[[ned:.*]] = addrspacecast %struct.NED addrspace(5)* %ned to %struct.NED* 194 // AMDGCN: %[[ncv:.*]] = addrspacecast %struct.NCV addrspace(5)* %ncv to %struct.NCV* 195 // AMDGCN: %[[vd:.*]] = addrspacecast %struct.VD addrspace(5)* %vd to %struct.VD* 196 // AMDGCN: %[[ncf:.*]] = addrspacecast %struct.NCF addrspace(5)* %ncf to %struct.NCF* 197 // AMDGCN: %[[ncfs:.*]] = addrspacecast %struct.NCFS addrspace(5)* %ncfs to %struct.NCFS* 198 // AMDGCN: %[[utc:.*]] = addrspacecast %struct.UTC addrspace(5)* %utc to %struct.UTC* 199 // AMDGCN: %[[netc:.*]] = addrspacecast %struct.NETC addrspace(5)* %netc to %struct.NETC* 200 // AMDGCN: %[[ec_i_ec:.*]] = addrspacecast %struct.EC_I_EC addrspace(5)* %ec_i_ec to %struct.EC_I_EC* 201 // AMDGCN: %[[ec_i_ec1:.*]] = addrspacecast %struct.EC_I_EC1 addrspace(5)* %ec_i_ec1 to %struct.EC_I_EC1* 202 // AMDGCN: %[[t_v_t:.*]] = addrspacecast %struct.T_V_T addrspace(5)* %t_v_t to %struct.T_V_T* 203 // AMDGCN: %[[t_b_nec:.*]] = addrspacecast %struct.T_B_NEC addrspace(5)* %t_b_nec to %struct.T_B_NEC* 204 // AMDGCN: %[[t_f_nec:.*]] = addrspacecast %struct.T_F_NEC addrspace(5)* %t_f_nec to %struct.T_F_NEC* 205 // AMDGCN: %[[t_fa_nec:.*]] = addrspacecast %struct.T_FA_NEC addrspace(5)* %t_fa_nec to %struct.T_FA_NEC* 206 // AMDGCN: %[[t_b_ned:.*]] = addrspacecast %struct.T_B_NED addrspace(5)* %t_b_ned to %struct.T_B_NED* 207 // AMDGCN: %[[t_f_ned:.*]] = addrspacecast %struct.T_F_NED addrspace(5)* %t_f_ned to %struct.T_F_NED* 208 // AMDGCN: %[[t_fa_ned:.*]] = addrspacecast %struct.T_FA_NED addrspace(5)* %t_fa_ned to %struct.T_FA_NED* 209 210 T t; 211 // DEVICE-NOT: call 212 EC ec; 213 // DEVICE: call void @_ZN2ECC1Ev(%struct.EC* %[[ec]]) 214 ED ed; 215 // DEVICE-NOT: call 216 ECD ecd; 217 // DEVICE: call void @_ZN3ECDC1Ev(%struct.ECD* %[[ecd]]) 218 ETC etc; 219 // DEVICE: call void @_ZN3ETCC1IJEEEDpT_(%struct.ETC* %[[etc]]) 220 UC uc; 221 // undefined constructor -- not allowed 222 // DEVICE: call void @_ZN2UCC1Ev(%struct.UC* %[[uc]]) 223 UD ud; 224 // undefined destructor -- not allowed 225 // DEVICE-NOT: call 226 ECI eci; 227 // empty constructor w/ initializer list -- not allowed 228 // DEVICE: call void @_ZN3ECIC1Ev(%struct.ECI* %[[eci]]) 229 NEC nec; 230 // non-empty constructor -- not allowed 231 // DEVICE: call void @_ZN3NECC1Ev(%struct.NEC* %[[nec]]) 232 // non-empty destructor -- not allowed 233 NED ned; 234 // no-constructor, virtual method -- not allowed 235 // DEVICE: call void @_ZN3NCVC1Ev(%struct.NCV* %[[ncv]]) 236 NCV ncv; 237 // DEVICE-NOT: call 238 VD vd; 239 // DEVICE: call void @_ZN2VDC1Ev(%struct.VD* %[[vd]]) 240 NCF ncf; 241 // DEVICE: call void @_ZN3NCFC1Ev(%struct.NCF* %[[ncf]]) 242 NCFS ncfs; 243 // DEVICE: call void @_ZN4NCFSC1Ev(%struct.NCFS* %[[ncfs]]) 244 UTC utc; 245 // DEVICE: call void @_ZN3UTCC1IJEEEDpT_(%struct.UTC* %[[utc]]) 246 NETC netc; 247 // DEVICE: call void @_ZN4NETCC1IJEEEDpT_(%struct.NETC* %[[netc]]) 248 T_B_T t_b_t; 249 // DEVICE-NOT: call 250 T_F_T t_f_t; 251 // DEVICE-NOT: call 252 T_FA_T t_fa_t; 253 // DEVICE-NOT: call 254 EC_I_EC ec_i_ec; 255 // DEVICE: call void @_ZN7EC_I_ECC1Ev(%struct.EC_I_EC* %[[ec_i_ec]]) 256 EC_I_EC1 ec_i_ec1; 257 // DEVICE: call void @_ZN8EC_I_EC1C1Ev(%struct.EC_I_EC1* %[[ec_i_ec1]]) 258 T_V_T t_v_t; 259 // DEVICE: call void @_ZN5T_V_TC1Ev(%struct.T_V_T* %[[t_v_t]]) 260 T_B_NEC t_b_nec; 261 // DEVICE: call void @_ZN7T_B_NECC1Ev(%struct.T_B_NEC* %[[t_b_nec]]) 262 T_F_NEC t_f_nec; 263 // DEVICE: call void @_ZN7T_F_NECC1Ev(%struct.T_F_NEC* %[[t_f_nec]]) 264 T_FA_NEC t_fa_nec; 265 // DEVICE: call void @_ZN8T_FA_NECC1Ev(%struct.T_FA_NEC* %[[t_fa_nec]]) 266 T_B_NED t_b_ned; 267 // DEVICE-NOT: call 268 T_F_NED t_f_ned; 269 // DEVICE-NOT: call 270 T_FA_NED t_fa_ned; 271 // DEVICE-NOT: call 272 static __shared__ EC s_ec; 273 // DEVICE-NOT: call void @_ZN2ECC1Ev(%struct.EC* addrspacecast (%struct.EC addrspace(3)* @_ZZ2dfvE4s_ec to %struct.EC*)) 274 static __shared__ ETC s_etc; 275 // DEVICE-NOT: call void @_ZN3ETCC1IJEEEDpT_(%struct.ETC* addrspacecast (%struct.ETC addrspace(3)* @_ZZ2dfvE5s_etc to %struct.ETC*)) 276 277 static const int const_array[] = {1, 2, 3, 4, 5}; 278 static const int const_int = 123; 279 280 // anchor point separating constructors and destructors 281 df(); // DEVICE: call void @_Z2dfv() 282 283 // Verify that we only call non-empty destructors 284 // DEVICE-NEXT: call void @_ZN8T_FA_NEDD1Ev(%struct.T_FA_NED* %[[t_fa_ned]]) 285 // DEVICE-NEXT: call void @_ZN7T_F_NEDD1Ev(%struct.T_F_NED* %[[t_f_ned]]) 286 // DEVICE-NEXT: call void @_ZN7T_B_NEDD1Ev(%struct.T_B_NED* %[[t_b_ned]]) 287 // DEVICE-NEXT: call void @_ZN2VDD1Ev(%struct.VD* %[[vd]]) 288 // DEVICE-NEXT: call void @_ZN3NEDD1Ev(%struct.NED* %[[ned]]) 289 // DEVICE-NEXT: call void @_ZN2UDD1Ev(%struct.UD* %[[ud]]) 290 // DEVICE-NEXT: call void @_ZN3ECDD1Ev(%struct.ECD* %[[ecd]]) 291 // DEVICE-NEXT: call void @_ZN2EDD1Ev(%struct.ED* %[[ed]]) 292 293 // DEVICE-NEXT: ret void 294 } 295 296 // We should not emit global init function. 297 // DEVICE-NOT: @_GLOBAL__sub_I 298