La formula classica dell'indice

Lo scopo di un kernel CUDA è fare in modo che ogni thread elabori un elemento di dati. Per farlo, ogni thread ha bisogno di un indice globale univoco.

All'interno di un block, threadIdx.x conta solo da 0 fino a blockDim meno uno. Molti block riutilizzano gli stessi numeri piccoli, quindi questo non può essere l'indice finale.

Immagini la grid come una serie di block disposti uno dopo l'altro. Ogni block possiede una porzione contigua dell'array e blockIdx.x indica in quale porzione ci si trova.

blockDim.x è il numero di thread per block. Rappresenta la larghezza di ogni porzione, quindi scala l'offset del block fino alla posizione corretta.

Combini i tre valori: salti i block precedenti, poi aggiungi la tua posizione all'interno del block corrente. Questa singola riga fornisce a ogni thread un indice univoco.

Con 4 thread per block, il block 0 copre da 0 a 3, il block 1 da 4 a 7 e il block 2 da 8 a 11. Gli offset non si sovrappongono mai.

Il thread 2 del block 3, con blockDim uguale a 256, si trova in 3 per 256 più 2, cioè 770. Questa è la sua posizione globale nei dati.

Una volta ottenuto i, lo si usa come indice di un array. Ogni thread legge e scrive solo il proprio elemento, senza sovrapposizioni.

Avviando n thread, la formula produce ogni valore da 0 a n meno 1 esattamente una volta. Si ottiene così una copertura uno a uno perfetta dell'array.

Moltiplichi sempre prima di sommare. blockIdx.x * blockDim.x indica l'inizio della porzione, mentre threadIdx.x indica lo spostamento al suo interno.

La stessa idea si estende a 2D e 3D usando i membri .y e .z, ma per gli array piatti è sufficiente la formula con .x. 🚀

Calcoli l'indice globale di un thread.

Ha imparato la formula usata da ogni kernel: blockIdx.x * blockDim.x + threadIdx.x. Essa assegna a ogni thread una posizione univoca nell'array. 🎉