308 lines · plain
1// REQUIRES: nvptx-registered-target2// REQUIRES: amdgpu-registered-target3 4// Make sure we don't allow dynamic initialization for device5// 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 %s9// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -std=c++11 \10// RUN: -fno-threadsafe-statics -emit-llvm -o - %s | FileCheck -check-prefixes=HOST %s11 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 %s14 15#ifdef __clang__16#include "Inputs/cuda.h"17#endif18 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 host37// side.38static __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 -- allowed47__device__ T d_t;48// DEVICE: @d_t ={{.*}} addrspace(1) externally_initialized global %struct.T zeroinitializer49// 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 constructor65__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 undef74 75// empty destructor76__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 arguments97__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 -- allowed115__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 -- allowed126__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 -- allowed137__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 OK149__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 undef160// DEVICE: @_ZZ2dfvE5s_etc = internal addrspace(3) global %struct.ETC undef161 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 123164 165// We should not emit global initializers for device-side variables.166// DEVICE-NOT: @__cxx_global_var_init167 168// Make sure that initialization restrictions do not apply to local169// variables.170__device__ void df() {171 // NVPTX: %[[ec:.*]] = alloca %struct.EC172 // NVPTX: %[[ed:.*]] = alloca %struct.ED173 // NVPTX: %[[ecd:.*]] = alloca %struct.ECD174 // NVPTX: %[[etc:.*]] = alloca %struct.ETC175 // NVPTX: %[[uc:.*]] = alloca %struct.UC176 // NVPTX: %[[ud:.*]] = alloca %struct.UD177 // NVPTX: %[[eci:.*]] = alloca %struct.ECI178 // NVPTX: %[[nec:.*]] = alloca %struct.NEC179 // NVPTX: %[[ned:.*]] = alloca %struct.NED180 // NVPTX: %[[ncv:.*]] = alloca %struct.NCV181 // NVPTX: %[[vd:.*]] = alloca %struct.VD182 // NVPTX: %[[ncf:.*]] = alloca %struct.NCF183 // NVPTX: %[[ncfs:.*]] = alloca %struct.NCFS184 // NVPTX: %[[utc:.*]] = alloca %struct.UTC185 // NVPTX: %[[netc:.*]] = alloca %struct.NETC186 // NVPTX: %[[ec_i_ec:.*]] = alloca %struct.EC_I_EC187 // NVPTX: %[[ec_i_ec1:.*]] = alloca %struct.EC_I_EC1188 // NVPTX: %[[t_v_t:.*]] = alloca %struct.T_V_T189 // NVPTX: %[[t_b_nec:.*]] = alloca %struct.T_B_NEC190 // NVPTX: %[[t_f_nec:.*]] = alloca %struct.T_F_NEC191 // NVPTX: %[[t_fa_nec:.*]] = alloca %struct.T_FA_NEC192 // NVPTX: %[[t_b_ned:.*]] = alloca %struct.T_B_NED193 // NVPTX: %[[t_f_ned:.*]] = alloca %struct.T_F_NED194 // NVPTX: %[[t_fa_ned:.*]] = alloca %struct.T_FA_NED195 // AMDGCN: %[[ec:.*]] ={{.*}} addrspacecast ptr addrspace(5) %ec to ptr196 // AMDGCN: %[[ed:.*]] ={{.*}} addrspacecast ptr addrspace(5) %ed to ptr197 // AMDGCN: %[[ecd:.*]] ={{.*}} addrspacecast ptr addrspace(5) %ecd to ptr198 // AMDGCN: %[[etc:.*]] ={{.*}} addrspacecast ptr addrspace(5) %etc to ptr199 // AMDGCN: %[[uc:.*]] ={{.*}} addrspacecast ptr addrspace(5) %uc to ptr200 // AMDGCN: %[[ud:.*]] ={{.*}} addrspacecast ptr addrspace(5) %ud to ptr201 // AMDGCN: %[[eci:.*]] ={{.*}} addrspacecast ptr addrspace(5) %eci to ptr202 // AMDGCN: %[[nec:.*]] ={{.*}} addrspacecast ptr addrspace(5) %nec to ptr203 // AMDGCN: %[[ned:.*]] ={{.*}} addrspacecast ptr addrspace(5) %ned to ptr204 // AMDGCN: %[[ncv:.*]] ={{.*}} addrspacecast ptr addrspace(5) %ncv to ptr205 // AMDGCN: %[[vd:.*]] ={{.*}} addrspacecast ptr addrspace(5) %vd to ptr206 // AMDGCN: %[[ncf:.*]] ={{.*}} addrspacecast ptr addrspace(5) %ncf to ptr207 // AMDGCN: %[[ncfs:.*]] ={{.*}} addrspacecast ptr addrspace(5) %ncfs to ptr208 // AMDGCN: %[[utc:.*]] ={{.*}} addrspacecast ptr addrspace(5) %utc to ptr209 // AMDGCN: %[[netc:.*]] ={{.*}} addrspacecast ptr addrspace(5) %netc to ptr210 // AMDGCN: %[[ec_i_ec:.*]] ={{.*}} addrspacecast ptr addrspace(5) %ec_i_ec to ptr211 // AMDGCN: %[[ec_i_ec1:.*]] ={{.*}} addrspacecast ptr addrspace(5) %ec_i_ec1 to ptr212 // AMDGCN: %[[t_v_t:.*]] ={{.*}} addrspacecast ptr addrspace(5) %t_v_t to ptr213 // AMDGCN: %[[t_b_nec:.*]] ={{.*}} addrspacecast ptr addrspace(5) %t_b_nec to ptr214 // AMDGCN: %[[t_f_nec:.*]] ={{.*}} addrspacecast ptr addrspace(5) %t_f_nec to ptr215 // AMDGCN: %[[t_fa_nec:.*]] ={{.*}} addrspacecast ptr addrspace(5) %t_fa_nec to ptr216 // AMDGCN: %[[t_b_ned:.*]] ={{.*}} addrspacecast ptr addrspace(5) %t_b_ned to ptr217 // AMDGCN: %[[t_f_ned:.*]] ={{.*}} addrspacecast ptr addrspace(5) %t_f_ned to ptr218 // AMDGCN: %[[t_fa_ned:.*]] ={{.*}} addrspacecast ptr addrspace(5) %t_fa_ned to ptr219 220 T t;221 // DEVICE-NOT: call222 EC ec;223 // DEVICE: call void @_ZN2ECC1Ev(ptr {{[^,]*}} %[[ec]])224 ED ed;225 // DEVICE-NOT: call226 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 allowed232 // DEVICE: call void @_ZN2UCC1Ev(ptr {{[^,]*}} %[[uc]])233 UD ud;234 // undefined destructor -- not allowed235 // DEVICE-NOT: call236 ECI eci;237 // empty constructor w/ initializer list -- not allowed238 // DEVICE: call void @_ZN3ECIC1Ev(ptr {{[^,]*}} %[[eci]])239 NEC nec;240 // non-empty constructor -- not allowed241 // DEVICE: call void @_ZN3NECC1Ev(ptr {{[^,]*}} %[[nec]])242 // non-empty destructor -- not allowed243 NED ned;244 // no-constructor, virtual method -- not allowed245 // DEVICE: call void @_ZN3NCVC1Ev(ptr {{[^,]*}} %[[ncv]])246 NCV ncv;247 // DEVICE-NOT: call248 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: call260 T_F_T t_f_t;261 // DEVICE-NOT: call262 T_FA_T t_fa_t;263 // DEVICE-NOT: call264 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: call278 T_F_NED t_f_ned;279 // DEVICE-NOT: call280 T_FA_NED t_fa_ned;281 // DEVICE-NOT: call282 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 destructors291 df(); // DEVICE: call void @_Z2dfv()292 293 // Verify that we only call non-empty destructors294 // 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 void304}305 306// We should not emit global init function.307// DEVICE-NOT: @_GLOBAL__sub_I308