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.
df()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