Streaming multiprocesadores, Bloques e Hilos (CUDA)


¿Cuál es la relación entre un núcleo CUDA, un multiprocesador de streaming y el modelo CUDA de bloques e hilos?

¿Qué se asigna a qué y qué se paraleliza y cómo? y lo que es más eficiente, maximizar el número de bloques o el número de hilos?


Mi entendimiento actual es que hay 8 núcleos cuda por multiprocesador. y que cada núcleo cuda será capaz de ejecutar un bloque cuda a la vez. y todos los hilos en ese bloque se ejecutan en ese núcleo en particular.

Es esto correcto?

 50
Author: talonmies, 2010-08-19

3 answers

El diseño del hilo / bloque se describe en detalle en la guía de programación de CUDA. En particular, el capítulo 4 dice:

La arquitectura CUDA está construida alrededor de una matriz escalable de multiprocesadores de streaming multiprocesador (SMs). Cuando un programa CUDA en la CPU del host invoca una cuadrícula del núcleo, los bloques de la cuadrícula se enumeran y se distribuyen a multiprocesadores con capacidad de ejecución disponible. Los subprocesos de un bloque de subprocesos se ejecutan simultáneamente en un multiprocesador, y múltiples bloques de subprocesos pueden ejecutarse simultáneamente en un multiprocesador. A medida que los bloques de subprocesos terminan, se lanzan nuevos bloques en los multiprocesadores desocupados.

Cada SM contiene 8 núcleos CUDA, y en cualquier momento están ejecutando una sola urdimbre de 32 hilos, por lo que se necesitan 4 ciclos de reloj para emitir una sola instrucción para toda la urdimbre. Puedes asumir que los hilos en cualquier warp dado se ejecutan en lock-step, pero para sincronizar entre warps, necesitas usar __syncthreads().

 51
Author: Edric,
Warning: date(): Invalid date.timezone value 'Europe/Kyiv', we selected the timezone 'UTC' for now. in /var/www/agent_stack/data/www/ajaxhispano.com/template/agent.layouts/content.php on line 61
2014-05-08 05:42:41

Para el GTX 970 hay 13 Multiprocesadores de streaming (SM) con 128 Núcleos Cuda cada uno. Los núcleos Cuda también se llaman Procesadores de flujo (SP).

Puede definir cuadrículas que asignan bloques a la GPU.

Puede definir bloques que mapean subprocesos a Procesadores de flujo (los 128 núcleos Cuda por SM).

Una urdimbre siempre está formada por 32 hilos y todos los hilos de una urdimbre se ejecutan simultáneamente.

Para utilizar toda la potencia posible de una GPU se necesitan muchos más hilos por SM que el SM tiene SPs. Para cada Capacidad de cómputo hay un cierto número de subprocesos que pueden residir en un SM a la vez. Todos los bloques que defina se ponen en cola y esperan a que un SM tenga los recursos (número de SPs libre), luego se carga. El SM comienza a cumplir las deformaciones. Dado que una deformación solo tiene 32 Hilos y un SM tiene, por ejemplo, 128 SPs, un SM puede ejecutar 4 deformaciones a la vez. La cosa es que si los hilos hacen acceso a la memoria, el hilo se bloqueará hasta que se satisfaga su solicitud de memoria. En números: Un cálculo aritmético en el SP tiene una latencia de 18-22 ciclos, mientras que un acceso a memoria global sin caché puede tomar hasta 300-400 ciclos. Esto significa que si los hilos de una urdimbre están esperando datos, solo un subconjunto de los 128 SPs funcionaría. Por lo tanto, el programador cambia para ejecutar otro warp si está disponible. Y si este warp bloquea ejecuta el siguiente y así sucesivamente. Este concepto se llama ocultación de latencia. El número de deformaciones y el tamaño del bloque determinan la ocupación (de cuántas deformaciones el SM puede elegir ejecutar). Si la ocupación es alta, es más improbable que no haya trabajo para el SPs.

Su declaración de que cada núcleo cuda ejecutará un bloque a la vez es incorrecta. Si se habla de Streaming Multiprocesadores que pueden ejecutar urdimbres de todos los subprocesos que residen en el SM. Si un bloque tiene un tamaño de 256 subprocesos y su GPU permite 2048 subprocesos a residente por SM, cada SM tendría 8 bloques residentes de los cuales el SM puede elegir deformaciones para ejecutar. Todos los hilos de la las urdimbres ejecutadas se ejecutan en paralelo.

Aquí encontrará los números de las diferentes Capacidades Informáticas y arquitecturas de GPU: https://en.wikipedia.org/wiki/CUDA#Limitations

Puede descargar una hoja de cálculo de ocupación de Nvidia Hoja de Cálculo de ocupación (por Nvidia).

 21
