CUDA: sincronizzazione dei thread

Question

Quasi ovunque leggo sulla programmazione con CUDA, si fa menzione dell'importanza che tutti i thread in un warp facciano la stessa cosa.
Nel mio codice ho una situazione in cui non posso evitare una determinata condizione. Ecco come si presenta:

// some math code, calculating d1, d2 
if (d1 < 0.5) 
{ 
    buffer[x1] += 1; // buffer is in the global memory 
} 
if (d2 < 0.5) 
{ 
    buffer[x2] += 1; 
} 
// some more math code. 

Alcuni dei fili potrebbe entrare in uno per le condizioni, alcuni potrebbero entrare in entrambi e altri potrebbero non entrare in uno.

Ora per fare in modo che tutto il thread ritorni a "fare la stessa cosa" di nuovo dopo le condizioni, dovrei sincronizzarli dopo le condizioni usando __syncthreads()? O questo in qualche modo accade automagicamente?
È possibile che due thread siano non facendo la stessa cosa a causa del fatto che uno di essi è un'operazione dietro, rovinandolo così per tutti? O c'è qualche sforzo dietro le quinte per convincerli a fare di nuovo la stessa cosa dopo una filiale?

Answer 1

All'interno di un ordito, nessun thread "anticipa" gli altri. Se c'è un ramo condizionale e viene preso da alcuni thread nel warp ma non da altri (aka warp "divergence"), gli altri thread rimarranno inattivi finché il ramo non sarà completo e tutti "convergeranno" di nuovo su un'istruzione comune . Quindi, se hai bisogno solo della sincronizzazione interna dei thread, ciò accade "automaticamente".

Ma i diversi orditi non sono sincronizzati in questo modo. Quindi se il tuo algoritmo richiede che certe operazioni vengano completate su molti warp, allora dovrai usare esplicite chiamate di sincronizzazione (vedi la Guida alla Programmazione CUDA, Sezione 5.4).

---

EDIT: riorganizzato i prossimi paragrafi per chiarire alcune cose.

Qui ci sono due problemi diversi: sincronizzazione delle istruzioni e visibilità della memoria.

• __syncthreads() fa rispettare sincronizzazione istruzione e assicura la visibilità memoria, ma solo all'interno di un blocco, non tutti i blocchi (CUDA Guida di programmazione, appendice B.6). È utile per scrivere e leggere sulla memoria condivisa, ma non è appropriato per sincronizzare l'accesso alla memoria globale.

• __threadfence() assicura la visibilità di memoria globale, ma non fa alcuna sincronizzazione di istruzioni, così nella mia esperienza è di uso limitato (ma si veda il codice di esempio nell'appendice B.5).

• La sincronizzazione delle istruzioni globali non è possibile all'interno di un kernel. Se è necessario eseguire f() su tutti i thread prima di chiamare g() su qualsiasi thread, dividere f() e g() in due diversi kernel e chiamarli in serie dall'host.

• Se è necessario incrementare i contatori globali o condivisi, considerare l'utilizzo della funzione di incremento atomico atomicInc() (Appendice B.10). Nel caso del tuo codice precedente, se x1 e x2 non sono globalmente univoci (attraverso tutti i thread nella griglia), gli incrementi non atomici daranno luogo a una condizione di competizione, simile all'ultimo paragrafo dell'Appendice B.2.4.

Infine, tenere presente che eventuali operazioni sulla memoria globale, e funzioni di sincronizzazione in particolare (compresi Atomics) sono male per le prestazioni.

Senza conoscere il problema che si sta risolvendo è difficile speculare, ma forse è possibile riprogettare il proprio algoritmo per utilizzare la memoria condivisa anziché la memoria globale in alcuni punti. Ciò ridurrà la necessità di sincronizzazione e ti darà un incremento delle prestazioni.

Answer 2

Dalla sezione 6.1 della migliore guida CUDA Practices:

Qualunque istruzione di controllo di flusso (se, interruttore, do, for, while) può influenzare significativamente il throughput istruzione provocando thread dello stesso ordito di divergere; ovvero, seguire diversi percorsi di esecuzione. Se ciò accade, i diversi percorsi di esecuzione devono essere serializzati, aumentando il numero totale di istruzioni eseguite per questo ordito . Una volta completati tutti i diversi percorsi di esecuzione, i thread convergono nello stesso percorso di esecuzione.

Quindi, non è necessario fare nulla di speciale.

Answer 3

La risposta alla tua domanda è no. Non devi fare nulla di speciale. In ogni caso, è possibile risolvere questo problema, anziché il vostro codice è possibile fare qualcosa di simile:

buffer[x1] += (d1 < 0.5); 
buffer[x2] += (d2 < 0.5); 

Si dovrebbe verificare se è possibile utilizzare la memoria condivisa e accesso alla memoria globale in un modello di coalescenza. Assicurarsi inoltre di NON scrivere sullo stesso indice in più di 1 thread.

Answer 4

In risposta di Gabriel:.

"sincronizzazione istruzione globale non è possibile all'interno di un kernel Se avete bisogno di f() fatto su tutte le discussioni prima di chiamare g() su qualsiasi thread, diviso f() eg() in due diversi kernel e chiamarli in serie dall'host. "

E se il motivo per cui è necessario f() e g() nella stessa discussione è perché si sta utilizzando la memoria di registro e si desidera registrare o dati condivisi da f per ottenere g? Cioè, per il mio problema, l'intera ragione per la sincronizzazione tra blocchi è dovuta al fatto che i dati di f sono necessari in g - e scoppiare con un kernel richiederebbe una grande quantità di memoria globale aggiuntiva per trasferire i dati di registro da f a g, che Mi piacerebbe evitare