title: Access CUDA global device variable from host
body: <p>I was wondering if there is an official source, why the following works:</p>
  
  <pre><code>#include &lt;iostream&gt;
  
  struct Array{
      int el[10000];
  };
  
  __device__ Array devAr;
  
  void test(Array&amp; ar = devAr){
      for(int i=0; i&lt;10000; i++)
          ar.el[i] = i;
      std::cout &lt;&lt; ar.el[0] + ar.el[9999] &lt;&lt; std::endl;
  }
  
  int main(){
      test();
  }
  </code></pre>
  
  <p>You get a warning "a __device__ variable "devAr" cannot be directly read in a host function" if you try to access devAr directly but through the reference there is no such warning (for good reason). But in both cases it is possible to access the variable from the host. So it seems, there is a host instance of that variable.</p>
  
  <p>What I need to know: Can I take this for granted?</p>
  
  <p>Other testcase showing the values of the pointers:</p>
  
  <pre><code>#include &lt;iostream&gt;
  #include &lt;cstdio&gt;
  
  __device__ int devAr[2];
  
  __global__ void foo(){
      printf("Device: %p\n", &amp;devAr);
      devAr[0] = 1337;
  }
  
  int main()
  {
      devAr[0] = 4;
      std::cout &lt;&lt; devAr[0] &lt;&lt; std::endl;
      void* ad;
      cudaGetSymbolAddress(&amp;ad, devAr);
      std::cout &lt;&lt; ad &lt;&lt; " " &lt;&lt; &amp;devAr &lt;&lt; std::endl;
      foo&lt;&lt;&lt;1,1&gt;&gt;&gt;();
      cudaDeviceSynchronize();
      int arHost[2];
      cudaMemcpyFromSymbol(arHost, devAr, sizeof(arHost), 0);
      std::cout &lt;&lt; "values: " &lt;&lt; arHost[0] &lt;&lt; std::endl;
  }
  </code></pre>
  
  <p>Output:</p>
  
  <blockquote>
    <p>4<br>
    0x500bc0000 0x66153c<br>
    Device: 0x500bc0000<br>
    values: 1337   </p>
  </blockquote>