Evitare i conflitti tra banchi

La shared memory è suddivisa in 32 bank, una per ogni thread del warp. Distribuire gli accessi tra queste bank permette a tutti i 32 thread di leggere alla massima velocità.

Le word consecutive da 4 byte finiscono in bank consecutive, ripartendo da capo dopo 32. La word 0 si trova quindi nella bank 0, così come la word 32.

Si verifica un conflitto tra bank quando due thread di un warp accedono alla stessa bank ma a word diverse. L'hardware deve serializzare questi accessi.

Un conflitto N-way trasforma un accesso veloce in N accessi lenti. Un conflitto 32-way può rendere la shared memory lenta come se gli accessi fossero completamente seriali. 🐢

Quando ogni thread legge tile[threadIdx.x], ogni thread accede a una bank distinta. Non ci sono conflitti e si raggiunge la larghezza di banda massima.

Leggere con uno stride pari a 32 manda ogni thread nella stessa bank. È il caso peggiore: un conflitto completo 32-way.

Qualsiasi stride pari che abbia un fattore in comune con 32 causa conflitti parziali. Uno stride di 2, per esempio, genera un conflitto a 2 vie all'interno del warp.

Buone notizie: se tutti i thread leggono lo stesso indirizzo, l'hardware lo trasmette in un'unica operazione. Leggere la stessa parola va bene; il conflitto si verifica solo quando si accede alla stessa bank per leggere parole diverse.

Per i tile 2D, aggiunga una colonna extra con un padding di +1. In questo modo ogni riga viene traslata e gli accessi alle colonne non finiscono più tutti nella stessa bank.

La colonna extra rende la lunghezza della riga coprima con 32. Ora, scendendo lungo una colonna, si visita ogni volta una bank diversa, eliminando il conflitto.

I conflitti tra bank non sono visibili nel codice sorgente. Utilizzi Nsight Compute per misurare i conflitti della shared memory prima di impegnarsi a risolverli.

Verifichiamo che cosa rende veloce l'accesso condiviso.

Ha imparato che la shared memory ha 32 bank, che gli accessi a parole diverse nella stessa bank vengono serializzati e che il padding di +1 risolve i conflitti sulle colonne. Prossimo argomento: la shared memory dinamica. 🎯