Clang Project

clang_source_code/test/SemaCUDA/device-var-init.cu
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.
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
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
229inline __host__ __device__ void hd_emitted_host_only() {
230  static int x = 42; // no error on device because this is never codegen'ed there.
231}
232void call_hd_emitted_host_only() { hd_emitted_host_only(); }
233
234// Verify that we also check field initializers in instantiated structs.
235struct NontrivialInitializer {
236  __host__ __device__ NontrivialInitializer() : x(43) {}
237  int x;
238};
239
240template <typename T>
241__global__ void bar() {
242  __shared__ T bad;
243// expected-error@-1 {{initialization is not supported for __shared__ variables.}}
244}
245
246void instantiate() {
247  bar<NontrivialInitializer><<<1, 1>>>();
248// expected-note@-1 {{in instantiation of function template specialization 'bar<NontrivialInitializer>' requested here}}
249}
250
CodeThinker, Enjoy your reading code