Costruire un istogramma
Utilizzi gli atomici con privatizzazione della memoria condivisa.
Costruire un istogramma è 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.
Che cosa conta un istogramma
Un istogramma conta quanti input rientrano in ciascun intervallo. Molti thread vorranno incrementare lo stesso intervallo, quindi questo è un problema da risolvere con operazioni atomiche. 📊
L'approccio ingenuo
Ogni thread legge un elemento, ne individua l'intervallo e incrementa quell'intervallo. Senza protezione, gli intervalli più popolari perdono conteggi a causa delle race.
La versione con atomic globale
La soluzione corretta più semplice consiste nell'eseguire un atomicAdd per ogni elemento direttamente nella memoria globale. Funziona, ma gli intervalli più contesi serializzano i thread.
atomicAdd(&hist[bin], 1);Il problema della contesa
Quando i dati si concentrano in pochi intervalli, migliaia di thread si accalcano sullo stesso indirizzo. Questa contesa può rendere le operazioni atomiche globali estremamente lente.
La privatizzazione viene in aiuto
La privatizzazione assegna a ogni blocco un istogramma privato nella veloce memoria condivisa. I thread entrano in conflitto solo all'interno del proprio blocco, non sull'intera griglia.
Dichiarare l'istogramma condiviso
Ogni blocco dichiara un array shared grande quanto il numero di intervalli. L'array risiede sul chip, quindi le operazioni atomiche sono molto meno costose di quelle globali.
__shared__ int local[NBINS];Passaggio 1: azzerare gli intervalli
I thread azzerano collaborativamente l'istogramma condiviso, quindi chiamano __syncthreads per impedire che qualcuno inizi a contare prima del termine dell'azzeramento.
local[tid] = 0;
__syncthreads();Passaggio 2: contare localmente
Ora ogni thread incrementa atomicamente il proprio intervallo nella memoria shared. È lo stesso atomicAdd, ma sulla copia veloce sul chip invece che nella memoria globale.
atomicAdd(&local[bin], 1);Passaggio 3: unire i risultati globalmente
Dopo una barriera, i thread aggiungono ogni intervallo condiviso all'istogramma globale con un atomicAdd per intervallo. Le operazioni atomiche globali sono molte meno di prima.
atomicAdd(&hist[i], local[i]);Perché è più veloce
Le operazioni atomiche sulla memoria condivisa sono rapide, mentre le costose operazioni atomiche globali vengono ora eseguite una volta per intervallo per blocco invece che una volta per elemento.
Controllare il numero di intervalli
L'istogramma privato deve rientrare nella memoria condivisa. Se gli intervalli sono troppi, suddivideteli in più passaggi oppure usate nuovamente le operazioni atomiche globali.
Verifica rapida
Una domanda sulla strategia per gli istogrammi.
Riepilogo: costruire un istogramma
Avete costruito un istogramma con operazioni atomiche globali, poi lo avete accelerato usando la privatizzazione nella memoria condivisa: azzerare, contare localmente, unire i risultati. ✅
Domande Frequenti
La lezione «Costruire un istogramma» è gratuita?
Sì — il testo completo di «Costruire un istogramma» è 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 «Costruire un istogramma»?
Utilizzi gli atomici con privatizzazione della memoria condivisa. 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 «Costruire un istogramma»?
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
- Condizioni di gara sulla GPU
- atomicAdd e funzioni correlate
- Costruire un istogramma
- Atomici personalizzati con atomicCAS