Streaming multiprocesadores, Bloques y roscas (CUDA)

Question

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

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

---

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

¿Es esto correcto?

Answer 1

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

  

La arquitectura CUDA se construye alrededor de una matriz escalable de multiproceso Streaming Multiprocesadores (SMS). Cuando un programa CUDA en la CPU host activa una red de núcleo, los bloques de la cuadrícula se enumeran y se distribuyen a los multiprocesadores con capacidad de ejecución disponible. Los hilos de un bloque de hilo ejecutan simultáneamente en un multiprocesador, y múltiples bloques de rosca pueden ejecutar simultáneamente en un multiprocesador. Como hilo se bloquea terminan, nuevos bloques se ponen en marcha en los multiprocesadores desocupadas.

Cada SM contiene 8 núcleos CUDA, y en cualquier momento en que sean Ejecución de una única urdimbre de 32 hilos - por lo que toma 4 ciclos de reloj para emitir una sola instrucción para toda la urdimbre. Se puede suponer que los hilos de urdimbre en cualquier dada ejecutar en el bloqueo de paso, pero para sincronizar a través de hilos de urdimbre, es necesario el uso __syncthreads().

Answer 2

Para la GTX 970 hay 13 Transmisión de multiprocesadores (SM) con 128 Cuda Núcleos de cada uno. Cuda Cores también son llamados Stream Processors (SP).

Se puede definir cuadrículas que mapea bloques para la GPU.

Se puede definir bloques que se asignan a las discusiones Procesadores de flujo (los 128 núcleos CUDA por SM).

Una urdimbre está siempre formada por 32 hilos y todos los hilos de un urdimbre se ejecutan simulaneously.

Para utilizar la posible potencia de una GPU que necesitan mucho más hilos por SM que el SM tiene SP. Para cada capacidad Compute hay un cierto número de hilos que pueden residir en un SM a la vez. Todos los bloques que defina se ponen en cola y esperar a que un SM para tener los recursos (número de SP gratis), entonces se carga. El SM comienza a ejecutar Deformaciones. Dado que uno Warp sólo tiene 32 hilos y un SM tiene por ejemplo 128 SPs un SM puede ejecutar 4 Deformaciones en un momento dado. La cosa es que si los hilos hacen acceso a la memoria el hilo se bloqueará hasta que su solicitud de memoria está satisfecho. En números: un cálculo aritmético en el SP tiene una latencia de 18-22 ciclos mientras que un acceso a memoria sin almacenamiento en caché global puede tomar hasta 300-400 ciclos. Esto significa que si los hilos de una urdimbre están esperando datos sólo un subconjunto de las 128 SPs funcionarían. Para ello el programador cambia a ejecutar otro urdimbre si está disponible. Y si esto bloquea urdimbre se ejecuta el siguiente y así sucesivamente. Este concepto se denomina ocultación latencia. El número de urdimbres y el tamaño del bloque determinan la ocupación (de cuántos urdimbres el SM puede optar por ejecutar). Si la ocupación es alta, es más probable que no hay trabajo para los SP.

Su afirmación de que cada núcleo CUDA ejecutará un bloque a la vez que está mal. Si se habla de Transmisión de multiprocesadores pueden ejecutar urdimbres de todo hilo que residen en el SM. Si un bloque tiene un tamaño de 256 hilos y sus allowes GPU 2048 Temas referente a residentes por cada SM SM tendría 8 bloques que residen desde la que el SM puede elegir urdimbres para su ejecución. Todos los hilos de la urdimbre ejecutados se ejecutan en paralelo.

Puedes encontrar números de las diferentes capacidades de calcular y GPU Arquitecturas aquí: https://en.wikipedia.org/wiki/CUDA#Limitations

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

Answer 3

El Distribuidor Compute trabajo programará una secuencia de rosca (CTA) en un SM sólo si el SM tiene recursos suficientes para el paso de rosca (memoria compartida, urdimbres, registros, barreras, ...). recursos a nivel de bloque hilo tan memoria compartida se asignan. El asignar crea urdimbres suficientes para todos los hilos en el bloque de hilo. Los Asigna gestor de recursos urdimbres utilizando ronda robin a los SM sub-particiones. Cada subpartición SM contiene un planificador de la urdimbre, banco de registros, y las unidades de ejecución. Una vez que una urdimbre se asigna a un subpartición que permanecerá en la subpartición hasta que se complete o se ha vaciado anteriormente por un cambio de contexto (arquitectura Pascal). En cambio de contexto restaurar la urdimbre será restaurado a la misma SM misma urdimbre-id.

