cuda - arquitectura de un gpu
Streaming de multiprocesadores, bloques e hilos(CUDA) (3)
¿Cuál es la relación entre un núcleo CUDA, un multiprocesador de transmisión y el modelo CUDA de bloques e hilos?
¿Qué se asigna a qué y qué se paraleliza y cómo? y, ¿qué es más eficiente, maximizar la cantidad de bloques o la cantidad de hilos?
Mi comprensión 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 serie en ese núcleo particular.
¿Es esto correcto?
El Distribuidor de Compute Work programará un bloque de hilos (CTA) en un SM solo si el SM tiene suficientes recursos para el bloque de hilos (memoria compartida, urdimbres, registros, barreras, ...). Los recursos de nivel de bloque de subprocesos se asignan. La asignación crea deformaciones suficientes para todos los hilos en el bloque de hilos. El administrador de recursos asigna warps round robin a las subparticiones SM. Cada subpartición SM contiene un planificador warp, un archivo de registro y unidades de ejecución. Una vez que un warp se asigna a una subpartición, permanecerá en la subpartición hasta que se complete o sea reemplazado por un cambio de contexto (arquitectura Pascal). Con el cambio de contexto, la urdimbre se restaurará en la misma SM id warp.
Cuando todos los hilos en warp se hayan completado, el planificador de warp espera a que se completen todas las instrucciones pendientes emitidas por el warp y luego el administrador de recursos libera los recursos del warp-level que incluyen warp-id y el archivo de registro.
Cuando se completan todas las distorsiones en un bloque de subprocesos, se liberan recursos de 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 un warp a una subpartición y se asignan todos los recursos, el warp se considera activo, lo que significa que el planificador de warp está rastreando activamente el estado del warp. En cada ciclo, el planificador de warp determina qué warps activos están estancados y cuáles son elegibles para emitir una instrucción. El planificador de warp selecciona la urdimbre elegible de mayor prioridad y emite 1-2 instrucciones consecutivas de warp. Las reglas para la doble emisión son específicas de cada arquitectura. Si un warp emite una carga de memoria, puede continuar ejecutando instrucciones independientes hasta que llegue a una instrucción dependiente. La urdimbre se informará detenida hasta que la carga se complete. Lo mismo es cierto para las instrucciones de matemáticas dependientes. La arquitectura SM está diseñada para ocultar la ALU y la latencia de memoria cambiando cada ciclo entre warps.
Esta respuesta no usa el término núcleo CUDA ya que esto introduce un modelo mental incorrecto. Los núcleos CUDA son unidades de ejecución de coma flotante / entero de precisión única segmentadas. La tasa de problemas y la latencia de dependencia es específica de cada arquitectura. Cada subpartición SM y SM tiene otras unidades de ejecución que incluyen unidades de carga / almacenamiento, unidades de punto flotante de precisión doble, unidades de coma flotante de media precisión, unidades de bifurcación, etc.
Para maximizar el rendimiento, el desarrollador debe entender el intercambio de bloques vs. warps vs. registros / hilo.
El término ocupación es la relación entre deformaciones activas y deformaciones máximas en un SM. Kepler: la arquitectura de Pascal (excepto GP100) tiene 4 planificadores de deformación por SM. El número mínimo de urdimbres por SM debe ser al menos igual al número de planificadores de warp. 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 cubrir 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 tales como cargas de memoria que pueden tomar 80-1000 ciclos. Esto requiere más warps activos por programador de warp para ocultar la latencia. Para cada núcleo hay un punto de intercambio entre el número de warps y otros recursos, como la memoria compartida o los registros, por lo que no se recomienda optimizar el 100% de ocupación, ya que probablemente se haga otro sacrificio. El generador de perfiles de CUDA 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 hilos puede afectar el rendimiento. Si el kernel tiene bloques grandes y usa barreras de sincronización, los puestos de barrera pueden ser un motivo de parada. Esto se puede aliviar reduciendo las deformaciones por bloque de hilo.
El diseño del hilo / bloque se describe en detalle en la guía de programación CUDA . En particular, el capítulo 4 dice:
La arquitectura CUDA se basa en una matriz escalable de Multiprocesadores de transmisión (SM) multiproceso. Cuando un programa CUDA en la CPU host invoca una grilla de kernel, los bloques de la grilla son enumerados y distribuidos a multiprocesadores con capacidad de ejecución disponible. Los hilos de un bloque de hilos se ejecutan simultáneamente en un multiprocesador, y múltiples bloques de hilos pueden ejecutarse simultáneamente en un multiprocesador. A medida que finalizan los bloques de subprocesos, 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 subprocesos, por lo que se requieren 4 ciclos de reloj para emitir una instrucción única para toda la disformidad. Puede suponer que los hilos en un warp determinado se ejecutan en el paso de bloqueo, pero para sincronizarlos a lo largo de los warps, debe usar __syncthreads() .
Para la GTX 970 hay 13 Multiprocesadores Streaming (SM) con 128 Cuda Cores cada uno. Los núcleos Cuda también se llaman Stream Processors (SP).
Puede definir las cuadrículas que asignan bloques a la GPU.
Puede definir bloques que mapeen hilos a Procesadores de flujo (los 128 Cuda Cores por SM).
Una urdimbre siempre está formada por 32 hilos y todos los hilos de una urdimbre se ejecutan de forma simultánea.
Para usar toda la potencia posible de una GPU necesita mucho más hilos por SM que SM tiene SP. Para cada capacidad de cómputo existe una cierta cantidad de subprocesos que pueden residir en un SM a la vez. Todos los bloques que defina están en cola y esperan que un SM tenga los recursos (número de SP libres), luego se carga. El SM comienza a ejecutar Warps. Como un Warp solo tiene 32 Threads y un SM tiene, por ejemplo, 128 SP, un SM puede ejecutar 4 Warps en un momento determinado. El problema es que si los subprocesos tienen 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 de memoria global no almacenado en caché puede tomar hasta 300-400 ciclos. Esto significa que si los hilos de un warp están esperando datos, solo un subconjunto de los 128 SP funcionaría. Por lo tanto, el planificador 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 urdimbres y el tamaño del bloque determinan la ocupación (de cuántas deformaciones puede elegir ejecutar el SM). Si la ocupación es alta, es más improbable que no haya trabajo para los SP.
Su afirmación de que cada núcleo cuda ejecutará un bloque a la vez es incorrecta. Si habla de multiprocesadores de transmisión, puede ejecutar warps desde todos los hilos que residen en el SM. Si un bloque tiene un tamaño de 256 subprocesos y su GPU permite que 2048 subprocesos sean residentes por SM, cada SM tendría 8 bloques residentes desde los cuales el SM puede elegir distorsiones para ejecutar. Todos los hilos de los warps ejecutados se ejecutan en paralelo.
Aquí encontrará números para las diferentes Capacidades de cálculo y Arquitecturas de GPU: https://en.wikipedia.org/wiki/CUDA#Limitations
Puede descargar una hoja de cálculo de ocupación de la hoja de Cálculo de ocupación de Nvidia (por Nvidia) .