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

Im obigen Beispiel denke ichx, y, offset werden in Registern gespeichert, während

nvcc -Xptxas -v gibt4 registers, 24+16 bytes smem

Profiler zeigt 4 Register

und der Kopf vonptx Datei:

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

$LDWbegin__Z3addPiPKiS1_:   
.loc    15  26  0  

Kann jemand die Verwendung von Registern klären? In Fermi beträgt die maximale Anzahl von Registern 63 für jeden Thread. In meinem Programm möchte ich den Fall testen, in dem ein Kernel zu viele Register belegt (daher müssen Variablen möglicherweise automatisch im lokalen Speicher abgelegt werden, was zu Leistungseinbußen führt). Dann kann ich an dieser Stelle einen Kernel in zwei aufteilen, so dass jeder Thread über genügend Register verfügt. Angenommen, die SM-Ressourcen sind für gleichzeitige Kernel ausreichend.

Ich bin mir nicht sicher, ob ich recht habe.