| 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 | |