CUDA: Quando utilizzare la memoria condivisa e quando fare affidamento sulla memorizzazione nella cache L1?

Question

Dopo che è stata rilasciata la funzione di calcolo 2.0 (Fermi), mi sono chiesto se per la memoria condivisa sono rimasti dei casi d'uso. Cioè, quando è meglio usare la memoria condivisa che lasciare che L1 esegua la sua magia sullo sfondo?

La memoria condivisa è semplicemente lì per consentire agli algoritmi progettati per CC < 2.0 di funzionare in modo efficiente senza modifiche?

Per collaborare tramite memoria condivisa, i thread in un blocco scrivono nella memoria condivisa e si sincronizzano con __syncthreads(). Perché non scrivere semplicemente nella memoria globale (tramite L1) e sincronizzarsi con __threadfence_block()? Quest'ultima opzione dovrebbe essere più semplice da implementare in quanto non deve riguardare due diverse posizioni di valori, e dovrebbe essere più veloce perché non esiste una copia esplicita dalla memoria globale a quella condivisa. Poiché i dati vengono memorizzati nella cache in L1, i thread non devono attendere che i dati raggiungano effettivamente la memoria globale.

Con la memoria condivisa, è garantito che un valore che è stato messo lì rimane lì per tutta la durata del blocco. Questo è in contrasto con i valori in L1, che vengono sfrattati se non vengono utilizzati abbastanza spesso. Ci sono casi in cui è meglio memorizzare nella cache i dati usati raramente nella memoria condivisa piuttosto che lasciare che la L1 li gestisca in base al modello di utilizzo effettivamente utilizzato dall'algoritmo?

Answer 1

Per quanto ne so, la cache L1 in una GPU si comporta come la cache di una CPU. Quindi il tuo commento che "Questo è in contrasto con i valori in L1, che vengono sfrattati se non vengono usati abbastanza spesso" non ha molto senso per me

I dati sulla cache L1 non sono sfrattati quando non è usato abbastanza spesso. Di solito è sfrattato quando viene fatta una richiesta per un'area di memoria che non era precedentemente nella cache e il cui indirizzo si risolve in uno già in uso. Non conosco l'algoritmo di caching esatto impiegato da NVidia, ma assumendo una normale associazione n-way, quindi ogni voce di memoria può essere memorizzata solo in un piccolo sottoinsieme dell'intera cache, basata sul suo indirizzo

Suppongo che questo può anche rispondere alla tua domanda. Con la memoria condivisa, hai il pieno controllo su ciò che viene memorizzato dove, mentre con la cache, tutto viene fatto automaticamente. Anche se il compilatore e la GPU possono ancora essere molto intelligenti nell'ottimizzare gli accessi alla memoria, a volte puoi ancora trovare un modo migliore, dato che sei quello che sa quale input sarà dato e quali thread faranno cosa (ad un certo portata ovviamente)

Answer 2

I carichi/negozi di memoria globale sono soggetti alle regole di coalescenza, anche se i dati sono in cache, ma la memoria condivisa è molto più flessibile in termini di accesso casuale. Ho provato ad usare il caching L1 per memorizzare/calcolare un istogramma e finisce per essere molto, molto più lento di usare la memoria condivisa a causa del pattern di accesso semi-casuale.

Inoltre, in base a uno NVIDIA employee, le cache L1 correnti sono write-through (scrive immediatamente nella cache L2), il che rallenterà il programma.

Quindi, in pratica, le cache si intromettono se si desidera davvero prestazioni.