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

W powyższym przykładzie chybax, y, offset są zapisywane w rejestrach

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

profiler pokazuje 4 rejestry

i głowaptx plik:

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

$LDWbegin__Z3addPiPKiS1_:   
.loc    15  26  0  

Czy każdy może wyjaśnić użycie rejestrów? W Fermi maksymalna liczba rejestrów wynosi 63 dla każdego wątku. W moim programie chcę przetestować przypadek, gdy jądro zużywa zbyt wiele rejestrów (więc zmienne mogą być automatycznie zapisywane w pamięci lokalnej, co prowadzi do zmniejszenia wydajności). W tym momencie mogę podzielić jedno jądro na dwa, aby każdy wątek miał wystarczającą liczbę rejestrów. Załóżmy, że zasoby SM są wystarczające dla współbieżnych jąder.

Nie jestem pewien, czy mam rację.