1 // REQUIRES: nvptx-registered-target
2 // REQUIRES: amdgpu-registered-target
4 // Make sure we don't allow dynamic initialization for device
5 // variables, but accept empty constructors allowed by CUDA.
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
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
16 #include "Inputs/cuda.h"
19 // Use the types we share with Sema tests.
20 #include "Inputs/cuda-initializers.h"
23 // DEVICE: @d_v = addrspace(1) externally_initialized global i32 0,
24 // HOST: @d_v = internal global i32 undef,
26 // DEVICE: @s_v = addrspace(3) global i32 undef,
27 // HOST: @s_v = internal global i32 undef,
29 // DEVICE: addrspace(4) externally_initialized global i32 0,
30 // HOST: @c_v = internal global i32 undef,
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,
36 // trivial constructor -- allowed
38 // DEVICE: @d_t = addrspace(1) externally_initialized global %struct.T zeroinitializer
39 // HOST: @d_t = internal global %struct.T undef,
41 // DEVICE: @s_t = addrspace(3) global %struct.T undef,
42 // HOST: @s_t = internal global %struct.T undef,
44 // DEVICE: @c_t = addrspace(4) externally_initialized global %struct.T zeroinitializer,
45 // HOST: @c_t = internal global %struct.T undef,
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,
56 // DEVICE: @d_ec = addrspace(1) externally_initialized global %struct.EC zeroinitializer,
57 // HOST: @d_ec = internal global %struct.EC undef,
59 // DEVICE: @s_ec = addrspace(3) global %struct.EC undef,
60 // HOST: @s_ec = internal global %struct.EC undef,
62 // DEVICE: @c_ec = addrspace(4) externally_initialized global %struct.EC zeroinitializer,
63 // HOST: @c_ec = internal global %struct.EC undef
67 // DEVICE: @d_ed = addrspace(1) externally_initialized global %struct.ED zeroinitializer,
68 // HOST: @d_ed = internal global %struct.ED undef,
70 // DEVICE: @s_ed = addrspace(3) global %struct.ED undef,
71 // HOST: @s_ed = internal global %struct.ED undef,
73 // DEVICE: @c_ed = addrspace(4) externally_initialized global %struct.ED zeroinitializer,
74 // HOST: @c_ed = internal global %struct.ED undef,
77 // DEVICE: @d_ecd = addrspace(1) externally_initialized global %struct.ECD zeroinitializer,
78 // HOST: @d_ecd = internal global %struct.ECD undef,
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,
86 // empty templated constructor -- allowed with no arguments
88 // DEVICE: @d_etc = addrspace(1) externally_initialized global %struct.ETC zeroinitializer,
89 // HOST: @d_etc = internal global %struct.ETC undef,
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,
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,
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,
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,
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,
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,
149 // DEVICE: @_ZZ2dfvE4s_ec = internal addrspace(3) global %struct.EC undef
150 // DEVICE: @_ZZ2dfvE5s_etc = internal addrspace(3) global %struct.ETC undef
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
155 // We should not emit global initializers for device-side variables.
156 // DEVICE-NOT: @__cxx_global_var_init
158 // Make sure that initialization restrictions do not apply to local
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*
213 // DEVICE: call void @_ZN2ECC1Ev(%struct.EC* %[[ec]])
217 // DEVICE: call void @_ZN3ECDC1Ev(%struct.ECD* %[[ecd]])
219 // DEVICE: call void @_ZN3ETCC1IJEEEDpT_(%struct.ETC* %[[etc]])
221 // undefined constructor -- not allowed
222 // DEVICE: call void @_ZN2UCC1Ev(%struct.UC* %[[uc]])
224 // undefined destructor -- not allowed
227 // empty constructor w/ initializer list -- not allowed
228 // DEVICE: call void @_ZN3ECIC1Ev(%struct.ECI* %[[eci]])
230 // non-empty constructor -- not allowed
231 // DEVICE: call void @_ZN3NECC1Ev(%struct.NEC* %[[nec]])
232 // non-empty destructor -- not allowed
234 // no-constructor, virtual method -- not allowed
235 // DEVICE: call void @_ZN3NCVC1Ev(%struct.NCV* %[[ncv]])
239 // DEVICE: call void @_ZN2VDC1Ev(%struct.VD* %[[vd]])
241 // DEVICE: call void @_ZN3NCFC1Ev(%struct.NCF* %[[ncf]])
243 // DEVICE: call void @_ZN4NCFSC1Ev(%struct.NCFS* %[[ncfs]])
245 // DEVICE: call void @_ZN3UTCC1IJEEEDpT_(%struct.UTC* %[[utc]])
247 // DEVICE: call void @_ZN4NETCC1IJEEEDpT_(%struct.NETC* %[[netc]])
255 // DEVICE: call void @_ZN7EC_I_ECC1Ev(%struct.EC_I_EC* %[[ec_i_ec]])
257 // DEVICE: call void @_ZN8EC_I_EC1C1Ev(%struct.EC_I_EC1* %[[ec_i_ec1]])
259 // DEVICE: call void @_ZN5T_V_TC1Ev(%struct.T_V_T* %[[t_v_t]])
261 // DEVICE: call void @_ZN7T_B_NECC1Ev(%struct.T_B_NEC* %[[t_b_nec]])
263 // DEVICE: call void @_ZN7T_F_NECC1Ev(%struct.T_F_NEC* %[[t_f_nec]])
265 // DEVICE: call void @_ZN8T_FA_NECC1Ev(%struct.T_FA_NEC* %[[t_fa_nec]])
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*))
277 static const int const_array[] = {1, 2, 3, 4, 5};
278 static const int const_int = 123;
280 // anchor point separating constructors and destructors
281 df(); // DEVICE: call void @_Z2dfv()
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]])
293 // DEVICE-NEXT: ret void
296 // We should not emit global init function.
297 // DEVICE-NOT: @_GLOBAL__sub_I