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