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__/__global__ function, only __shared__ variables may be marked "static"}} 211 static __constant__ int dc; 212 // expected-error@-1 {{Within a __device__/__global__ function, only __shared__ variables may be marked "static"}} 213 static int v; 214 // expected-error@-1 {{Within a __device__/__global__ function, only __shared__ variables may be marked "static"}} 215 } 216
