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 -verify %s -triple nvptx64-nvidia-cuda -fcuda-is-device -std=c++11 %s
7
8 #ifdef __clang__
9 #include "Inputs/cuda.h"
10 #endif
11
12 // Use the types we share with CodeGen tests.
13 #include "Inputs/cuda-initializers.h"
14
15 __shared__ int s_v_i = 1;
16 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
17
18 __device__ int d_v_f = f();
19 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
20 __shared__ int s_v_f = f();
21 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
22 __constant__ int c_v_f = f();
23 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
24
25 __shared__ T s_t_i = {2};
26 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
27
28 __device__ EC d_ec_i(3);
29 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
30 __shared__ EC s_ec_i(3);
31 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
32 __constant__ EC c_ec_i(3);
33 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
34
35 __device__ EC d_ec_i2 = {3};
36 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
37 __shared__ EC s_ec_i2 = {3};
38 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
39 __constant__ EC c_ec_i2 = {3};
40 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
41
42 __device__ ETC d_etc_i(3);
43 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
44 __shared__ ETC s_etc_i(3);
45 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
46 __constant__ ETC c_etc_i(3);
47 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
48
49 __device__ ETC d_etc_i2 = {3};
50 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
51 __shared__ ETC s_etc_i2 = {3};
52 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
53 __constant__ ETC c_etc_i2 = {3};
54 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
55
56 __device__ UC d_uc;
57 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
58 __shared__ UC s_uc;
59 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
60 __constant__ UC c_uc;
61 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
62
63 __device__ UD d_ud;
64 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
65 __shared__ UD s_ud;
66 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
67 __constant__ UD c_ud;
68 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
69
70 __device__ ECI d_eci;
71 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
72 __shared__ ECI s_eci;
73 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
74 __constant__ ECI c_eci;
75 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
76
77 __device__ NEC d_nec;
78 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
79 __shared__ NEC s_nec;
80 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
81 __constant__ NEC c_nec;
82 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
83
84 __device__ NED d_ned;
85 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
86 __shared__ NED s_ned;
87 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
88 __constant__ NED c_ned;
89 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
90
91 __device__ NCV d_ncv;
92 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
93 __shared__ NCV s_ncv;
94 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
95 __constant__ NCV c_ncv;
96 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
97
98 __device__ VD d_vd;
99 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
100 __shared__ VD s_vd;
101 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
102 __constant__ VD c_vd;
103 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
104
105 __device__ NCF d_ncf;
106 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
107 __shared__ NCF s_ncf;
108 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
109 __constant__ NCF c_ncf;
110 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
111
112 __shared__ NCFS s_ncfs;
113 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
114
115 __device__ UTC d_utc;
116 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
117 __shared__ UTC s_utc;
118 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
119 __constant__ UTC c_utc;
120 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
121
122 __device__ UTC d_utc_i(3);
123 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
124 __shared__ UTC s_utc_i(3);
125 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
126 __constant__ UTC c_utc_i(3);
127 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
128
129 __device__ NETC d_netc;
130 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
131 __shared__ NETC s_netc;
132 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
133 __constant__ NETC c_netc;
134 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
135
136 __device__ NETC d_netc_i(3);
137 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
138 __shared__ NETC s_netc_i(3);
139 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
140 __constant__ NETC c_netc_i(3);
141 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
142
143 __device__ EC_I_EC1 d_ec_i_ec1;
144 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
145 __shared__ EC_I_EC1 s_ec_i_ec1;
146 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
147 __constant__ EC_I_EC1 c_ec_i_ec1;
148 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
149
150 __device__ T_V_T d_t_v_t;
151 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
152 __shared__ T_V_T s_t_v_t;
153 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
154 __constant__ T_V_T c_t_v_t;
155 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
156
157 __device__ T_B_NEC d_t_b_nec;
158 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
159 __shared__ T_B_NEC s_t_b_nec;
160 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
161 __constant__ T_B_NEC c_t_b_nec;
162 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
163
164 __device__ T_F_NEC d_t_f_nec;
165 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
166 __shared__ T_F_NEC s_t_f_nec;
167 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
168 __constant__ T_F_NEC c_t_f_nec;
169 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
170
171 __device__ T_FA_NEC d_t_fa_nec;
172 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
173 __shared__ T_FA_NEC s_t_fa_nec;
174 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
175 __constant__ T_FA_NEC c_t_fa_nec;
176 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
177
178 __device__ T_B_NED d_t_b_ned;
179 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
180 __shared__ T_B_NED s_t_b_ned;
181 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
182 __constant__ T_B_NED c_t_b_ned;
183 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
184
185 __device__ T_F_NED d_t_f_ned;
186 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
187 __shared__ T_F_NED s_t_f_ned;
188 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
189 __constant__ T_F_NED c_t_f_ned;
190 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
191
192 __device__ T_FA_NED d_t_fa_ned;
193 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
194 __shared__ T_FA_NED s_t_fa_ned;
195 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
196 __constant__ T_FA_NED c_t_fa_ned;
197 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
198
199 // Verify that only __shared__ local variables may be static on device
200 // side and that they are not allowed to be initialized.
df_sema()201 __device__ void df_sema() {
202 static __shared__ NCFS s_ncfs;
203 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
204 static __shared__ UC s_uc;
205 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
206 static __shared__ NED s_ned;
207 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
208
209 static __device__ int ds;
210 // expected-error@-1 {{within a __device__ function, only __shared__ variables or const variables without device memory qualifier may be marked 'static'}}
211 static __constant__ int dc;
212 // expected-error@-1 {{within a __device__ function, only __shared__ variables or const variables without device memory qualifier may be marked 'static'}}
213 static int v;
214 // expected-error@-1 {{within a __device__ function, only __shared__ variables or const variables without device memory qualifier may be marked 'static'}}
215 static const int cv = 1;
216 static const __device__ int cds = 1;
217 // expected-error@-1 {{within a __device__ function, only __shared__ variables or const variables without device memory qualifier may be marked 'static'}}
218 static const __constant__ int cdc = 1;
219 // expected-error@-1 {{within a __device__ function, only __shared__ variables or const variables without device memory qualifier may be marked 'static'}}
220 }
221
hd_sema()222 __host__ __device__ void hd_sema() {
223 static int x = 42;
224 #ifdef __CUDA_ARCH__
225 // expected-error@-2 {{within a __host__ __device__ function, only __shared__ variables or const variables without device memory qualifier may be marked 'static'}}
226 #endif
227 }
228
hd_emitted_host_only()229 inline __host__ __device__ void hd_emitted_host_only() {
230 static int x = 42; // no error on device because this is never codegen'ed there.
231 }
call_hd_emitted_host_only()232 void call_hd_emitted_host_only() { hd_emitted_host_only(); }
233
234 // Verify that we also check field initializers in instantiated structs.
235 struct NontrivialInitializer {
NontrivialInitializerNontrivialInitializer236 __host__ __device__ NontrivialInitializer() : x(43) {}
237 int x;
238 };
239
240 template <typename T>
bar()241 __global__ void bar() {
242 __shared__ T bad;
243 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
244 }
245
instantiate()246 void instantiate() {
247 bar<NontrivialInitializer><<<1, 1>>>();
248 // expected-note@-1 {{in instantiation of function template specialization 'bar<NontrivialInitializer>' requested here}}
249 }
250