2012-02-08 22 views
7

No encuentro una mejora en la velocidad con la memoria compartida en un NVIDIA Tesla M2050 con aproximadamente 49K de memoria compartida por bloque. De hecho, si asigno una gran matriz de caracteres en la memoria compartida, ralentiza mi programa. Por ejemploCUDA: ¿El acceso global a la memoria global es más rápido que la memoria compartida? Además, ¿la asignación de una gran matriz de memoria compartida ralentiza el programa?

__shared__ char database[49000]; 

me da tiempos de funcionamiento más lento que

__shared__ char database[4900]; 

El programa accede sólo los primeros 100 caracteres de la base de datos por lo que el espacio extra es innecesaria. No puedo entender por qué está pasando esto. Cualquier ayuda sería apreciada. Gracias.

Respuesta

28

La razón del rendimiento relativamente pobre de la memoria compartida CUDA cuando se utilizan matrices más grandes puede tener que ver con el hecho de que cada multiprocesador tiene una cantidad limitada de memoria compartida disponible.

Cada multiprocesador aloja varios procesadores; para dispositivos modernos, generalmente 32, la cantidad de hilos en una urdimbre. Esto significa que, en ausencia de divergencia o paradas de memoria, la tasa promedio de procesamiento es de 32 instrucciones por ciclo (la latencia es alta debido a la canalización).

CUDA planifica varios bloques para un multiprocesador. Cada bloque consta de varias urdimbres. Cuando una urdimbre se detiene en un acceso de memoria global (incluso los accesos fusionados tienen una alta latencia), se procesan otras distorsiones. Esto efectivamente oculta la latencia, por lo que la memoria global de alta latencia es aceptable en las GPU. Para ocultar eficazmente la latencia, necesitas suficientes warps adicionales para ejecutar hasta que el warp estancado pueda continuar. Si todos los warps se detienen en los accesos a la memoria, ya no se puede ocultar la latencia.

La memoria compartida se asigna a bloques en CUDA y se almacena en un solo multiprocesador en el dispositivo GPU. Cada multiprocesador tiene una cantidad relativamente pequeña y fija de espacio de memoria compartida. CUDA no puede programar más bloques para multiprocesadores de los que los multiprocesadores pueden admitir en términos de memoria compartida y uso de registros. En otras palabras, si la cantidad de memoria compartida en un multiprocesador es X y cada bloque requiere Y de memoria compartida, CUDA no programará más bloques de piso (X/Y) a la vez para cada multiprocesador (podría ser menor ya que hay otras restricciones, como el uso del registro).

Ergo, al aumentar el uso de la memoria compartida de un bloque, es posible que esté reduciendo el número de warps activos (la ocupación) de su núcleo, lo que perjudica el rendimiento. Debería examinar su código de kernel compilando con el indicador -Xptxas = "- v"; esto debería darle registro y compartir el uso de la memoria constante & para cada kernel. Utilice esta información y los parámetros de inicio del kernel, así como otra información requerida, en la versión más reciente de la Calculadora de ocupación CUDA para determinar si puede verse afectado por la ocupación.

EDIT:

Para hacer frente a la otra parte de su pregunta, suponiendo que no haya conflictos de bancos de memoria compartida y perfecta coalescencia de memoria global accesos ... hay dos dimensiones a esta respuesta: la latencia y ancho de banda. La latencia de la memoria compartida será menor que la de la memoria global, ya que la memoria compartida está en chip. El ancho de banda será muy similar. Ergo, si puede ocultar la latencia de acceso a la memoria global mediante coalescencia, no hay penalización (nota: el patrón de acceso es importante aquí, en esa memoria compartida permite patrones de acceso potencialmente más diversos con poca o ninguna pérdida de rendimiento, por lo que puede ser beneficioso al uso de la memoria compartida, incluso si puede ocultar toda la latencia de la memoria global).

+0

Excelente descripción. – harrism

+0

Gracias. Esto es muy útil. – Ross

+0

¿Por qué el ancho de banda sería el mismo? ¿O lo dijiste solo por el M2050? – einpoklum

2

Además, si aumenta la memoria compartida por bloque, CUDA programará las cuadrículas con menos bloques simultáneos, por lo que todos tienen suficiente memoria compartida, por lo que reduce el paralelismo y aumenta el tiempo de ejecución.

+2

La ocupación no es lo mismo que la eficiencia. Estoy evitando la tentación de rechazar esto, ya que solo es incompleto, no necesariamente incorrecto ... pero podría ser engañoso. Puede obtener el 100% de eficiencia ejecutando tantas deformaciones completas como multiprocesadores. – Patrick87

1

La cantidad de recursos disponibles en el gpu es limitada. El número de bloques que se ejecutan simultáneamente es inversamente proporcional al tamaño de la memoria compartida por bloque.

Esto explica por qué el tiempo de ejecución es más lento cuando se inicia el núcleo que utiliza una gran cantidad de memoria compartida.

+3

Ocupación no es lo mismo que eficiencia. Estoy evitando la tentación de rechazar esto, ya que solo es incompleto, no necesariamente incorrecto ... pero podría ser engañoso. Puede obtener el 100% de eficiencia ejecutando tantas deformaciones completas como multiprocesadores. – Patrick87

Cuestiones relacionadas