Riduzione finale su più blocchi

Una riduzione all'interno di un blocco è semplice, ma i blocchi vengono eseguiti indipendentemente e non possono sincronizzarsi tra loro durante l'esecuzione del kernel. Perciò un singolo lancio non può sommare tutto.

Ogni blocco riduce quindi il proprio segmento a un solo numero, una somma parziale, e la scrive in un piccolo array di output indicizzato da blockIdx.

Con 1000 blocchi si passa da un milione di input a 1000 somme parziali. La parte difficile è terminata: resta solo un piccolo array da combinare.

Il modo più semplice per concludere consiste in un secondo lancio dello stesso kernel sulle somme parziali. Ripetete l'operazione finché non rimane un solo valore.

A ogni passaggio l'array si riduce della dimensione del blocco. Alcuni lanci ricorsivi riducono milioni di valori a un'unica somma finale.

In alternativa, il thread 0 di ogni blocco può aggiungere direttamente il proprio parziale a un unico totale globale usando atomicAdd, evitando un secondo kernel.

Le operazioni atomiche sono semplici e richiedono un solo lancio, ma molti blocchi in contesa sullo stesso indirizzo possono essere serializzati. Con pochi parziali, di solito funzionano bene.

Un ciclo grid-stride consente innanzitutto a ogni thread di sommare molti elementi in un registro, così prima del passaggio finale servono molti meno blocchi.

Eseguire più lavoro per thread all'inizio significa avere meno parziali e meno lanci. Spesso è più efficiente che creare un thread per ogni elemento.

Se usate operazioni atomiche, ricordatevi di azzerare il totale globale prima del lancio; altrimenti la somma partirà dai dati residui presenti in quella memoria.

Per gli input piccoli, le operazioni atomiche sono vantaggiose per la loro semplicità; per quelli enormi è preferibile un design a due passaggi o grid-stride. Misurate le prestazioni sui vostri dati per scegliere.

Riflettete sul motivo per cui un singolo lancio del kernel non può sommare direttamente l'intero array.

Ogni blocco produce una somma parziale, che poi combinate con un secondo lancio, con operazioni atomiche oppure con un approccio grid-stride. Ora potete ridurre array di qualsiasi dimensione. 🏁