Un bloque de subprocesos es una abstracción de programación que representa un grupo de subprocesos que se pueden ejecutar en serie o en paralelo. Para un mejor mapeo de procesos y datos, los subprocesos se agrupan en bloques de subprocesos. Anteriormente, la arquitectura limitaba el número de subprocesos en un bloque de subprocesos a un total de 512 subprocesos por bloque, pero a partir de marzo de 2010, con capacidad de cálculo 2.x y superior, los bloques pueden contener hasta 1024 subprocesos. Los subprocesos del mismo bloque de subprocesos se ejecutan en el mismo procesador de flujo. [1] Los subprocesos en el mismo bloque pueden comunicarse entre sí a través de memoria compartida , sincronización de barrera u otras primitivas de sincronización, como operaciones atómicas.
Se combinan varios bloques para formar una cuadrícula. Todos los bloques de la misma cuadrícula contienen el mismo número de hilos. La cantidad de subprocesos en un bloque es limitada, pero las cuadrículas se pueden usar para cálculos que requieren una gran cantidad de bloques de subprocesos para operar en paralelo y usar todos los multiprocesadores disponibles.
CUDA es una plataforma informática paralela y un modelo de programación que los lenguajes de nivel superior pueden utilizar para explotar el paralelismo. En CUDA, el kernel se ejecuta con la ayuda de subprocesos. El hilo es una entidad abstracta que representa la ejecución del kernel. Un kernel es una función que se compila para ejecutarse en un dispositivo especial. Las aplicaciones de subprocesos múltiples utilizan muchos de estos subprocesos que se ejecutan al mismo tiempo para organizar el cálculo paralelo. Cada hilo tiene un índice, que se utiliza para calcular las ubicaciones de las direcciones de memoria y también para tomar decisiones de control.
CUDA opera en un modelo de programación heterogéneo que se utiliza para ejecutar programas de aplicaciones del dispositivo host. Tiene un modelo de ejecución similar a OpenCL . En este modelo, comenzamos a ejecutar una aplicación en el dispositivo host, que suele ser un núcleo de CPU . El dispositivo es un dispositivo orientado al rendimiento, es decir, un núcleo GPU que realiza cálculos paralelos. Las funciones del kernel se utilizan para realizar estas ejecuciones paralelas. Una vez que se ejecutan estas funciones del kernel, el control se devuelve al dispositivo host que reanuda la ejecución en serie.
Como muchas aplicaciones paralelas involucran datos multidimensionales, es conveniente organizar bloques de subprocesos en matrices de subprocesos 1D, 2D o 3D. Los bloques de una cuadrícula deben poder ejecutarse de forma independiente, ya que la comunicación o cooperación entre bloques de una cuadrícula no es posible. 'Cuando se inicia un kernel, se especifica el número de subprocesos por bloque de subprocesos y el número de bloques de subprocesos, esto, a su vez, define el número total de subprocesos CUDA iniciados. [2] ' Las dimensiones máximas x, y y z de un bloque son 1024, 1024 y 64, y deben asignarse de modo que x × y × z ≤ 1024, que es el número máximo de subprocesos por bloque. [3] Los bloques se pueden organizar en cuadrículas de una, dos o tres dimensiones de hasta 2 31 -1, 65,535 y 65,535 bloques en las dimensiones x, y y z respectivamente. [3] A diferencia del máximo de hilos por bloque, no existe un límite de bloques por cuadrícula distinto de las dimensiones máximas de la cuadrícula.
Cada hilo en CUDA está asociado con un índice particular para que pueda calcular y acceder a ubicaciones de memoria en una matriz.
Considere un ejemplo en el que hay una matriz de 512 elementos. Una de las estructuras de la organización es tomar una cuadrícula con un solo bloque que tiene 512 subprocesos. Considere que hay una matriz C de 512 elementos que está hecha de la multiplicación por elementos de dos matrices A y B, que tienen 512 elementos cada una. Cada hilo tiene un índice i y realiza la multiplicación del i ésimo elemento de A y B y luego almacena el resultado en el i ésimo elemento de C. i se calcula usando blockIdx (que es 0 en este caso ya que solo hay uno block), blockDim (512 en este caso ya que el bloque tiene 512 elementos) y threadIdx que varía de 0 a 511 para cada bloque.
El índice del hilo i se calcula mediante la siguiente fórmula:
blockIdx.x es el identificador del bloque de dimensión x
blockDim.x es la dimensión x de la dimensión del bloque
threadIdx.x es la dimensión x del identificador del hilo
Por lo tanto, 'i' tendrá valores que oscilan entre 0 y 511 y cubren toda la matriz.
Si queremos considerar cálculos para una matriz mayor que 1024, podemos tener varios bloques con 1024 subprocesos cada uno. Considere un ejemplo con 2048 elementos de matriz. En este caso tenemos 2 bloques de hilos con 1024 hilos cada uno. Por lo tanto, los valores de los identificadores de subprocesos variarán de 0 a 1023, el identificador de bloque variará de 0 a 1 y la dimensión del bloque será 1024. Por lo tanto, el primer bloque obtendrá valores de índice de 0 a 1023 y el último tendrá valores de índice. de 1024 a 2047.
Así, cada hilo calculará primero el índice de memoria al que tiene que acceder y luego procederá con el cálculo. Considere un ejemplo en el que los elementos de las matrices A y B se agregan en paralelo mediante el uso de subprocesos y los resultados se almacenan en una matriz C. El código correspondiente en un subproceso se muestra a continuación: [5]
__global__ void vecAddKernel ( flotante * A , flotante * B , flotante * C , int n ) { int index = blockIdx . x * bloqueDim . x + hiloIdx . X ; si ( índice < n ) { C [ índice ] = A [ índice ] + B [ índice ] ; } } De la misma manera, en grillas particularmente complejas, cada hilo debe calcular el blockId y el threadId dependiendo de la geometría de la grilla. Considere una cuadrícula bidimensional con bloques bidimensionales. El threadId y el blockId se calcularán mediante las siguientes fórmulas:
[6]
Aunque hemos establecido la jerarquía de subprocesos, debemos tener en cuenta que los subprocesos, los bloques de subprocesos y la cuadrícula son esencialmente la perspectiva de un programador. Para obtener una idea completa del bloque de subprocesos, es fundamental conocerlo desde la perspectiva del hardware. El hardware agrupa subprocesos que ejecutan la misma instrucción en deformaciones. Varias urdimbres constituyen un bloque de hilo. Se asignan varios bloques de subprocesos a un Streaming Multiprocessor (SM). Varios SM constituyen toda la unidad GPU (que ejecuta todo el Kernel Grid). [ cita necesaria ]
Cada arquitectura en GPU (digamos Kepler o Fermi ) consta de varios SM o multiprocesadores de streaming. Estos son procesadores de uso general con un objetivo de frecuencia de reloj baja y un caché pequeño. Un SM es capaz de ejecutar varios bloques de subprocesos en paralelo. Tan pronto como uno de sus bloques de subprocesos haya completado su ejecución, ocupará el siguiente bloque de subprocesos en serie. En general, los SM admiten el paralelismo a nivel de instrucción pero no la predicción de bifurcaciones . [8]
Para lograr este propósito, un SM contiene lo siguiente: [8]
El hardware programa bloques de subprocesos en un SM. En general, un SM puede manejar múltiples bloques de subprocesos al mismo tiempo. Un SM puede contener hasta 8 bloques de subprocesos en total. Un ID de subproceso es asignado a un subproceso por su SM respectivo.
Siempre que un SM ejecuta un bloque de subprocesos, todos los subprocesos dentro del bloque de subprocesos se ejecutan al mismo tiempo. Por lo tanto, para liberar la memoria de un bloque de subprocesos dentro del SM, es fundamental que todo el conjunto de subprocesos del bloque haya concluido su ejecución. Cada bloque de hilo se divide en unidades programadas conocidas como urdimbre. Estos se analizan en detalle en la siguiente sección.
El programador warp de SM decide cuál de los warp tiene prioridad durante la emisión de instrucciones. [11] Algunas de las políticas de priorización de warp también se han discutido en las siguientes secciones.
En el lado del hardware, un bloque de hilos se compone de 'urdimbres'. (Este término proviene del tejido . [12] ) Una urdimbre es un conjunto de 32 hilos dentro de un bloque de hilos de manera que todos los hilos en una urdimbre ejecutan la misma instrucción. Estos hilos son seleccionados en serie por el SM. [13]
Una vez que se lanza un bloque de subprocesos en un multiprocesador (SM), todos sus warps permanecen residentes hasta que finaliza su ejecución. Por lo tanto, no se lanza un nuevo bloque en un SM hasta que haya suficiente número de registros libres para todos los warps del nuevo bloque y hasta que haya suficiente memoria compartida libre para el nuevo bloque.
Considere una urdimbre de 32 hilos ejecutando una instrucción. Si uno o ambos operandos no están listos (por ejemplo, aún no han sido recuperados de la memoria global), tiene lugar un proceso llamado ' cambio de contexto ' que transfiere el control a otro warp. [14] Al salir de una deformación en particular, todos los datos de esa deformación permanecen en el archivo de registro para que pueda reanudarse rápidamente cuando sus operandos estén listos. Cuando una instrucción no tiene dependencias de datos pendientes, es decir, ambos operandos están listos, se considera que el warp respectivo está listo para su ejecución. Si más de un warp es elegible para ejecución, el SM principal usa una política de programación de warp para decidir qué warp recibe la siguiente instrucción recuperada.
A continuación se analizan diferentes políticas para programar warps que son elegibles para ejecución: [15]
El "cambio" de contexto de subproceso de CPU tradicional requiere guardar y restaurar los valores de registro asignados y el contador del programa en la memoria fuera del chip (o caché) y, por lo tanto, es una operación mucho más pesada que con el cambio de contexto warp. Todos los valores de registro de un warp (incluido su contador de programa) permanecen en el archivo de registro, y la memoria compartida (y el caché) también permanecen en su lugar, ya que se comparten entre todos los warps en el bloque de subprocesos.
Para aprovechar la arquitectura warp, los lenguajes de programación y los desarrolladores deben comprender cómo fusionar los accesos a la memoria y cómo gestionar la divergencia del flujo de control. Si cada hilo en un warp toma una ruta de ejecución diferente o si cada hilo accede a una memoria significativamente divergente, entonces los beneficios de la arquitectura warp se pierden y el rendimiento se degradará significativamente.
Las GPU NVIDIA ejecutan grupos de subprocesos conocidos como warps en forma SIMT (Instrucción única, subprocesos múltiples)
