pression de registre cuda

Question

J'ai un noyau qui a un ajustement linéaire des moindres carrés. Il s'avère que les threads utilisent trop de registres, par conséquent, l'occupation est faible. Voici le noyau,

__global__ 
void strainAxialKernel(
    float* d_dis, 
    float* d_str 
){ 
    int i = threadIdx.x; 
    float a = 0; 
    float c = 0; 
    float e = 0; 
    float f = 0; 
    int shift = (int)((float)(i*NEIGHBOURS)/(float)WINDOW_PER_LINE); 
    int j; 
    __shared__ float dis[WINDOW_PER_LINE]; 
    __shared__ float str[WINDOW_PER_LINE]; 

    // fetch data from global memory 
    dis[i] = d_dis[blockIdx.x*WINDOW_PER_LINE+i]; 
    __syncthreads(); 

    // least square fit 
    for (j=-shift; j<NEIGHBOURS-shift; j++)          
    {                   
     a += j;                 
     c += j*j;                
     e += dis[i+j];               
     f += (float(j))*dis[i+j];            
    }                  
    str[i] = AMP*(a*e-NEIGHBOURS*f)/(a*a-NEIGHBOURS*c)/(float)BLOCK_SPACING;  

    // compensate attenuation 
    if (COMPEN_EXP>0 && COMPEN_BASE>0)           
    {                   
     str[i]                 
     = (float)(str[i]*pow((float)i/(float)COMPEN_BASE+1.0f,COMPEN_EXP));  
    } 

    // write back to global memory 
    if (!SIGN_PRESERVE && str[i]<0)            
    {                   
     d_str[blockIdx.x*WINDOW_PER_LINE+i] = -str[i];       
    }                   
    else                   
    {                   
     d_str[blockIdx.x*WINDOW_PER_LINE+i] = str[i];       
    } 
} 

J'ai 32x404 blocs avec 96 threads dans chaque bloc. Sur le GTS 250, le SM doit pouvoir gérer 8 blocs. Pourtant, visual profileer montre que j'ai 11 registres par thread, en conséquence, l'occupation est de 0,625 (5 blocs par SM). BTW, la mémoire partagée utilisée par chaque bloc est 792 B, de sorte que le registre est le problème. La performance n'est pas la fin du monde. Je suis simplement curieux de savoir si je peux contourner cela. Merci.

Answer 1

Il existe toujours un compromis entre les registres rapides/limités, la mémoire partagée et la mémoire globale lente mais volumineuse. Il n'y a aucun moyen de «contourner» ce compromis. Si vous utilisez l'utilisation du registre réduit en utilisant la mémoire globale, vous devriez obtenir une occupation plus élevée mais un accès mémoire plus lent.

Cela dit, voici quelques idées pour utiliser moins de registres:

1. peuvent modifier être précalculées et stockées dans la mémoire constante? Ensuite, chaque thread doit juste rechercher shift [i].
2. Est-ce que a et c doivent être flottants?
3. Ou, a et c peuvent-ils être supprimés de la boucle et calculés une seule fois? Et donc complètement supprimé?

un est calculé comme une séquence arithmétique simple, donc le réduire ... (quelque chose comme ça)

a = ((NEIGHBORS-shift) - (-shift) + 1) * ((NEIGHBORS-shift) + (-shift))/2 

ou

a = (NEIGHBORS + 1) * ((NEIGHBORS - 2*shift))/2 

donc au lieu, faire quelque chose comme ce qui suit (vous pouvez probablement réduire ces expressions plus loin):

str[i] = AMP*((NEIGHBORS + 1) * ((NEIGHBORS - 2*shift))/2*e-NEIGHBOURS*f) 
str[i] /= ((NEIGHBORS + 1) * ((NEIGHBORS - 2*shift))/2*(NEIGHBORS + 1) * ((NEIGHBORS - 2*shift))/2-NEIGHBOURS*c) 
str[i] /= (float)BLOCK_SPACING; 

Answer 2

L'occupation n'est pas un problème. Le SM dans GTS 250 (capacité de calcul 1.1) peut contenir simultanément 8 blocs (8x96 threads) dans ses registres, mais il n'a que 8 unités d'exécution, ce qui signifie que seulement 8 sur 8x96 (ou, dans votre cas, 5x96) les discussions avanceraient à n'importe quel moment donné. Il y a très peu de valeur à essayer de presser plus de blocs sur le SM surchargé. En fait, vous pouvez essayer de jouer avec l'option -maxrregcount pour AUGMENTER le nombre de registres, ce qui pourrait avoir un effet positif sur les performances.

Answer 3

Vous pouvez utiliser les limites de lancement pour demander au compilateur de générer un mappage de registre pour un nombre maximal de threads et un nombre minimum de blocs par multiprocesseur. Cela peut réduire le nombre de registres afin que vous puissiez atteindre l'occupation souhaitée. Pour votre cas, le calculateur d'occupation de Nvidia montre une occupation maximale théorique de 63%, ce qui semble être ce que vous réalisez. Cela est dû à votre nombre d'enregistrements, comme vous l'avez mentionné, mais cela est également dû au nombre de threads par bloc. L'augmentation du nombre de threads par bloc à 128 et la diminution du nombre de registres à 10 donnent 100% d'occupation maximale théorique.

Pour contrôler les limites de lancement pour votre noyau:

__global__ void 
__launch_bounds__(128, 6) 
MyKernel(...) 
{ 
    ... 
} 

Ensuite, il suffit de lancer avec une taille de bloc de 128 threads et profiter de votre occupation. Le compilateur devrait générer votre noyau de telle sorte qu'il utilise 10 registres ou moins.