Cuando todos los hilos en la urdimbre han completado las esperas urdimbre planificador para todas las instrucciones emitidas por la urdimbre para completar y luego el director de recursos libera los recursos a nivel de la urdimbre que incluyen urdimbre Identificación y registro de archivo.

Cuando todos los hilos de urdimbre en una secuencia de rosca se liberan recursos completos a nivel de bloque y luego los notifica SM el Distribuidor Calcular el trabajo que el bloque ha finalizado.

Una vez que una urdimbre se asigna a un subpartición y todos los recursos se asignan de la urdimbre se considera activa, el sentido de que el programador de urdimbre está siguiendo activamente el estado de la urdimbre. En cada ciclo de la urdimbre planificador determinar qué activos urdimbres están estancadas y los que son elegibles para emitir una instrucción. El planificador de urdimbre recoge la más alta prioridad elegibles urdimbre y emite instrucciones 1-2 consecutivos a partir de la urdimbre. Las reglas para la doble cuestión son específicos de cada arquitectura. Si una urdimbre emite una carga de memoria que puede seguir las instrucciones independientes ejecutados hasta que llega a una instrucción dependiente. La urdimbre a su vez informará paralizado hasta que se complete la carga. Lo mismo es cierto para las instrucciones matemáticas dependientes. La arquitectura SM está diseñado para ocultar tanto ALU y latencia de la memoria al cambiar por ciclo entre urdimbres.

Esta respuesta no usa el término núcleo CUDA como Esto introduce un modelo mental incorrecta. núcleos CUDA son flotantes unidades de ejecución punto / enteros de precisión simple pipeline. El tipo de emisión de latencia y la dependencia es específico para cada arquitectura. Cada subpartición SM y SM tiene otras unidades de ejecución incluidas las unidades de carga / almacenamiento, unidades de punto flotante de doble precisión, unidades de punto flotante de precisión media, conjuntos de derivación, etc.

Con el fin de maximizar el rendimiento del desarrollador tiene que entender el comercio fuera de los bloques frente a frente urdimbres registros / hilo.

El término ocupación es la relación de urdimbres activas para urdimbres máximo sobre un SM. Kepler - arquitectura Pascal (excepto GP100) tienen 4 programadores de urdimbre por SM. El número mínimo de urdimbres por SM debería al menos ser igual al número de programadores de urdimbre. Si la arquitectura tiene una latencia ejecución dependiente de 6 ciclos (Maxwell y Pascal) a continuación, que se necesita al menos 6 urdimbres por planificador que es 24 por SM (24/64 = 37,5% de ocupación) para cubrir la latencia. Si los hilos tienen paralelismo a nivel entonces esto podría ser reducida. Casi todos los granos dicten las instrucciones de latencia variables tales como cargas de memoria que pueden tomar 80-1000 ciclos. Esto requiere urdimbres más activas por urdimbre planificador para ocultar la latencia. Para cada núcleo no es una solución de compromiso entre el número de punto de urdimbres y otros recursos como la memoria o registros optimizando así el 100% de ocupación compartida no se aconseja que es probable que se hizo algún otro sacrificio. El generador de perfiles CUDA puede ayudar a identificar la tasa de instrucciones tema, la ocupación, y las razones de puesto con el fin de ayudar a los desarrolladores a determinar ese equilibrio.

El tamaño de un bloque de hilo puede afectar al rendimiento. Si el núcleo tiene grandes bloques y utiliza barreras de sincronización entonces puestos de barrera pueden ser una vienen razones de pérdida. Esto puede ser aliviado mediante la reducción delurdimbres por bloque hilo.

Answer 4

Hay múltiples streaming de varios procesadores en un solo dispositivo.
A SM puede contener múltiples bloques. Cada bloque puede contener varios hilos.
Un SM tiene múltiples núcleos CUDA (como desarrollador, usted no debe preocuparse por esto, ya que es abstraído por urdimbre), que trabajarán en hilo. SM siempre trabajando en urdimbre de hilos (siempre 32). Una urdimbre sólo se trabaja en la rosca del mismo bloque.
SM y bloquear ambos tienen límites a la cantidad de hilo, número de registro y la memoria compartida.