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 // 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.
foo()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 global %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 global %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 global %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 global %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 global %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 global %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 global %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 global %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 global %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 global %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 global %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.
df()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 %struct.EC addrspace(5)* %ec to %struct.EC*
196 // AMDGCN: %[[ed:.*]] ={{.*}} addrspacecast %struct.ED addrspace(5)* %ed to %struct.ED*
197 // AMDGCN: %[[ecd:.*]] ={{.*}} addrspacecast %struct.ECD addrspace(5)* %ecd to %struct.ECD*
198 // AMDGCN: %[[etc:.*]] ={{.*}} addrspacecast %struct.ETC addrspace(5)* %etc to %struct.ETC*
199 // AMDGCN: %[[uc:.*]] ={{.*}} addrspacecast %struct.UC addrspace(5)* %uc to %struct.UC*
200 // AMDGCN: %[[ud:.*]] ={{.*}} addrspacecast %struct.UD addrspace(5)* %ud to %struct.UD*
201 // AMDGCN: %[[eci:.*]] ={{.*}} addrspacecast %struct.ECI addrspace(5)* %eci to %struct.ECI*
202 // AMDGCN: %[[nec:.*]] ={{.*}} addrspacecast %struct.NEC addrspace(5)* %nec to %struct.NEC*
203 // AMDGCN: %[[ned:.*]] ={{.*}} addrspacecast %struct.NED addrspace(5)* %ned to %struct.NED*
204 // AMDGCN: %[[ncv:.*]] ={{.*}} addrspacecast %struct.NCV addrspace(5)* %ncv to %struct.NCV*
205 // AMDGCN: %[[vd:.*]] ={{.*}} addrspacecast %struct.VD addrspace(5)* %vd to %struct.VD*
206 // AMDGCN: %[[ncf:.*]] ={{.*}} addrspacecast %struct.NCF addrspace(5)* %ncf to %struct.NCF*
207 // AMDGCN: %[[ncfs:.*]] ={{.*}} addrspacecast %struct.NCFS addrspace(5)* %ncfs to %struct.NCFS*
208 // AMDGCN: %[[utc:.*]] ={{.*}} addrspacecast %struct.UTC addrspace(5)* %utc to %struct.UTC*
209 // AMDGCN: %[[netc:.*]] ={{.*}} addrspacecast %struct.NETC addrspace(5)* %netc to %struct.NETC*
210 // AMDGCN: %[[ec_i_ec:.*]] ={{.*}} addrspacecast %struct.EC_I_EC addrspace(5)* %ec_i_ec to %struct.EC_I_EC*
211 // AMDGCN: %[[ec_i_ec1:.*]] ={{.*}} addrspacecast %struct.EC_I_EC1 addrspace(5)* %ec_i_ec1 to %struct.EC_I_EC1*
212 // AMDGCN: %[[t_v_t:.*]] ={{.*}} addrspacecast %struct.T_V_T addrspace(5)* %t_v_t to %struct.T_V_T*
213 // AMDGCN: %[[t_b_nec:.*]] ={{.*}} addrspacecast %struct.T_B_NEC addrspace(5)* %t_b_nec to %struct.T_B_NEC*
214 // AMDGCN: %[[t_f_nec:.*]] ={{.*}} addrspacecast %struct.T_F_NEC addrspace(5)* %t_f_nec to %struct.T_F_NEC*
215 // AMDGCN: %[[t_fa_nec:.*]] ={{.*}} addrspacecast %struct.T_FA_NEC addrspace(5)* %t_fa_nec to %struct.T_FA_NEC*
216 // AMDGCN: %[[t_b_ned:.*]] ={{.*}} addrspacecast %struct.T_B_NED addrspace(5)* %t_b_ned to %struct.T_B_NED*
217 // AMDGCN: %[[t_f_ned:.*]] ={{.*}} addrspacecast %struct.T_F_NED addrspace(5)* %t_f_ned to %struct.T_F_NED*
218 // AMDGCN: %[[t_fa_ned:.*]] ={{.*}} addrspacecast %struct.T_FA_NED addrspace(5)* %t_fa_ned to %struct.T_FA_NED*
219
220 T t;
221 // DEVICE-NOT: call
222 EC ec;
223 // DEVICE: call void @_ZN2ECC1Ev(%struct.EC* {{[^,]*}} %[[ec]])
224 ED ed;
225 // DEVICE-NOT: call
226 ECD ecd;
227 // DEVICE: call void @_ZN3ECDC1Ev(%struct.ECD* {{[^,]*}} %[[ecd]])
228 ETC etc;
229 // DEVICE: call void @_ZN3ETCC1IJEEEDpT_(%struct.ETC* {{[^,]*}} %[[etc]])
230 UC uc;
231 // undefined constructor -- not allowed
232 // DEVICE: call void @_ZN2UCC1Ev(%struct.UC* {{[^,]*}} %[[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(%struct.ECI* {{[^,]*}} %[[eci]])
239 NEC nec;
240 // non-empty constructor -- not allowed
241 // DEVICE: call void @_ZN3NECC1Ev(%struct.NEC* {{[^,]*}} %[[nec]])
242 // non-empty destructor -- not allowed
243 NED ned;
244 // no-constructor, virtual method -- not allowed
245 // DEVICE: call void @_ZN3NCVC1Ev(%struct.NCV* {{[^,]*}} %[[ncv]])
246 NCV ncv;
247 // DEVICE-NOT: call
248 VD vd;
249 // DEVICE: call void @_ZN2VDC1Ev(%struct.VD* {{[^,]*}} %[[vd]])
250 NCF ncf;
251 // DEVICE: call void @_ZN3NCFC1Ev(%struct.NCF* {{[^,]*}} %[[ncf]])
252 NCFS ncfs;
253 // DEVICE: call void @_ZN4NCFSC1Ev(%struct.NCFS* {{[^,]*}} %[[ncfs]])
254 UTC utc;
255 // DEVICE: call void @_ZN3UTCC1IJEEEDpT_(%struct.UTC* {{[^,]*}} %[[utc]])
256 NETC netc;
257 // DEVICE: call void @_ZN4NETCC1IJEEEDpT_(%struct.NETC* {{[^,]*}} %[[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(%struct.EC_I_EC* {{[^,]*}} %[[ec_i_ec]])
266 EC_I_EC1 ec_i_ec1;
267 // DEVICE: call void @_ZN8EC_I_EC1C1Ev(%struct.EC_I_EC1* {{[^,]*}} %[[ec_i_ec1]])
268 T_V_T t_v_t;
269 // DEVICE: call void @_ZN5T_V_TC1Ev(%struct.T_V_T* {{[^,]*}} %[[t_v_t]])
270 T_B_NEC t_b_nec;
271 // DEVICE: call void @_ZN7T_B_NECC1Ev(%struct.T_B_NEC* {{[^,]*}} %[[t_b_nec]])
272 T_F_NEC t_f_nec;
273 // DEVICE: call void @_ZN7T_F_NECC1Ev(%struct.T_F_NEC* {{[^,]*}} %[[t_f_nec]])
274 T_FA_NEC t_fa_nec;
275 // DEVICE: call void @_ZN8T_FA_NECC1Ev(%struct.T_FA_NEC* {{[^,]*}} %[[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(%struct.EC* addrspacecast (%struct.EC addrspace(3)* @_ZZ2dfvE4s_ec to %struct.EC*))
284 static __shared__ ETC s_etc;
285 // DEVICE-NOT: call void @_ZN3ETCC1IJEEEDpT_(%struct.ETC* addrspacecast (%struct.ETC addrspace(3)* @_ZZ2dfvE5s_etc to %struct.ETC*))
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(%struct.T_FA_NED* {{[^,]*}} %[[t_fa_ned]])
295 // DEVICE-NEXT: call void @_ZN7T_F_NEDD1Ev(%struct.T_F_NED* {{[^,]*}} %[[t_f_ned]])
296 // DEVICE-NEXT: call void @_ZN7T_B_NEDD1Ev(%struct.T_B_NED* {{[^,]*}} %[[t_b_ned]])
297 // DEVICE-NEXT: call void @_ZN2VDD1Ev(%struct.VD* {{[^,]*}} %[[vd]])
298 // DEVICE-NEXT: call void @_ZN3NEDD1Ev(%struct.NED* {{[^,]*}} %[[ned]])
299 // DEVICE-NEXT: call void @_ZN2UDD1Ev(%struct.UD* {{[^,]*}} %[[ud]])
300 // DEVICE-NEXT: call void @_ZN3ECDD1Ev(%struct.ECD* {{[^,]*}} %[[ecd]])
301 // DEVICE-NEXT: call void @_ZN2EDD1Ev(%struct.ED* {{[^,]*}} %[[ed]])
302
303 // DEVICE-NEXT: ret void
304 }
305
306 // We should not emit global init function.
307 // DEVICE-NOT: @_GLOBAL__sub_I
308