Letture coalescenti contro strided

La coalescenza dipende dall'indirizzo toccato da ciascun thread. La mappatura da thread a indirizzo determina se una sola transazione serve l'intero warp.

Quando i thread adiacenti leggono indirizzi adiacenti, il warp copre un'unica sequenza contigua. Questa disposizione ordinata è un accesso coalescente.

Il thread 0 accede all'indice 0, il thread 1 all'indice 1 e così via. Tutti i 32 indirizzi rientrano in un'unica linea allineata, quindi alla GPU serve una sola transazione.

Uno stride è la distanza fissa tra gli indirizzi toccati da thread successivi. Non appena questa distanza supera un elemento, la coalescenza inizia a deteriorarsi.

Con uno stride pari a 2, il thread 0 legge 0, il thread 1 legge 2 e il thread 2 legge 4. Ora il warp copre il doppio della memoria, quindi richiede più transazioni.

Raddoppiando lo stride, si raddoppiano all'incirca le linee interessate. Stride elevati possono portare un warp a effettuare molte transazioni usando solo una piccola parte di ogni linea.

Assegnare a ogni thread un'intera colonna di una matrice memorizzata per righe crea uno stride enorme. Nel codice sembra ordinato, ma disperde silenziosamente ogni warp.

Spesso basta scambiare l'indice che varia più velocemente con il thread. Se i thread consecutivi percorrono posizioni consecutive della memoria, le letture tornano a essere coalescenti.

Anche un pattern coalescente ne risente se inizia a metà di una linea. Un offset non allineato sposta il warp oltre un confine, dividendo una lettura in due.

Per valutare un pattern, non immagini un solo thread. Immagini tutte e 32 le lane contemporaneamente e si chieda quante linee coprono complessivamente i loro indirizzi. 🔍

Ha visto che le letture coalescenti mantengono vicini gli elementi adiacenti, mentre uno stride li disperde tra più linee. Lasci che threadIdx.x controlli l'indice più veloce. 🎉