title: Global device variables in CUDA: bad practice?
body: <p>I am designing a library that has a large contingent of CUDA kernels to perform parallel computations.  All the kernels will be acting on a common object, say a computational grid, which is defined using C++ style objects.  The computational domain doesn't necessarily need to be accessed from the host side, so creating it on the device side and keeping it there makes sense for now.  I'm wondering if the following is considered "good practice": </p>
  
  <ul>
  <li><p>Suppose my computational grid class is called <code>Domain</code>.  First I Define a global device-side variable to store the computational domain:</p>
  
  <pre><code>__device__ Domain* D
  </code></pre></li>
  <li><p>Then I Initialize the computational domain using a CUDA kernel </p>
  
  <pre><code>__global__ void initDomain(paramType P){
      D = new Domain(P);
  }
  </code></pre></li>
  <li><p>Then, I perform computations using this domain with other kernels: </p>
  
  <pre><code>__global__ void doComputation(double *x,double *y){
      D-&gt;doThing(x,y);
      //...
  }
  </code></pre></li>
  </ul>
  
  <p>If my domain remains fixed (i.e. kernels don't modify the domain once it's created), is this OK?  Is there a better way?  I initially tried creating the <code>Domain</code> object on the host side and copying it over to the device, but this turned out to be a hassle because <code>Domain</code> is a relatively complex type that makes it a pain to copy over using e.g. <code>cudaMemCpy</code> or even <code>Thrust::device_new</code> (at least, I couldn't get it to work nicely).</p>