Author: JoeFox,
Warning: date(): Invalid date.timezone value 'Europe/Kyiv', we selected the timezone 'UTC' for now. in /var/www/agent_stack/data/www/ajaxhispano.com/template/agent.layouts/content.php on line 61
2016-06-15 03:56:01

El Distribuidor de Trabajo de Cómputo programará un bloque de subprocesos (CTA) en un SM solo si el SM tiene recursos suficientes para el bloque de subprocesos (memoria compartida, deformaciones, registros, barreras, ...). Se asignan recursos a nivel de bloque de subprocesos como la memoria compartida. La asignación crea suficientes deformaciones para todos los subprocesos en el bloque de subprocesos. El administrador de recursos asigna warps round robin a las sub-particiones SM. Cada subpartición SM contiene un planificador warp, un archivo de registro y unidades de ejecución. Una vez un warp se asigna a una subpartición permanecerá en la subpartición hasta que se complete o sea precedida por un cambio de contexto (arquitectura Pascal). En context switch restore la deformación se restaurará al mismo SM mismo warp-id.

Cuando todos los subprocesos en warp han completado el planificador warp espera a que se completen todas las instrucciones pendientes emitidas por warp y luego el administrador de recursos libera los recursos de nivel warp que incluyen warp-id y el archivo de registro.

Cuando todas las deformaciones en un bloque de subprocesos completado, se liberan recursos a nivel de bloque y el SM notifica al Distribuidor de Trabajo de Cómputo que el bloque se ha completado.

Una vez que se asigna una urdimbre a una subpartición y se asignan todos los recursos, la urdimbre se considera activa, lo que significa que el planificador de urdimbres está rastreando activamente el estado de la urdimbre. En cada ciclo, el planificador de urdimbre determina qué urdimbres activos están estancados y cuáles son elegibles para emitir una instrucción. El planificador de urdimbre elige el más alto warp elegible prioritario y emite 1-2 instrucciones consecutivas de warp. Las reglas para la emisión dual son específicas para cada arquitectura. Si un warp emite una carga de memoria, puede continuar ejecutando instrucciones independientes hasta que alcance una instrucción dependiente. La urdimbre se paralizará hasta que se complete la carga. Lo mismo es cierto para las instrucciones matemáticas dependientes. La arquitectura SM está diseñada para ocultar tanto la latencia de ALU como la de memoria cambiando por ciclo entre deformaciones.

Esto answer no utiliza el término CUDA core ya que introduce un modelo mental incorrecto. Los núcleos CUDA son unidades de ejecución de punto flotante/entero de precisión simple canalizadas. La tasa de emisión y la latencia de dependencia son específicas de cada arquitectura. Cada subpartición SM y SM tiene otras unidades de ejecución, incluidas unidades de carga / almacenamiento, unidades de punto flotante de doble precisión, unidades de punto flotante de media precisión, unidades de rama, etc.

Para maximizar el rendimiento el desarrollador tiene que entender la el comercio fuera de los bloques vs deformación vs registros/hilo.

El término ocupación es la relación de urdimbres activos a urdimbres máximos en un SM. La arquitectura Kepler - Pascal (excepto GP100) tiene 4 programadores warp por SM. El número mínimo de urdimbres por SM debe ser al menos igual al número de programadores de urdimbre. Si la arquitectura tiene una latencia de ejecución dependiente de 6 ciclos (Maxwell y Pascal), entonces necesitaría al menos 6 warps por programador, que es 24 por SM (24 / 64 = 37.5% de ocupación) para cubre la latencia. Si los hilos tienen paralelismo de nivel de instrucción, esto podría reducirse. Casi todos los núcleos emiten instrucciones de latencia variable, como cargas de memoria que pueden tomar entre 80 y 1000 ciclos. Esto requiere urdimbres más activos por planificador de urdimbre para ocultar la latencia. Para cada núcleo hay un punto de compensación entre el número de deformaciones y otros recursos, como la memoria compartida o los registros, por lo que no se recomienda optimizar para una ocupación del 100%, ya que es probable que se haga algún otro sacrificio. El CUDA el generador de perfiles puede ayudar a identificar la tasa de problemas de instrucción, la ocupación y las razones de pérdida para ayudar al desarrollador a determinar ese equilibrio.

El tamaño de un bloque de hilo puede afectar el rendimiento. Si el núcleo tiene bloques grandes y usa barreras de sincronización, entonces los bloqueos de barreras pueden ser razones de bloqueo. Esto se puede aliviar reduciendo las deformaciones por bloque de rosca.

 2
Author: Greg Smith,
Warning: date(): Invalid date.timezone value 'Europe/Kyiv', we selected the timezone 'UTC' for now. in /var/www/agent_stack/data/www/ajaxhispano.com/template/agent.layouts/content.php on line 61
2017-05-26 00:59:15