Eliminare la divergenza dei warp
Reindicizzi per mantenere attivi i warp.
Eliminare la divergenza dei warp è una lezione CUDA Academy gratuita su CoddyKit. Questa è la lezione 2 di 4. Puoi leggere la lezione completa qui gratuitamente — poi esercitati direttamente nel browser con un editor di codice integrato e un tutor IA disponibile 24/7. Fa parte del percorso di apprendimento CUDA Academy, e i tuoi progressi si sincronizzano tra il web e l'app CoddyKit. Il corso CUDA Academy include 4 lezioni in totale.
I warp eseguono in lockstep
Un warp è composto da 32 thread che eseguono insieme la stessa istruzione. Quando i loro percorsi coincidono, l’hardware lavora alla massima velocità.
Il costo della divergenza
Se i thread di un warp seguono rami diversi, si verifica una divergenza. L’hardware esegue ogni percorso in serie, lasciando inattivi alcuni lane e sprecando cicli.
La riduzione ingenua diverge
La versione semplice usa tid % (2*s) per scegliere i thread attivi. I thread attivi e inattivi si alternano all’interno di ogni warp, causando una forte divergenza.
if (tid % (2 * s) == 0)
data[tid] += data[tid + s];Anche i lane inattivi hanno un costo
Anche se metà dei thread non fa nulla, continua a occupare il warp. Il warp non può terminare finché non vengono gestiti sia il percorso attivo sia quello inattivo.
Reindicizzare in base all’ID del thread
La soluzione consiste nell’associare il lavoro attivo agli ID di thread più bassi, anziché distribuirlo in modo sparso. Calcoli un indice a partire da tid e dallo stride.
int index = 2 * s * tid;
if (index < blockDim.x)
data[index] += data[index + s];Perché è utile
Ora i thread attivi sono contigui: tid 0,1,2,... lavorano tutti, mentre gli altri sono tutti inattivi. I warp sono quindi interamente attivi oppure interamente inattivi.
I warp completamente inattivi non costano nulla
Un warp in cui ogni lane è inattivo termina semplicemente senza eseguire lavoro. Non c’è serializzazione per lane, quindi il costo della divergenza scompare in gran parte.
La trappola del modulo
Il responsabile nascosto era la condizione con modulo. Distribuiva i thread attivi in tutto il warp, creando esattamente la divergenza.
Stesso lavoro, associazione migliore
Non ha modificato la matematica né il numero di addizioni. Ha solo rimappato quale thread esegue ciascuna addizione, e i warp ne traggono vantaggio.
Il vantaggio cresce con la scala
Su migliaia di blocchi e numerosi passaggi, eliminare la divergenza produce uno speedup concreto, spesso di un paio di volte rispetto al kernel ingenuo.
Resta ancora un problema
Questa versione legge vicini distribuiti in modo alternato nella memoria condivisa, causando potenzialmente conflitti tra banchi. Anche questo problema verrà risolto nella prossima lezione.
Verifica rapida
Rifletta su cosa causa la divergenza dei warp nella riduzione ingenua.
Riepilogo
Avete eliminato la divergenza assegnando il lavoro agli ID dei thread più bassi, così i warp sono tutti attivi o tutti inattivi. Stessa matematica, riduzione più rapida. Prossimo argomento: i conflitti tra banchi. 🚀
Domande Frequenti
La lezione «Eliminare la divergenza dei warp» è gratuita?
Sì — il testo completo di «Eliminare la divergenza dei warp» è gratuito qui sul web. Per esercitarvi in modo interattivo (un editor di codice integrato e un tutor IA 24/7) e sbloccare il resto del corso CUDA Academy, passa a CoddyKit PRO. Il corso CUDA Academy include 4 lezioni in totale.
Cosa imparerò in «Eliminare la divergenza dei warp»?
Reindicizzi per mantenere attivi i warp. Eserciti CUDA Academy con codice pratico che esegui direttamente nel browser, e un tutor IA 24/7 risponde alle tue domande mentre lavori sulla lezione.
Ho bisogno di esperienza per iniziare CUDA Academy?
Non è richiesta alcuna esperienza precedente. CUDA Academy su CoddyKit è strutturato per principianti e studenti avanzati, quindi puoi iniziare da qui o dall'inizio e procedere al tuo ritmo. Questa è la lezione 2 di 4.
Quanto tempo richiede la lezione «Eliminare la divergenza dei warp»?
La maggior parte delle lezioni CoddyKit richiede circa 5–10 minuti. Ogni lezione è breve e interattiva, quindi fai progressi costanti e riprendi esattamente da dove hai lasciato su web e app.
Posso scrivere ed eseguire codice in questa lezione CUDA Academy?
Sì. Ogni lezione CUDA Academy include un editor di codice integrato, quindi scrivi ed esegui codice reale direttamente nel tuo browser e ricevi feedback istantaneo dall'IA — nessuna configurazione locale necessaria.
Tutte le lezioni di questo corso
- L'idea dell'albero di riduzione
- Eliminare la divergenza dei warp
- Indirizzamento sequenziale
- Riduzione finale su più blocchi