__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];
}

No exemplo acima, eu achox, y, offset são salvos em registros enquanto

nvcc -Xptxas -v dá4 registers, 24+16 bytes smem

perfil mostra 4 registros

e o chefe deptx Arquivo:

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

$LDWbegin__Z3addPiPKiS1_:   
.loc    15  26  0  

Alguém pode esclarecer o uso de registros? No Fermi, o número máximo de registros é 63 para cada thread. No meu programa, quero testar o caso quando um kernel consome muitos registradores (portanto, as variáveis ​​podem ter que ser armazenadas na memória local automaticamente e, assim, levar à diminuição do desempenho). Então neste ponto eu posso dividir um kernel em dois para que cada thread tenha registros suficientes. Suponha que os recursos do SM sejam suficientes para kernels simultâneos.

Não tenho certeza se estou certo.