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 9 #ifdef __clang__ 10 #include "Inputs/cuda.h" 11 #endif 12 13 // Use the types we share with Sema tests. 14 #include "Inputs/cuda-initializers.h" 15 16 __device__ int d_v; 17 // CHECK: @d_v = addrspace(1) externally_initialized global i32 0, 18 __shared__ int s_v; 19 // CHECK: @s_v = addrspace(3) global i32 undef, 20 __constant__ int c_v; 21 // CHECK: addrspace(4) externally_initialized global i32 0, 22 23 __device__ int d_v_i = 1; 24 // CHECK: @d_v_i = addrspace(1) externally_initialized global i32 1, 25 26 // trivial constructor -- allowed 27 __device__ T d_t; 28 // CHECK: @d_t = addrspace(1) externally_initialized global %struct.T zeroinitializer 29 __shared__ T s_t; 30 // CHECK: @s_t = addrspace(3) global %struct.T undef, 31 __constant__ T c_t; 32 // CHECK: @c_t = addrspace(4) externally_initialized global %struct.T zeroinitializer, 33 34 __device__ T d_t_i = {2}; 35 // CHECK: @d_t_i = addrspace(1) externally_initialized global %struct.T { i32 2 }, 36 __constant__ T c_t_i = {2}; 37 // CHECK: @c_t_i = addrspace(4) externally_initialized global %struct.T { i32 2 }, 38 39 // empty constructor 40 __device__ EC d_ec; 41 // CHECK: @d_ec = addrspace(1) externally_initialized global %struct.EC zeroinitializer, 42 __shared__ EC s_ec; 43 // CHECK: @s_ec = addrspace(3) global %struct.EC undef, 44 __constant__ EC c_ec; 45 // CHECK: @c_ec = addrspace(4) externally_initialized global %struct.EC zeroinitializer, 46 47 // empty destructor 48 __device__ ED d_ed; 49 // CHECK: @d_ed = addrspace(1) externally_initialized global %struct.ED zeroinitializer, 50 __shared__ ED s_ed; 51 // CHECK: @s_ed = addrspace(3) global %struct.ED undef, 52 __constant__ ED c_ed; 53 // CHECK: @c_ed = addrspace(4) externally_initialized global %struct.ED zeroinitializer, 54 55 __device__ ECD d_ecd; 56 // CHECK: @d_ecd = addrspace(1) externally_initialized global %struct.ECD zeroinitializer, 57 __shared__ ECD s_ecd; 58 // CHECK: @s_ecd = addrspace(3) global %struct.ECD undef, 59 __constant__ ECD c_ecd; 60 // CHECK: @c_ecd = addrspace(4) externally_initialized global %struct.ECD zeroinitializer, 61 62 // empty templated constructor -- allowed with no arguments 63 __device__ ETC d_etc; 64 // CHECK: @d_etc = addrspace(1) externally_initialized global %struct.ETC zeroinitializer, 65 __shared__ ETC s_etc; 66 // CHECK: @s_etc = addrspace(3) global %struct.ETC undef, 67 __constant__ ETC c_etc; 68 // CHECK: @c_etc = addrspace(4) externally_initialized global %struct.ETC zeroinitializer, 69 70 __device__ NCFS d_ncfs; 71 // CHECK: @d_ncfs = addrspace(1) externally_initialized global %struct.NCFS { i32 3 } 72 __constant__ NCFS c_ncfs; 73 // CHECK: @c_ncfs = addrspace(4) externally_initialized global %struct.NCFS { i32 3 } 74 75 // Regular base class -- allowed 76 __device__ T_B_T d_t_b_t; 77 // CHECK: @d_t_b_t = addrspace(1) externally_initialized global %struct.T_B_T zeroinitializer, 78 __shared__ T_B_T s_t_b_t; 79 // CHECK: @s_t_b_t = addrspace(3) global %struct.T_B_T undef, 80 __constant__ T_B_T c_t_b_t; 81 // CHECK: @c_t_b_t = addrspace(4) externally_initialized global %struct.T_B_T zeroinitializer, 82 83 // Incapsulated object of allowed class -- allowed 84 __device__ T_F_T d_t_f_t; 85 // CHECK: @d_t_f_t = addrspace(1) externally_initialized global %struct.T_F_T zeroinitializer, 86 __shared__ T_F_T s_t_f_t; 87 // CHECK: @s_t_f_t = addrspace(3) global %struct.T_F_T undef, 88 __constant__ T_F_T c_t_f_t; 89 // CHECK: @c_t_f_t = addrspace(4) externally_initialized global %struct.T_F_T zeroinitializer, 90 91 // array of allowed objects -- allowed 92 __device__ T_FA_T d_t_fa_t; 93 // CHECK: @d_t_fa_t = addrspace(1) externally_initialized global %struct.T_FA_T zeroinitializer, 94 __shared__ T_FA_T s_t_fa_t; 95 // CHECK: @s_t_fa_t = addrspace(3) global %struct.T_FA_T undef, 96 __constant__ T_FA_T c_t_fa_t; 97 // CHECK: @c_t_fa_t = addrspace(4) externally_initialized global %struct.T_FA_T zeroinitializer, 98 99 100 // Calling empty base class initializer is OK 101 __device__ EC_I_EC d_ec_i_ec; 102 // CHECK: @d_ec_i_ec = addrspace(1) externally_initialized global %struct.EC_I_EC zeroinitializer, 103 __shared__ EC_I_EC s_ec_i_ec; 104 // CHECK: @s_ec_i_ec = addrspace(3) global %struct.EC_I_EC undef, 105 __constant__ EC_I_EC c_ec_i_ec; 106 // CHECK: @c_ec_i_ec = addrspace(4) externally_initialized global %struct.EC_I_EC zeroinitializer, 107 108 // We should not emit global initializers for device-side variables. 109 // CHECK-NOT: @__cxx_global_var_init 110 111 // Make sure that initialization restrictions do not apply to local 112 // variables. 113 __device__ void df() { 114 T t; 115 // CHECK-NOT: call 116 EC ec; 117 // CHECK: call void @_ZN2ECC1Ev(%struct.EC* %ec) 118 ED ed; 119 // CHECK-NOT: call 120 ECD ecd; 121 // CHECK: call void @_ZN3ECDC1Ev(%struct.ECD* %ecd) 122 ETC etc; 123 // CHECK: call void @_ZN3ETCC1IJEEEDpT_(%struct.ETC* %etc) 124 UC uc; 125 // undefined constructor -- not allowed 126 // CHECK: call void @_ZN2UCC1Ev(%struct.UC* %uc) 127 UD ud; 128 // undefined destructor -- not allowed 129 // CHECK-NOT: call 130 ECI eci; 131 // empty constructor w/ initializer list -- not allowed 132 // CHECK: call void @_ZN3ECIC1Ev(%struct.ECI* %eci) 133 NEC nec; 134 // non-empty constructor -- not allowed 135 // CHECK: call void @_ZN3NECC1Ev(%struct.NEC* %nec) 136 // non-empty destructor -- not allowed 137 NED ned; 138 // no-constructor, virtual method -- not allowed 139 // CHECK: call void @_ZN3NCVC1Ev(%struct.NCV* %ncv) 140 NCV ncv; 141 // CHECK-NOT: call 142 VD vd; 143 // CHECK: call void @_ZN2VDC1Ev(%struct.VD* %vd) 144 NCF ncf; 145 // CHECK: call void @_ZN3NCFC1Ev(%struct.NCF* %ncf) 146 NCFS ncfs; 147 // CHECK: call void @_ZN4NCFSC1Ev(%struct.NCFS* %ncfs) 148 UTC utc; 149 // CHECK: call void @_ZN3UTCC1IJEEEDpT_(%struct.UTC* %utc) 150 NETC netc; 151 // CHECK: call void @_ZN4NETCC1IJEEEDpT_(%struct.NETC* %netc) 152 T_B_T t_b_t; 153 // CHECK-NOT: call 154 T_F_T t_f_t; 155 // CHECK-NOT: call 156 T_FA_T t_fa_t; 157 // CHECK-NOT: call 158 EC_I_EC ec_i_ec; 159 // CHECK: call void @_ZN7EC_I_ECC1Ev(%struct.EC_I_EC* %ec_i_ec) 160 EC_I_EC1 ec_i_ec1; 161 // CHECK: call void @_ZN8EC_I_EC1C1Ev(%struct.EC_I_EC1* %ec_i_ec1) 162 T_V_T t_v_t; 163 // CHECK: call void @_ZN5T_V_TC1Ev(%struct.T_V_T* %t_v_t) 164 T_B_NEC t_b_nec; 165 // CHECK: call void @_ZN7T_B_NECC1Ev(%struct.T_B_NEC* %t_b_nec) 166 T_F_NEC t_f_nec; 167 // CHECK: call void @_ZN7T_F_NECC1Ev(%struct.T_F_NEC* %t_f_nec) 168 T_FA_NEC t_fa_nec; 169 // CHECK: call void @_ZN8T_FA_NECC1Ev(%struct.T_FA_NEC* %t_fa_nec) 170 T_B_NED t_b_ned; 171 // CHECK-NOT: call 172 T_F_NED t_f_ned; 173 // CHECK-NOT: call 174 T_FA_NED t_fa_ned; 175 // CHECK-NOT: call 176 static __shared__ EC s_ec; 177 // CHECK-NOT: call void @_ZN2ECC1Ev(%struct.EC* addrspacecast (%struct.EC addrspace(3)* @_ZZ2dfvE4s_ec to %struct.EC*)) 178 static __shared__ ETC s_etc; 179 // CHECK-NOT: call void @_ZN3ETCC1IJEEEDpT_(%struct.ETC* addrspacecast (%struct.ETC addrspace(3)* @_ZZ2dfvE5s_etc to %struct.ETC*)) 180 181 // anchor point separating constructors and destructors 182 df(); // CHECK: call void @_Z2dfv() 183 184 // Verify that we only call non-empty destructors 185 // CHECK-NEXT: call void @_ZN8T_FA_NEDD1Ev(%struct.T_FA_NED* %t_fa_ned) #6 186 // CHECK-NEXT: call void @_ZN7T_F_NEDD1Ev(%struct.T_F_NED* %t_f_ned) #6 187 // CHECK-NEXT: call void @_ZN7T_B_NEDD1Ev(%struct.T_B_NED* %t_b_ned) #6 188 // CHECK-NEXT: call void @_ZN2VDD1Ev(%struct.VD* %vd) 189 // CHECK-NEXT: call void @_ZN3NEDD1Ev(%struct.NED* %ned) 190 // CHECK-NEXT: call void @_ZN2UDD1Ev(%struct.UD* %ud) 191 // CHECK-NEXT: call void @_ZN3ECDD1Ev(%struct.ECD* %ecd) 192 // CHECK-NEXT: call void @_ZN2EDD1Ev(%struct.ED* %ed) 193 194 // CHECK-NEXT: ret void 195 } 196 197 // We should not emit global init function. 198 // CHECK-NOT: @_GLOBAL__sub_I 199
