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 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 }