CUDA: ¿Cuándo usar la memoria compartida y cuándo confiar en el almacenamiento en caché L1?

Question

Después de liberar Compute Capability 2.0 (Fermi), me he preguntado si queda algún caso de uso para la memoria compartida. Es decir, ¿cuándo es mejor usar memoria compartida que simplemente dejar que L1 realice su magia en el fondo?

¿La memoria compartida está ahí para permitir que los algoritmos diseñados para CC < 2.0 se ejecuten eficientemente sin modificaciones?

Para colaborar a través de la memoria compartida, los hilos en un bloque escriben en la memoria compartida y sincronizan con __syncthreads(). ¿Por qué no simplemente escribir en la memoria global (a través de L1) y sincronizar con __threadfence_block()? La última opción debería ser más fácil de implementar, ya que no tiene que relacionarse con dos ubicaciones diferentes de valores, y debería ser más rápida porque no hay una copia explícita de la memoria global a compartida. Dado que los datos se almacenan en caché en L1, los subprocesos no tienen que esperar que los datos lleguen realmente a la memoria global.

Con memoria compartida, se garantiza que un valor que se puso allí permanece allí durante toda la duración del bloque. Esto es opuesto a los valores en L1, que son desalojados si no se usan con la suficiente frecuencia. ¿Hay algún caso en el que también es mejor almacenar en caché esos datos que se usan con poca frecuencia en la memoria compartida que dejar que el L1 los administre en función del patrón de uso que tiene el algoritmo en realidad?

Answer 1

Por lo que sé, la memoria caché L1 en una GPU se comporta de forma muy similar a la memoria caché en una CPU. Entonces su comentario de que "Esto es opuesto a los valores en L1, que son desalojados si no se usan con suficiente frecuencia" no tiene mucho sentido para mí

Los datos en la memoria caché L1 no se desalojan cuando no está usado con suficiente frecuencia Por lo general, se desaloja cuando se realiza una solicitud para una región de memoria que no estaba previamente en la memoria caché y cuya dirección se resuelve a una que ya está en uso. No sé el algoritmo de almacenamiento en caché exacta empleada por Nvidia, pero asumiendo una asociativo regular de n vías, a continuación, cada entrada de la memoria sólo puede ser almacenado en caché en un pequeño subconjunto de toda la memoria caché, basado en su dirección

supongo que esto también puede responder tu pregunta. Con la memoria compartida, obtienes el control total de lo que se almacena en el lugar, mientras que con la caché, todo se hace automáticamente. Aunque el compilador y la GPU aún pueden ser muy hábiles para optimizar los accesos a la memoria, a veces puede encontrar una mejor manera, ya que usted es quien sabe qué entrada se le dará, y qué hilos harán qué (para un cierto extensión por supuesto)

Answer 2

Las cargas/tiendas de memoria global están sujetas a reglas de fusión, incluso si los datos están en caché, pero la memoria compartida es mucho más flexible en términos de acceso aleatorio. Intenté usar el almacenamiento en caché L1 para almacenar/calcular un histograma y terminó siendo mucho, mucho más lento que el uso de memoria compartida debido al patrón de acceso semialeatorio.

Además, de acuerdo con NVIDIA employee, las memorias caché L1 actuales son de escritura directa (escribe de inmediato en la memoria caché L2), lo que ralentizará su programa.

Así que, básicamente, las cachés se interponen en el camino si realmente desea rendimiento.