title: Forcing CUDA to use register for a variable
body: <p>I have many unused registers in my kernel.  I'd like to tell CUDA to use a few registers to hold some data, rather than doing a global data read every time I need it.  (I'm not able to use shared mem.)</p>
  
  <pre><code>__global__ void simple(float *gData) {
  float rData[1024];
  for(int i=0; i&lt;1024; i++) {
    rData[i]=gData[i];
    }
  // work on the data here
  }
  </code></pre>
  
  <p>compile w/: nvcc -arch sm_20 --ptxas-options=-v simple.cu, and I get<br>
    0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads<br>
    Used 2 registers, 40 bytes cmem[0]</p>
  
  <pre><code>__global__ void simple(float *gData) {
  register float rData[1024];
  for(int i=0; i&lt;1024; i++) {
    rData[i]=gData[i];
    }
  // work on the data here
  }
  </code></pre>
  
  <p><strong>register</strong> declaration does nothing.<br>
    0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads<br>
    Used 2 registers, 40 bytes cmem[0]</p>
  
  <pre><code>__global__ void simple(float *gData) {
  volatile float rData[1024];
  for(int i=0; i&lt;1024; i++) {
    rData[i]=gData[i];
    }
  // work on the data here
  }
  </code></pre>
  
  <p><strong>volatile</strong> declaration creates stack storage:<br>
    4096 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads<br>
    Used 21 resisters, 40 bytes cmem[0]  </p>
  
  <p>1) Is there an easy way to tell the compiler to use register space for a variable?<br>
  2) Where is 'stack frame': register, global mem, local mem,...?  What is a stack frame?  (Since when does the GPU have a stack?  A virtual stack?)<br>
  3) The simple.ptx file is basically empty: (nvcc -arch sm_20 -ptx simple.cu)</p>
  
  <pre><code>.loc 2 14 2
  ret;
  </code></pre>
  
  <p>Any idea where I can find the real machine/compiled code?</p>