Funzioni ballot e vote
Valuti predicati tra i thread di un warp.
Funzioni ballot e vote è una lezione CUDA Academy gratuita su CoddyKit. Questa è la lezione 3 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.
Le lane possono votare
Oltre a spostare dati, un warp può rispondere insieme a domande con risposta sì o no. Le intrinsic Vote consentono a ogni lane di condividere un predicato vero o falso in un'unica operazione poco costosa.
Conoscere ballot_sync
Il voto più flessibile è ballot_sync. Restituisce un numero a 32 bit in cui il bit i vale 1 esattamente quando il predicato della lane i è vero.
unsigned bits = __ballot_sync(mask, pred);Un bit per lane
L'intero restituito è una piccola bitmap del warp. Leggendo il bit i, si scopre se la lane i ha votato vero: tutto è contenuto in un unico registro.
Contare i valori veri
Vuole sapere quante lane hanno risposto sì? Calcoli il pop-count del risultato del ballot. popc conta i bit impostati e fornisce immediatamente il numero di voti veri nel warp.
int yes = __popc(__ballot_sync(mask, pred));Tutte o almeno una
Due scorciatoie rispondono alle domande più comuni. all_sync restituisce true solo se tutte le lane concordano, mentre any_sync restituisce true se lo fa almeno una lane.
bool every = __all_sync(mask, pred);Any significa almeno una
Usi any_sync per le uscite anticipate: se una qualsiasi lane trova una corrispondenza, l'intero warp può reagire insieme senza esaminare le lane una alla volta.
bool found = __any_sync(mask, pred);Compattazione del flusso
Il ballot abilita la compaction: ogni lane decide se conservare il proprio elemento, vota e la bitmap indica alle lane dove scrivere in un array di output compattato.
Trovare la propria posizione
Per collocare il proprio valore, una lane conta i bit veri al di sotto della propria posizione. Questo prefix count ottenuto dal ballot è il suo indice nel risultato compattato.
int slot = __popc(bits & ((1u << lane) - 1));Scegliere una lane leader
A volte una sola lane deve eseguire un'operazione condivisa. Il bit impostato meno significativo del ballot, trovato con ffs, è un leader naturale su cui l'intero warp può concordare.
Passare sempre una maschera
Come gli shuffle, le intrinsic di voto accettano una mask attiva. Le lane inattive contribuiscono con un bit 0, quindi passi l'insieme corretto per evitare di contare lane assenti.
Coordinamento economico
I voti costano una sola istruzione e non accedono alla memoria. Consentono a un warp di coordinate le decisioni a un costo molto inferiore rispetto agli atomic o ai flag nella memoria condivisa.
Verifica rapida
Ricordi che cosa restituisce ballot_sync a ogni lane.
Riepilogo
Ha interrogato un warp con ballot e vote: contare i valori veri, verificare se tutti o almeno uno hanno votato sì e compattare i flussi a basso costo. Prossimo argomento: cooperative groups per una sincronizzazione flessibile. ✨
Domande Frequenti
La lezione «Funzioni ballot e vote» è gratuita?
Sì — il testo completo di «Funzioni ballot e vote» è 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 «Funzioni ballot e vote»?
Valuti predicati tra i thread di un 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 3 di 4.
Quanto tempo richiede la lezione «Funzioni ballot e vote»?
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