memory - ¿Cuál es la diferencia entre CUDA compartido y la memoria global?

shared-memory

Question

Me estoy confundiendo acerca de cómo usar la memoria compartida y global en CUDA, especialmente con respecto a lo siguiente:

• Cuando usamos cudaMalloc() , ¿obtenemos un puntero a la memoria compartida o global?
• ¿Reside la memoria global en el host o dispositivo?
• ¿Hay un límite de tamaño para cualquiera de los dos?
• ¿Cuál es más rápido de acceder?

• ¿Es lo mismo almacenar una variable en la memoria compartida que pasar su dirección a través del kernel? Es decir, en lugar de tener

  __global__ void kernel() {
     __shared__ int i;
     foo(i);
  }

  por qué no hacer equivalentemente

  __global__ void kernel(int *i_ptr) {
     foo(*i_ptr);
  }
  
  int main() {
     int *i_ptr;
     cudaMalloc(&i_ptr, sizeof(int));
     kernel<<<blocks,threads>>>(i_ptr);
  }

Ha habido muchas preguntas sobre problemas específicos de velocidad en la memoria global frente a la memoria compartida, pero ninguna abarca una descripción general de cuándo usar una en la práctica.

Muchas gracias

Answer 1

La memoria compartida CUDA es memoria compartida entre los hilos dentro de un bloque, es decir, entre bloques en una cuadrícula, los contenidos de la memoria compartida no están definidos. Se puede pensar como un caché L2 administrado manualmente.

Por lo general, la memoria global reside en el dispositivo, pero las versiones recientes de CUDA (si el dispositivo lo admite) pueden mapear la memoria del host en el espacio de direcciones del dispositivo, desencadenando una transferencia de DMA in situ desde el host a la memoria del dispositivo en tales ocasiones.

Hay un límite de tamaño en la memoria compartida, según el dispositivo. Se informa en las capacidades del dispositivo, recuperadas al enumerar los dispositivos CUDA. La memoria global está limitada por la memoria total disponible para la GPU. Por ejemplo, una GTX680 ofrece 48 kB de memoria compartida y memoria de dispositivo 2 G iB.

La memoria compartida es más rápida de acceder que la memoria global, pero los patrones de acceso deben alinearse cuidadosamente (tanto para memoria compartida como global) para ser eficiente. Si no puede hacer que sus patrones de acceso estén alineados correctamente, use texturas (también memoria global, pero a la que se accede a través de un caché y un entorno diferente, que puede manejar mejor el acceso desalineado).

¿Es lo mismo almacenar una variable en la memoria compartida que pasar su dirección a través del kernel?

No, definitivamente no. El código que ha propuesto sería un caso en el que utilizaría la memoria global transferida in situ. La memoria compartida no se puede pasar entre los núcleos, ya que los contenidos de un bloque compartido se definen dentro de un bloque de ejecución de hilos solamente.

Answer 2

Los contenidos de la memoria global son visibles para todos los hilos de la grilla. Cualquier hilo puede leer y escribir en cualquier ubicación de la memoria global.

La memoria compartida está separada para cada bloque de la grilla. Cualquier hilo de un bloque puede leer y escribir en la memoria compartida de ese bloque. Un hilo en un bloque no puede acceder a la memoria compartida de otro bloque.

1. cudaMalloc siempre asigna memoria global.
2. La memoria global reside en el dispositivo.
3. Obviamente, cada memoria tiene un límite de tamaño. La memoria global es la cantidad total de DRAM de la GPU que está utilizando. Por ejemplo, uso GTX460M que tiene 1536 MB DRAM, por lo tanto, 1536 MB de memoria global. La memoria compartida es especificada por la arquitectura del dispositivo y se mide en base a cada bloque. Los dispositivos de la capacidad de cómputo 1.0 a 1.3 tienen 16 KB/Block , computan 2.0 a 7.0 tienen 48 KB/Block memoria compartida.
4. La memoria compartida tiene magnitudes más rápidas de acceder que la memoria global. Es como un caché local compartido entre los hilos de un bloque.
5. No. Solo las direcciones de memoria global se pueden pasar a un kernel iniciado desde el host. En su primer ejemplo, la variable se lee desde la memoria compartida, mientras que en la segunda se lee desde la memoria global.

