CUDA: ¿En qué espacio de memoria se almacena una matriz de tamaño fijo?

Question

Al configurar una matriz de tamaño fijo en un núcleo, tales como:

int my_array[100]; 

en que el espacio de memoria termina la matriz de arriba?

En particular, me gustaría saber si una matriz de este tipo se puede almacenar en el archivo de registro o en la memoria compartida en dispositivos> = 2.0 y, de ser así, cuáles son los requisitos.

Answer 1

Para Fermi (y arquitecturas probablemente anteriores), para una una matriz que se almacena en el archivo de registro, el siguiente condiciones deben cumplirse:

1. el array sólo está indexado con constantes
2. Hay registros disponibles
3. de esperar, el compilador también hace algunos análisis para determinar el impacto en el rendimiento general

El motivo de (1) es que los índices de registro están codificados directamente dentro de las instrucciones de SASS. No hay forma de abordar los registros de forma indirecta.

Los principales factores que limita el número de registros encontrados para (2) son: instrucciones

• El SASS contienen sólo 6 bits para indexación registro, lo que limita el número de registros que se puede utilizar en un núcleo para 64. El número real es 63 entonces uno está reservado para algo.
• Un SM tiene un bloque de registros que comparten todos los hilos que se encuentran simultáneamente en vuelo.
• Los registros también son necesarios para contener variables, por lo que el compilador debe equilibrar el uso de registros para obtener el mejor rendimiento general.

Una posible solución alternativa para (1) es el despliegue del bucle. Si un bucle utiliza un contador de bucle como un índice en una matriz, desenrollar el bucle (con #pragma unroll o manualmente) hace que los índices de la matriz se vuelvan constantes ya que ahora hay una instrucción SASS separada para cada acceso de la matriz.

Basado en parte en esta presentación de NVIDIA: Local Memory and Register Spilling. El documento también detalla cómo la ubicación de las variables y las matrices afectan el rendimiento.

Answer 2

Arrays locales dentro de un núcleo, ya que el que ha definido se asigna en los registros y en la memoria local cuando no hay suficiente registro.

Si desea asignar la matriz en la memoria compartida se debe especificar como sigue:

__shared__ int my_array[100];