__shfl_down_sync per le riduzioni
Esegua riduzioni tra warp senza barriere.
__shfl_down_sync per le riduzioni è 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.
Lo shuffle sposta i registri
Uno shuffle permette a una lane di leggere direttamente il valore contenuto nel registro di un'altra lane. I dati passano da un thread all'altro senza memoria condivisa e senza una barriera intermedia.
Conoscere shfl_down_sync
Lo strumento principale per le riduzioni è shfl_down_sync. Ogni lane recupera un valore da una lane situata un numero fisso di posizioni più avanti nel warp.
float v = __shfl_down_sync(mask, val, offset);Leggere gli argomenti
La chiamata riceve una mask attiva, il valore da condividere e un offset. La lane i riceve il valore contenuto nella lane i più offset.
Sommare ciò che si riceve
Una somma del warp consiste semplicemente nel sommare nuovamente il valore ottenuto con lo shuffle. La lane i prende il numero del vicino e lo aggiunge al proprio totale parziale.
val += __shfl_down_sync(mask, val, offset);Dimezzare l'offset
Come in una riduzione ad albero, l'offset parte da 16 e si dimezza a ogni passaggio: 16, 8, 4, 2, 1. Dopo cinque passaggi l'intero warp è stato sommato.
for (int o = 16; o > 0; o >>= 1)
val += __shfl_down_sync(mask, val, o);Perché cinque passaggi
Un warp ha 32 lane e 2 elevato alla quinta è 32. Quindi cinque passaggi di dimezzamento sono esattamente sufficienti per combinare tutti i 32 valori in uno solo.
Il risultato arriva nella lane 0
Dopo il ciclo, il totale completo della warp si trova in lane 0. Le altre lane contengono risultati parziali inutili, quindi solo lane 0 deve scrivere il risultato.
Nessuna barriera necessaria
Poiché lo scambio avviene tramite i registri in lockstep, può evitare del tutto syncthreads. Il suffisso sync coordina già le lane incluse nella maschera.
Più veloce della memoria condivisa
La riduzione tramite warp shuffle è più veloce della versione con memoria condivisa: richiede meno istruzioni, non presenta conflitti di banco e non subisce stalli dovuti alle barrier. È il percorso rapido per le ultime 32 lane.
Passare la maschera corretta
Se alcune lane sono già uscite, usare una maschera completa è errato. Acquisisca l'insieme delle lane attive con activemask, così ogni shuffle coinvolge solo le lane presenti.
Riduzioni a due livelli
Nelle riduzioni di grandi dimensioni, spesso si riduce ogni warp con gli shuffle e poi si combinano i risultati dei warp nella shared memory. Gli shuffle gestiscono egregiamente il livello interno, poco costoso.
Verifica rapida
Rifletta su quale lane contiene il risultato quando il ciclo termina.
Riepilogo
Ha sommato un warp con shfl_down_sync: dimezzi l'offset cinque volte e lane 0 contiene il totale, senza bisogno di barriere. Prossimo argomento: ballot e vote. ✨
Domande Frequenti
La lezione «__shfl_down_sync per le riduzioni» è gratuita?
Sì — il testo completo di «__shfl_down_sync per le riduzioni» è 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 «__shfl_down_sync per le riduzioni»?
Esegua riduzioni tra warp senza barriere. 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 «__shfl_down_sync per le riduzioni»?
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