clang/test/SemaCUDA/device-use-host-var.cu

   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 varables in class or file scope is not allowed.
 115   *out = A::host_var; // dev-error {{reference to __host__ variable 'host_var' in __device__ function}}
 116   *out = static_host_var; // dev-error {{reference to __host__ variable 'static_host_var' in __device__ function}}
 117
 118   // Check function-scope static variable is allowed.
 119   static int static_var;
 120   *out = static_var;
 121
 122   // Check non-ODR use of host varirables are allowed.
 123   *out = sizeof(global_host_var);
 124   *out = sizeof(global_host_struct_var.x);
 125   decltype(global_host_var) var1;
 126   decltype(global_host_struct_var.x) var2;
 127 }
 128
 129 __global__ void global_fun(int *out) {
 130   int &ref_host_var = global_host_var; // dev-error {{reference to __host__ variable 'global_host_var' in __global__ function}}
 131   int &ref_dev_var = global_dev_var;
 132   int &ref_constant_var = global_constant_var;
 133   int &ref_shared_var = global_shared_var;
 134   const int &ref_constexpr_var = global_constexpr_var;
 135   const int &ref_const_var = global_const_var;
 136
 137   *out = global_host_var; // dev-error {{reference to __host__ variable 'global_host_var' in __global__ function}}
 138   *out = global_dev_var;
 139   *out = global_constant_var;
 140   *out = global_shared_var;
 141   *out = global_constexpr_var;
 142   *out = global_const_var;
 143
 144   *out = ref_host_var;
 145   *out = ref_dev_var;
 146   *out = ref_constant_var;
 147   *out = ref_shared_var;
 148   *out = ref_constexpr_var;
 149   *out = ref_const_var;
 150 }
 151
 152 __host__ __device__ void host_dev_fun(int *out) {
 153   int &ref_host_var = global_host_var; // dev-error {{reference to __host__ variable 'global_host_var' in __host__ __device__ function}}
 154   int &ref_dev_var = global_dev_var;
 155   int &ref_constant_var = global_constant_var;
 156   int &ref_shared_var = global_shared_var;
 157   const int &ref_constexpr_var = global_constexpr_var;
 158   const int &ref_const_var = global_const_var;
 159
 160   *out = global_host_var; // dev-error {{reference to __host__ variable 'global_host_var' in __host__ __device__ function}}
 161   *out = global_dev_var;
 162   *out = global_constant_var;
 163   *out = global_shared_var;
 164   *out = global_constexpr_var;
 165   *out = global_const_var;
 166
 167   *out = ref_host_var;
 168   *out = ref_dev_var;
 169   *out = ref_constant_var;
 170   *out = ref_shared_var;
 171   *out = ref_constexpr_var;
 172   *out = ref_const_var;
 173 }
 174
 175 inline __host__ __device__ void inline_host_dev_fun(int *out) {
 176   int &ref_host_var = global_host_var;
 177   int &ref_dev_var = global_dev_var;
 178   int &ref_constant_var = global_constant_var;
 179   int &ref_shared_var = global_shared_var;
 180   const int &ref_constexpr_var = global_constexpr_var;
 181   const int &ref_const_var = global_const_var;
 182
 183   *out = global_host_var;
 184   *out = global_dev_var;
 185   *out = global_constant_var;
 186   *out = global_shared_var;
 187   *out = global_constexpr_var;
 188   *out = global_const_var;
 189
 190   *out = ref_host_var;
 191   *out = ref_dev_var;
 192   *out = ref_constant_var;
 193   *out = ref_shared_var;
 194   *out = ref_constexpr_var;
 195   *out = ref_const_var;
 196 }
 197
 198 void dev_lambda_capture_by_ref(int *out) {
 199   int &ref_host_var = global_host_var;
 200   kernel<<<1,1>>>([&]() {
 201   int &ref_dev_var = global_dev_var;
 202   int &ref_constant_var = global_constant_var;
 203   int &ref_shared_var = global_shared_var;
 204   const int &ref_constexpr_var = global_constexpr_var;
 205   const int &ref_const_var = global_const_var;
 206
 207   *out = global_host_var; // dev-error {{reference to __host__ variable 'global_host_var' in __host__ __device__ function}}
 208                           // dev-error@-1 {{capture host variable 'out' by reference in device or host device lambda function}}
 209   *out = global_dev_var;
 210   *out = global_constant_var;
 211   *out = global_shared_var;
 212   *out = global_constexpr_var;
 213   *out = global_const_var;
 214
 215   *out = ref_host_var; // dev-error {{capture host variable 'ref_host_var' by reference in device or host device lambda function}}
 216   *out = ref_dev_var;
 217   *out = ref_constant_var;
 218   *out = ref_shared_var;
 219   *out = ref_constexpr_var;
 220   *out = ref_const_var;
 221   });
 222 }
 223
 224 void dev_lambda_capture_by_copy(int *out) {
 225   int &ref_host_var = global_host_var;
 226   kernel<<<1,1>>>([=]() {
 227   int &ref_dev_var = global_dev_var;
 228   int &ref_constant_var = global_constant_var;
 229   int &ref_shared_var = global_shared_var;
 230   const int &ref_constexpr_var = global_constexpr_var;
 231   const int &ref_const_var = global_const_var;
 232
 233   *out = global_host_var; // dev-error {{reference to __host__ variable 'global_host_var' in __host__ __device__ function}}
 234   *out = global_dev_var;
 235   *out = global_constant_var;
 236   *out = global_shared_var;
 237   *out = global_constexpr_var;
 238   *out = global_const_var;
 239
 240   *out = ref_host_var;
 241   *out = ref_dev_var;
 242   *out = ref_constant_var;
 243   *out = ref_shared_var;
 244   *out = ref_constexpr_var;
 245   *out = ref_const_var;
 246   });
 247 }
 248
 249 // Texture references are special. As far as C++ is concerned they are host
 250 // variables that are referenced from device code. However, they are handled
 251 // very differently by the compiler under the hood and such references are
 252 // allowed. Compiler should produce no warning here, but it should diagnose the
 253 // same case without the device_builtin_texture_type attribute.
 254 template <class, int = 1, int = 1>
 255 struct __attribute__((device_builtin_texture_type)) texture {
 256   static texture<int> ref;
 257   __device__ void c() {
 258     auto &x = ref;
 259   }
 260 };
 261
 262 template <class, int = 1, int = 1>
 263 struct  not_a_texture {
 264   static not_a_texture<int> ref;
 265   __device__ void c() {
 266     auto &x = ref; // dev-error {{reference to __host__ variable 'ref' in __device__ function}}
 267   }
 268 };
 269
 270 template<>
 271 not_a_texture<int> not_a_texture<int>::ref; // dev-note {{host variable declared here}}
 272
 273 __device__ void test_not_a_texture() {
 274   not_a_texture<int> inst;
 275   inst.c(); // dev-note {{in instantiation of member function 'not_a_texture<int>::c' requested here}}
 276 }
 277
 278 // Test static variable in host function used by device function.
 279 void test_static_var_host() {
 280   for (int i = 0; i < 10; i++) {
 281     static int x; // dev-note {{host variable declared here}}
 282     struct A {
 283       __device__ int f() {
 284         return x; // dev-error{{reference to __host__ variable 'x' in __device__ function}}
 285       }
 286     };
 287   }
 288 }
 289
 290 // Test static variable in device function used by device function.
 291 __device__ void test_static_var_device() {
 292   for (int i = 0; i < 10; i++) {
 293     static int x;
 294     int y = x;
 295     struct A {
 296       __device__ int f() {
 297         return x;
 298       }
 299     };
 300   }
 301 }