1 // RUN: %clang_cc1 -std=c++11 -fcuda-is-device -fsyntax-only -verify=dev %s 2 // RUN: %clang_cc1 -std=c++11 -fsyntax-only -verify=host %s 3 4 // host-no-diagnostics 5 6 #include "Inputs/cuda.h" 7 8 int func(); 9 10 struct A { 11 int x; 12 static int host_var; 13 }; 14 15 int A::host_var; // dev-note {{host variable declared here}} 16 17 namespace X { 18 int host_var; // dev-note {{host variable declared here}} 19 } 20 21 // struct with non-empty ctor. 22 struct B1 { 23 int x; 24 B1() { x = 1; } 25 }; 26 27 // struct with non-empty dtor. 28 struct B2 { 29 int x; 30 B2() {} 31 ~B2() { x = 0; } 32 }; 33 34 static int static_host_var; // dev-note {{host variable declared here}} 35 36 __device__ int global_dev_var; 37 __constant__ int global_constant_var; 38 __shared__ int global_shared_var; 39 40 int global_host_var; // dev-note 8{{host variable declared here}} 41 const int global_const_var = 1; 42 constexpr int global_constexpr_var = 1; 43 44 int global_host_array[2] = {1, 2}; // dev-note {{host variable declared here}} 45 const int global_const_array[2] = {1, 2}; 46 constexpr int global_constexpr_array[2] = {1, 2}; 47 48 A global_host_struct_var{1}; // dev-note 2{{host variable declared here}} 49 const A global_const_struct_var{1}; 50 constexpr A global_constexpr_struct_var{1}; 51 52 // Check const host var initialized with non-empty ctor is not allowed in 53 // device function. 54 const B1 b1; // dev-note {{const variable cannot be emitted on device side due to dynamic initialization}} 55 56 // Check const host var having non-empty dtor is not allowed in device function. 57 const B2 b2; // dev-note {{const variable cannot be emitted on device side due to dynamic initialization}} 58 59 // Check const host var initialized by non-constant initializer is not allowed 60 // in device function. 61 const int b3 = func(); // dev-note {{const variable cannot be emitted on device side due to dynamic initialization}} 62 63 template<typename F> 64 __global__ void kernel(F f) { f(); } // dev-note2 {{called by 'kernel<(lambda}} 65 66 __device__ void dev_fun(int *out) { 67 // Check access device variables are allowed. 68 int &ref_dev_var = global_dev_var; 69 int &ref_constant_var = global_constant_var; 70 int &ref_shared_var = global_shared_var; 71 *out = ref_dev_var; 72 *out = ref_constant_var; 73 *out = ref_shared_var; 74 *out = global_dev_var; 75 *out = global_constant_var; 76 *out = global_shared_var; 77 78 // Check access of non-const host variables are not allowed. 79 *out = global_host_var; // dev-error {{reference to __host__ variable 'global_host_var' in __device__ function}} 80 *out = global_const_var; 81 *out = global_constexpr_var; 82 *out = b1.x; // dev-error {{reference to __host__ variable 'b1' in __device__ function}} 83 *out = b2.x; // dev-error {{reference to __host__ variable 'b2' in __device__ function}} 84 *out = b3; // dev-error {{reference to __host__ variable 'b3' in __device__ function}} 85 global_host_var = 1; // dev-error {{reference to __host__ variable 'global_host_var' in __device__ function}} 86 87 // Check reference of non-constexpr host variables are not allowed. 88 int &ref_host_var = global_host_var; // dev-error {{reference to __host__ variable 'global_host_var' in __device__ function}} 89 const int &ref_const_var = global_const_var; 90 const int &ref_constexpr_var = global_constexpr_var; 91 *out = ref_host_var; 92 *out = ref_constexpr_var; 93 *out = ref_const_var; 94 95 // Check access member of non-constexpr struct type host variable is not allowed. 96 *out = global_host_struct_var.x; // dev-error {{reference to __host__ variable 'global_host_struct_var' in __device__ function}} 97 *out = global_const_struct_var.x; 98 *out = global_constexpr_struct_var.x; 99 global_host_struct_var.x = 1; // dev-error {{reference to __host__ variable 'global_host_struct_var' in __device__ function}} 100 101 // Check address taking of non-constexpr host variables is not allowed. 102 int *p = &global_host_var; // dev-error {{reference to __host__ variable 'global_host_var' in __device__ function}} 103 const int *cp = &global_const_var; 104 const int *cp2 = &global_constexpr_var; 105 106 // Check access elements of non-constexpr host array is not allowed. 107 *out = global_host_array[1]; // dev-error {{reference to __host__ variable 'global_host_array' in __device__ function}} 108 *out = global_const_array[1]; 109 *out = global_constexpr_array[1]; 110 111 // Check ODR-use of host variables in namespace is not allowed. 112 *out = X::host_var; // dev-error {{reference to __host__ variable 'host_var' in __device__ function}} 113 114 // Check ODR-use of static host variables in class or file scope is not 115 // allowed. 116 *out = A::host_var; // dev-error {{reference to __host__ variable 'host_var' in __device__ function}} 117 *out = static_host_var; // dev-error {{reference to __host__ variable 'static_host_var' in __device__ function}} 118 119 // Check function-scope static variable is allowed. 120 static int static_var; 121 *out = static_var; 122 123 // Check non-ODR use of host varirables are allowed. 124 *out = sizeof(global_host_var); 125 *out = sizeof(global_host_struct_var.x); 126 decltype(global_host_var) var1; 127 decltype(global_host_struct_var.x) var2; 128 } 129 130 __global__ void global_fun(int *out) { 131 int &ref_host_var = global_host_var; // dev-error {{reference to __host__ variable 'global_host_var' in __global__ function}} 132 int &ref_dev_var = global_dev_var; 133 int &ref_constant_var = global_constant_var; 134 int &ref_shared_var = global_shared_var; 135 const int &ref_constexpr_var = global_constexpr_var; 136 const int &ref_const_var = global_const_var; 137 138 *out = global_host_var; // dev-error {{reference to __host__ variable 'global_host_var' in __global__ function}} 139 *out = global_dev_var; 140 *out = global_constant_var; 141 *out = global_shared_var; 142 *out = global_constexpr_var; 143 *out = global_const_var; 144 145 *out = ref_host_var; 146 *out = ref_dev_var; 147 *out = ref_constant_var; 148 *out = ref_shared_var; 149 *out = ref_constexpr_var; 150 *out = ref_const_var; 151 } 152 153 __host__ __device__ void host_dev_fun(int *out) { 154 int &ref_host_var = global_host_var; // dev-error {{reference to __host__ variable 'global_host_var' in __host__ __device__ function}} 155 int &ref_dev_var = global_dev_var; 156 int &ref_constant_var = global_constant_var; 157 int &ref_shared_var = global_shared_var; 158 const int &ref_constexpr_var = global_constexpr_var; 159 const int &ref_const_var = global_const_var; 160 161 *out = global_host_var; // dev-error {{reference to __host__ variable 'global_host_var' in __host__ __device__ function}} 162 *out = global_dev_var; 163 *out = global_constant_var; 164 *out = global_shared_var; 165 *out = global_constexpr_var; 166 *out = global_const_var; 167 168 *out = ref_host_var; 169 *out = ref_dev_var; 170 *out = ref_constant_var; 171 *out = ref_shared_var; 172 *out = ref_constexpr_var; 173 *out = ref_const_var; 174 } 175 176 inline __host__ __device__ void inline_host_dev_fun(int *out) { 177 int &ref_host_var = global_host_var; 178 int &ref_dev_var = global_dev_var; 179 int &ref_constant_var = global_constant_var; 180 int &ref_shared_var = global_shared_var; 181 const int &ref_constexpr_var = global_constexpr_var; 182 const int &ref_const_var = global_const_var; 183 184 *out = global_host_var; 185 *out = global_dev_var; 186 *out = global_constant_var; 187 *out = global_shared_var; 188 *out = global_constexpr_var; 189 *out = global_const_var; 190 191 *out = ref_host_var; 192 *out = ref_dev_var; 193 *out = ref_constant_var; 194 *out = ref_shared_var; 195 *out = ref_constexpr_var; 196 *out = ref_const_var; 197 } 198 199 void dev_lambda_capture_by_ref(int *out) { 200 int &ref_host_var = global_host_var; 201 kernel<<<1,1>>>([&]() { 202 int &ref_dev_var = global_dev_var; 203 int &ref_constant_var = global_constant_var; 204 int &ref_shared_var = global_shared_var; 205 const int &ref_constexpr_var = global_constexpr_var; 206 const int &ref_const_var = global_const_var; 207 208 *out = global_host_var; // dev-error {{reference to __host__ variable 'global_host_var' in __host__ __device__ function}} 209 // dev-error@-1 {{capture host variable 'out' by reference in device or host device lambda function}} 210 *out = global_dev_var; 211 *out = global_constant_var; 212 *out = global_shared_var; 213 *out = global_constexpr_var; 214 *out = global_const_var; 215 216 *out = ref_host_var; // dev-error {{capture host variable 'ref_host_var' by reference in device or host device lambda function}} 217 *out = ref_dev_var; 218 *out = ref_constant_var; 219 *out = ref_shared_var; 220 *out = ref_constexpr_var; 221 *out = ref_const_var; 222 }); 223 } 224 225 void dev_lambda_capture_by_copy(int *out) { 226 int &ref_host_var = global_host_var; 227 kernel<<<1,1>>>([=]() { 228 int &ref_dev_var = global_dev_var; 229 int &ref_constant_var = global_constant_var; 230 int &ref_shared_var = global_shared_var; 231 const int &ref_constexpr_var = global_constexpr_var; 232 const int &ref_const_var = global_const_var; 233 234 *out = global_host_var; // dev-error {{reference to __host__ variable 'global_host_var' in __host__ __device__ function}} 235 *out = global_dev_var; 236 *out = global_constant_var; 237 *out = global_shared_var; 238 *out = global_constexpr_var; 239 *out = global_const_var; 240 241 *out = ref_host_var; 242 *out = ref_dev_var; 243 *out = ref_constant_var; 244 *out = ref_shared_var; 245 *out = ref_constexpr_var; 246 *out = ref_const_var; 247 }); 248 } 249 250 // Texture references are special. As far as C++ is concerned they are host 251 // variables that are referenced from device code. However, they are handled 252 // very differently by the compiler under the hood and such references are 253 // allowed. Compiler should produce no warning here, but it should diagnose the 254 // same case without the device_builtin_texture_type attribute. 255 template <class, int = 1, int = 1> 256 struct __attribute__((device_builtin_texture_type)) texture { 257 static texture<int> ref; 258 __device__ void c() { 259 auto &x = ref; 260 } 261 }; 262 263 template <class, int = 1, int = 1> 264 struct not_a_texture { 265 static not_a_texture<int> ref; 266 __device__ void c() { 267 auto &x = ref; // dev-error {{reference to __host__ variable 'ref' in __device__ function}} 268 } 269 }; 270 271 template<> 272 not_a_texture<int> not_a_texture<int>::ref; // dev-note {{host variable declared here}} 273 274 __device__ void test_not_a_texture() { 275 not_a_texture<int> inst; 276 inst.c(); // dev-note {{in instantiation of member function 'not_a_texture<int>::c' requested here}} 277 } 278 279 // Test static variable in host function used by device function. 280 void test_static_var_host() { 281 for (int i = 0; i < 10; i++) { 282 static int x; // dev-note {{host variable declared here}} 283 struct A { 284 __device__ int f() { 285 return x; // dev-error{{reference to __host__ variable 'x' in __device__ function}} 286 } 287 }; 288 } 289 } 290 291 // Test static variable in device function used by device function. 292 __device__ void test_static_var_device() { 293 for (int i = 0; i < 10; i++) { 294 static int x; 295 int y = x; 296 struct A { 297 __device__ int f() { 298 return x; 299 } 300 }; 301 } 302 } 303