Warp, lane e maschere
L'unità da 32 thread e le maschere degli elementi attivi.
Warp, lane e maschere è una lezione CUDA Academy gratuita su CoddyKit. Questa è la lezione 1 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 thread vengono eseguiti in warp
La GPU non pianifica i thread uno alla volta. Li raggruppa in un warp di 32 thread che procedono insieme ed eseguono la stessa istruzione in lockstep.
Perché 32 è importante
Sull'hardware NVIDIA, un warp è sempre composto da 32 thread. È la vera unità di esecuzione, quindi i kernel efficienti ragionano in gruppi di 32, non in singoli thread.
Ogni thread è una lane
All'interno di un warp, ogni thread ha una posizione da 0 a 31 chiamata lane. L'ID della lane permette alle primitive del warp di sapere chi comunica con chi.
int lane = threadIdx.x % 32;Individuare l'ID della lane
Può leggere direttamente la lane da un registro speciale invece di calcolarla. Il registro laneid contiene sempre un valore da 0 a 31 per il thread corrente.
Il lockstep ha un inconveniente
I thread condividono un program counter per warp. Quando un if indirizza le lane verso percorsi diversi, il warp diverge ed esegue ogni percorso a turno, riducendo le prestazioni.
La maschera attiva
Non tutte le lane sono sempre in esecuzione. Una mask a 32 bit indica quali lane sono attive in quel momento, con un bit per lane impostato a 1 quando partecipa.
Perché esistono le maschere
Dopo una divergenza, solo alcune lane sono attive. Le primitive del warp devono sapere esattamente chi è presente, quindi si passa loro una active mask per mantenerne il corretto funzionamento.
La maschera completa
Quando è certo che tutte le 32 lane siano attive, la mask è 0xffffffff, con tutti i bit impostati. Questa mask completa è il valore passato più spesso.
unsigned mask = 0xffffffff;Costruire una mask in sicurezza
All'interno di un branch, non indovini la mask. Chiami activemask per acquisire esattamente quali lane hanno raggiunto questo punto in quel preciso momento.
unsigned mask = __activemask();Le lane comunicano senza memoria
Il grande vantaggio è che le lane di uno stesso warp possono scambiarsi direttamente i dati tramite i registri. Per questo scambio locale al warp non servono né memoria condivisa né barriere.
Sync significa a livello di warp
Il suffisso sync di questi intrinsic indica una sincronizzazione a livello di warp, non una barriera del blocco. Coordina soltanto le lane indicate nella mask fornita.
Verifica rapida
Ricordi cos'è un warp e qual è la sua dimensione sulle GPU NVIDIA.
Riepilogo
Un warp è formato da 32 lane eseguite in lockstep e una mask tiene traccia di quelle attive. Questa base consente alle lane di condividere rapidamente i dati. Prossimo argomento: gli shuffle per le riduzioni. ✨
Impara C++ con un tutor IA — gratis
Scrivi ed esegui vero codice nel tuo browser, ricevi aiuto istantaneo da un tutor IA disponibile 24/7, e riprendi da dove hai lasciato sul web o nell'app.
- Corsi
- 30
- Lezioni
- 120
Domande Frequenti
La lezione «Warp, lane e maschere» è gratuita?
Sì — il testo completo di «Warp, lane e maschere» è 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 «Warp, lane e maschere»?
L'unità da 32 thread e le maschere degli elementi attivi. 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 1 di 4.
Quanto tempo richiede la lezione «Warp, lane e maschere»?
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
- Warp, lane e maschere
- __shfl_down_sync per le riduzioni
- Funzioni ballot e vote
- Cooperative Groups