cuda contrôle divergence

Question

disons que j'ai 3 matrice de mémoire partagée: a b c. Je ne sais pas si, après accord de fil provoque une divergence de contrôle ou non,

if (threadIdx < 64) 
{ 
    if (threadIdx == 1) 
     for (int i = 0; i < N; i++) 
      c += a[threadIdx]*a[threadIdx]; 
    else 
     for (int i = 0; i < N; i++) 
      c += a[threadIdx]*b[threadIdx]; 
} 

si elle le fait, à quel point est-il affecter la performance va? existe-t-il un moyen efficace de gérer le problème? merci

Answer 1

S'il y a plus d'un fil par bloc, je m'attendrais à une divergence dans une chaîne de chaque bloc (quel que soit le bloc qui contient le fil 1). Mais, la différence entre vos deux boucles est seulement dans quelle mémoire accéder, pas dans les instructions. Donc, je le ferais à la place ...

if (threadIdx.x < 64) 
{ 
    //this conditional might diverge 
    if (threadIdx.x == 1) 
     ptr = a; 
    else 
     ptr = b; 

    //but obviously this part will not 
    for (int i = 0; i < N; i++) 
     c += a[threadIdx]*ptr[threadIdx]; 
} 

Answer 2

En fonction des dimensions de votre bloc, la première condition threadIdx.x < 64 (notez le .x) ne doit pas provoquer de divergence. Par exemple, si vous avez un bloc avec les cotes (128,1,1) alors les deux premiers warps (groupes de 32 threads qui s'exécutent en lock-step) entreront dans le bloc if tandis que les deux derniers le contourneront. Depuis la chaîne entière va dans un sens ou dans l'autre, il n'y a pas de divergence.

Un conditionnel comme threadIdx.x == 1 provoquera une divergence, mais il aura un coût très modeste. En effet, dans de nombreux cas, CUDA pourra implémenter l'expression conditionnelle avec une seule instruction. Par exemple, des opérations telles que min, max et abs seront généralement implémentées avec une seule instruction et ne provoqueront absolument aucune divergence. Vous pouvez lire à propos de ces instructions dans le PTX Manual.

En général, vous ne devriez pas trop vous inquiéter des quantités modestes de divergence d'écoulement de contrôle comme ci-dessus. Où divergence vous mordre dans des situations comme

if (threadIdx.x % 4 == 0) 
    // do expensive operation 
else if (threadIdx.x % 4 == 1) 
    // do expensive operation 
else if (threadIdx.x % 4 == 2) 
    // do expensive operation 
else 
    // do expensive operation 

où une « opération coûteuse » serait celui qui doit 10s ou 100s d'instructions. Dans ce cas, la divergence provoquée par les instructions if réduirait l'efficacité de 75%. Gardez à l'esprit que la divergence des threads est beaucoup moins importante que (1) les choix d'algorithmes de haut niveau et (2) la localisation/coalescence de la mémoire. Très peu de programmeurs CUDA devraient jamais être concernés par le genre de divergence dans vos exemples.