Actualizar:

Los dispositivos de Compute Capability 7.0 (Volta Architecture) permiten asignar memoria compartida de hasta 96 KB por bloque, siempre que se cumplan las siguientes condiciones.

• La memoria compartida se asigna dinámicamente
• Antes de iniciar el núcleo, el tamaño máximo de la memoria compartida dinámica se especifica utilizando la función cudaFuncSetAttribute siguiente manera.

__global__ void MyKernel(...)
{
    extern __shared__ float shMem[];
}

int bytes = 98304; //96 KB
cudaFuncSetAttribute(MyKernel, cudaFuncAttributeMaxDynamicSharedMemorySize, bytes);

MyKernel<<<gridSize, blockSize, bytes>>>(...);

Answer 3

• Cuando usamos cudaMalloc ()

  Para almacenar datos en la GPU que pueden ser comunicados nuevamente al host, necesitamos tener una memoria alojada que viva hasta que se libere, ver la memoria global como el espacio de montón con vida hasta que la aplicación se cierre o se libere, es visible a cualquier hilo y bloque que tenga un puntero a esa región de memoria. La memoria compartida se puede considerar como espacio de pila con vida hasta que termina un bloque de un kernel, la visibilidad está limitada a solo hilos dentro del mismo bloque. Entonces, cudaMalloc se usa para asignar espacio en la memoria global.

• ¿Obtenemos un puntero a la memoria compartida o global?

  Obtendrá un puntero a una dirección de memoria que resida en la memoria global.

• ¿Reside la memoria global en el host o dispositivo?

  La memoria global reside en el dispositivo. Sin embargo, hay formas de utilizar la memoria del host como memoria "global" utilizando la memoria asignada; consulte: CUDA Zero Copy Consideraciones sobre la memoria , sin embargo, pueden ser velocidades lentas debido a las limitaciones de velocidad de transferencia del bus.

• ¿Hay un límite de tamaño para cualquiera de los dos?

  El tamaño de la memoria global depende de una tarjeta a otra, desde ninguna hasta 8GB. Mientras que la memoria compartida depende de la capacidad de cálculo. Cualquier cosa por debajo de la capacidad de cálculo 2.x tiene un máximo de 16 KB de memoria compartida por multiprocesador (donde la cantidad de multiprocesadores varía de una tarjeta a otra). Y las tarjetas con una capacidad de cálculo de 2.xy mayor tienen un máximo de 48 KB de memoria compartida por multiprocesador.

  Si está utilizando la memoria mapeada, la única limitación es cuánto tiene la máquina host en la memoria.

• ¿Cuál es más rápido de acceder?

  En términos de números brutos, la memoria compartida es mucho más rápida (memoria compartida ~ 1.7TB / s, mientras que la memoria global ~ 150GB / s). Sin embargo, para hacer cualquier cosa que necesites para llenar la memoria compartida con algo, usualmente sacas de la memoria global. Si el acceso a la memoria de la memoria global se combina (no aleatorio), obtendrá velocidades de hasta 150-200 GB / s dependiendo de la tarjeta y su interfaz de memoria.

  El uso de la memoria compartida es cuando necesita dentro de un bloque de hilos, reutilizar datos ya extraídos o evaluados desde la memoria global. Por lo tanto, en lugar de volver a extraer de la memoria global, la coloca en la memoria compartida para que otros hilos dentro del mismo bloque puedan verse y reutilizarse.

• ¿Es lo mismo almacenar una variable en la memoria compartida que pasar su dirección a través del kernel?

  No, si pasas una dirección de algo, siempre es una dirección para la memoria global. Desde el host no puede establecer la memoria compartida, a menos que la pase como una constante donde el núcleo establece la memoria compartida a esa constante, o le pasa una dirección a la memoria global donde es extraída por el kernel cuando sea necesario.