CUDA: sincronizzazione dei thread

Question

Quasi ovunque leggo della programmazione con CUDA si fa menzione dell'importanza che tutti i thread in un ordito facciano la stessa cosa.
Nel mio codice ho una situazione in cui non posso evitare una determinata condizione. Sembra così:

// 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 thread potrebbero entrare in uno per le condizioni, alcuni potrebbero entrare in entrambi e altri potrebbero non entrare in nessuno dei due.

Ora per far tornare tutto il thread a " facendo la stessa cosa " di nuovo dopo le condizioni, dovrei sincronizzarle dopo le condizioni usando __syncthreads () ? O questo succede in qualche modo automagicamente?
Due thread possono non fare la stessa cosa a causa del fatto che uno di essi è un'operazione dietro, rovinandola così per tutti? O c'è qualche sforzo dietro le quinte per indurli a fare di nuovo la stessa cosa dopo un ramo?

Answer 1

All'interno di una curvatura, nessun thread andrà avanti " di tutti gli altri. Se esiste un ramo condizionale ed è preso da alcuni fili nell'ordito ma non da altri (a.k.a. ordito "divergenza"), gli altri fili resteranno inattivi fino a quando il ramo non sarà completo e tutti "convergeranno". di nuovo insieme su un'istruzione comune. Quindi, se hai solo bisogno della sincronizzazione all'interno dei thread di warp, ciò accade "automagicamente".

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

---

EDIT: riorganizzati i prossimi paragrafi per chiarire alcune cose.

Ci sono davvero due diversi problemi qui: sincronizzazione delle istruzioni e visibilità della memoria.

• __syncthreads () applica la sincronizzazione delle istruzioni e garantisce la visibilità della memoria, ma solo all'interno di un blocco, non attraverso i blocchi (Guida alla programmazione CUDA, Appendice B.6). È utile per scrivere-quindi-leggere sulla memoria condivisa, ma non è appropriato per sincronizzare l'accesso alla memoria globale.

• __threadfence () garantisce la visibilità della memoria globale ma non esegue alcuna sincronizzazione delle istruzioni, quindi nella mia esperienza è di utilità limitata (ma vedi il codice di esempio nell'Appendice B.5).

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

• Se hai solo bisogno di incrementare i contatori condivisi o globali, considera l'utilizzo della funzione di incremento atomico atomicInc () (Appendice B.10). Nel caso del codice sopra riportato, se x1 e x2 non sono univoci a livello globale (su tutti i thread della griglia), gli incrementi non atomici comporteranno una condizione di competizione , simile all'ultimo paragrafo dell'appendice B.2.4.

Infine, tieni presente che qualsiasi operazione sulla memoria globale, e in particolare le funzioni di sincronizzazione (inclusa l'atomica) sono dannose per le prestazioni.

Senza conoscere il problema che stai risolvendo, è difficile speculare, ma forse puoi riprogettare il tuo algoritmo per usare la memoria condivisa invece della memoria globale in alcuni punti. Ciò ridurrà la necessità di sincronizzazione e ti darà un aumento delle prestazioni.

Answer 2

Dalla sezione 6.1 della Guida alle migliori pratiche CUDA:

  

Qualsiasi istruzione di controllo del flusso (if, switch, do, for, while) può influire in modo significativo   il throughput delle istruzioni facendo divergere i thread dello stesso ordito; questo è,   seguire percorsi di esecuzione diversi. In questo caso, i diversi percorsi di esecuzione   deve essere serializzato, aumentando il numero totale di istruzioni eseguite per questo   ordito. Quando tutti i diversi percorsi di esecuzione sono stati completati, i thread convergono   tornare allo stesso percorso di esecuzione.

Quindi, non devi fare niente di speciale.

Answer 3

Nella risposta di Gabriel:

" La sincronizzazione globale delle istruzioni non è possibile all'interno di un kernel. Se hai bisogno di f () fatto su tutti i thread prima di chiamare g () su qualsiasi thread, dividi f () e g () in due kernel diversi e chiamali in serie dall'host. & Quot;

Cosa succede se il motivo per cui hai bisogno di f () e g () nello stesso thread è perché stai usando la memoria del registro e vuoi che i dati registrati o condivisi da f arrivino a g? Questo è, per il mio problema, l'intero motivo della sincronizzazione tra i blocchi è perché i dati di f sono necessari in g - e la distribuzione in un kernel richiederebbe una grande quantità di memoria globale aggiuntiva per trasferire i dati del registro da f a g, che I vorrei evitare

Answer 4

La risposta alla tua domanda è no. Non devi fare nulla di speciale. Ad ogni modo, puoi risolvere questo problema, invece del tuo codice puoi fare qualcosa del genere:

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

È necessario verificare se è possibile utilizzare la memoria condivisa e accedere alla memoria globale in uno schema a coalescenza. Inoltre, assicurati di NON voler scrivere nello stesso indice in più di 1 thread.