What kind of variables consume registers in CUDA?

__global__ void add( int *c, const int* a, const int* b )
{
    int x = blockIdx.x;
    int y = blockIdx.y;
    int offset = x + y * gridDim.x;
    c[offset] = a[offset] + b[offset];
}

In the above example, I guess x, y, offset are saved in registers while

• nvcc -Xptxas -v gives 4 registers, 24+16 bytes smem

• profiler shows 4 registers

• and the head of ptx file:

  .reg .u16 %rh<4>;
  .reg .u32 %r<9>;    
  .reg .u64 %rd<10>;  
  .loc    15  21  0   
  
  $LDWbegin__Z3addPiPKiS1_:   
  .loc    15  26  0  
  

Can anyone clarify the usage of registers? In Fermi, the maximum number of registers is 63 for each thread. In my program I want to test the case when a kernel consumes too many registers (so variables may have to be stored in local memory automatically and thus leads to performance decrease). Then at this point I can split one kernel into two so that each thread has enough registers. Assume that the SM resources are sufficient for concurrent kernels.

I'm not sure if I am right.