performance cuda - La divergenza delle filiali è davvero così grave?




samples make (2)

Ho visto molte domande sparse su Internet sulla divergenza delle filiali e su come evitarlo . Tuttavia, anche dopo aver letto dozzine di articoli su come funziona CUDA, non riesco a vedere come evitare la divergenza delle filiali nella maggior parte dei casi . Prima che qualcuno mi salti addosso con le unghie aperte, permettimi di descrivere ciò che considero essere "la maggior parte dei casi".

Mi sembra che la maggior parte delle istanze di divergenza delle filiali coinvolga un numero di blocchi di codice veramente distinti. Ad esempio, abbiamo il seguente scenario:

if (A):
  foo(A)
else:
  bar(B)

Se abbiamo due thread che incontrano questa divergenza, il thread 1 verrà eseguito per primo, prendendo il percorso A. Successivamente, il thread 2 prenderà il percorso B. Per rimuovere la divergenza, potremmo cambiare il blocco sopra in modo da leggere in questo modo:

foo(A)
bar(B)

Supponendo che sia sicuro chiamare foo(A) sul thread 2 e bar(B) sul thread 1, ci si potrebbe aspettare che le prestazioni migliorino. Tuttavia, ecco come vedo io:

Nel primo caso, i thread 1 e 2 vengono eseguiti in serie. Chiama questi due cicli di clock.

Nel secondo caso, i thread 1 e 2 eseguono foo(A) in parallelo, quindi eseguono la bar(B) in parallelo. Questo mi sembra ancora come due cicli di clock, la differenza è che nel primo caso, se foo(A) implica una lettura dalla memoria, immagino che il thread 2 possa iniziare l'esecuzione durante quella latenza, il che si traduce in latenza nascosta. Se questo è il caso, il codice divergente del ramo è più veloce.


Answers

Si presume (almeno è l'esempio che si dà e l'unico riferimento che si fa) che l'unico modo per evitare la divergenza delle diramazioni è consentire a tutti i thread di eseguire tutto il codice.

In tal caso sono d'accordo non c'è molta differenza.

Ma evitare la divergenza delle diramazioni probabilmente ha più a che fare con la riorganizzazione dell'algoritmo ad un livello superiore rispetto alla sola aggiunta o rimozione di alcune istruzioni if ​​e rendendo il codice "sicuro" da eseguire in tutti i thread.

Offrirò un esempio. Supponiamo che sappia che i thread dispari dovranno gestire il componente blu di un pixel e anche i thread dovranno gestire il componente verde:

#define N 2 // number of pixel components
#define BLUE 0
#define GREEN 1
// pixel order: px0BL px0GR px1BL px1GR ...


if (threadIdx.x & 1)  foo(pixel(N*threadIdx.x+BLUE));
else                  bar(pixel(N*threadIdx.x+GREEN));

Ciò significa che ogni thread alternativo sta seguendo un determinato percorso, che sia foo o bar . Quindi ora il mio warp impiega il doppio del tempo per essere eseguito.

Tuttavia, se riorganizzo i dati dei pixel in modo che i componenti del colore siano contigui, forse in porzioni di 32 pixel: BL0 BL1 BL2 ... GR0 GR1 GR2 ...

Posso scrivere un codice simile:

if (threadIdx.x & 32)  foo(pixel(threadIdx.x));
else                   bar(pixel(threadIdx.x));

Sembra ancora che io abbia la possibilità di divergenze. Ma dal momento che la divergenza avviene sui limiti di curvatura, un warp di tipo esegue il percorso if o il percorso else , quindi non si verifica alcuna divergenza effettiva.

Questo è un esempio banale, e probabilmente stupido, ma illustra come ci possano essere modi per aggirare la divergenza del warp che non implicano l'esecuzione di tutto il codice di tutti i percorsi divergenti.


Aggiungilo a PS1 usando Mac:

PS1='\[email protected]\u >`[ -d .git ] && git branch | grep  ^*|cut -d" " -f2`> $ '

Prima di eseguire il comando sopra:

Dopo aver eseguito questo comando:

Non preoccuparti, se non è un repository GIT, non visualizzerà l'errore a causa di [-d .git] che controlla se la cartella .git esiste o meno.





performance cuda branch