Bloque de hilos (programación CUDA)

En CUDA, cada hilo está asociado a un índice específico para poder calcular y acceder a las ubicaciones de memoria en un array.

Consideremos un ejemplo con una matriz de 512 elementos. Una de las estructuras organizativas consiste en una cuadrícula con un único bloque de 512 hilos. Supongamos que existe una matriz C de 512 elementos, formada por la multiplicación elemento a elemento de dos matrices A y B, ambas de 512 elementos cada una. Cada hilo tiene un índice i y realiza la multiplicación del i- ^ésimo elemento de A y B, almacenando el resultado en el i ^-ésimo elemento de C. El índice i se calcula utilizando blockIdx (que es 0 en este caso, ya que solo hay un bloque), blockDim (512 en este caso, dado 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  :

${\displaystyle i=blockIdx.x*blockDim.x+threadIdx.x}$

blockIdx.x es el identificador del bloque de la 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 van desde 0 hasta 511, lo que cubre todo el array.

Si queremos realizar cálculos con un array de más de 1024 elementos, podemos usar varios bloques con 1024 hilos cada uno. Consideremos un ejemplo con un array de 2048 elementos. En este caso, tenemos dos bloques de hilos con 1024 hilos cada uno. Por lo tanto, los identificadores de los hilos variarán de 0 a 1023, el identificador del bloque variará de 0 a 1 y la dimensión del bloque será de 1024. Así, el primer bloque tendrá índices de 0 a 1023 y el último tendrá índices de 1024 a 2047.

Thus each thread will first calculate the index of memory that it has to access and then proceed with the calculation. Consider an example in which elements from arrays A and B are added in parallel by using threads and the results is stored in an array C. The corresponding code in a thread is shown below :^[5]

__global__voidvecAddKernel(float*A,float*B,float*C,intn){intindex=blockIdx.x*blockDim.x+threadIdx.x;if(index<n){C[index]=A[index]+B[index];}}

In the same way in particularly complex grids, the blockId as well as the threadId need to be calculated by each thread depending on geometry of the grid. Consider, a 2-dimensional Grid with 2-dimensional blocks. The threadId and the blockId will be calculated by the following formulae :

${\displaystyle blockId=blockIdx.x+blockIdx.y*gridDim.x;}$${\displaystyle threadId=blockId*(blockDim.x*blockDim.y)+(threadIdx.y*blockDim.x)+threadIdx.x;}$^[6]

Each architecture in GPU (say Kepler or Fermi) consists of several SM or Streaming Multiprocessors. These are general purpose processors with a low clock rate target and a small cache. An SM is able to execute several thread blocks in parallel. As soon as one of its thread blocks has completed execution, it takes up the serially next thread block. In general, SMs support instruction-level parallelism but not branch prediction.^[8]

Para lograr este propósito, un SM contiene lo siguiente: ^{[ 8 ]}

• Núcleos de ejecución. (unidades de punto flotante de precisión simple, unidades de punto flotante de doble precisión, unidades de funciones especiales (SFU)).
• Cachés:

1. Caché L1 (para reducir la latencia de acceso a la memoria).
2. Memoria compartida (para datos compartidos entre hilos).
3. Caché constante (para la difusión de lecturas desde una memoria de solo lectura ).
4. Caché de texturas (para agregar ancho de banda desde la memoria de texturas).

• Planificadores para warps. (Estos sirven para dar instrucciones a los warps según políticas de planificación específicas).
• Un número considerable de registros. (Un SM puede estar ejecutando una gran cantidad de hilos activos simultáneamente, por lo que es imprescindible contar con miles de registros).

El hardware programa los bloques de subprocesos en un SM. En general, un SM puede gestionar varios bloques de subprocesos simultáneamente. Un SM puede contener hasta 8 bloques de subprocesos en total. A cada subproceso se le asigna un ID mediante su SM correspondiente.

Cuando un SM ejecuta un bloque de hilos, todos los hilos que contiene se ejecutan simultáneamente. Por lo tanto, para liberar la memoria de un bloque de hilos dentro del SM, es fundamental que todos los hilos del bloque hayan finalizado su ejecución. Cada bloque de hilos se divide en unidades programadas conocidas como warp. Estas se describen en detalle en la siguiente sección.

El planificador de warps de SM decide cuál de los warps se prioriza durante la emisión de instrucciones. ^{[ 11 ]} Algunas de las políticas de priorización de warps también se han analizado en las siguientes secciones.

En el lado del hardware, un bloque de hilos se compone de "warps". (Este término proviene de " weaving" [ ^{12 ] )} . Un warp es un conjunto de 32 hilos dentro de un bloque de hilos. En el pasado, se garantizaba que estos hilos se ejecutaran "en bloque" (todos los hilos dentro del warp ejecutaban una instrucción simultáneamente) y, fundamentalmente, que accedieran a cada posición de memoria con todos o ningún hilo del warp. Este comportamiento podía provocar fácilmente interbloqueos (por ejemplo, al usar bifurcaciones if en bucles). Sin embargo, desde la arquitectura Volta , es posible el intercambio de datos dentro del warp mediante bloqueos de grano más fino ^{[ 13 ] [ 14 ] . Estos hilos son seleccionados en serie por el SM [ 15 ]} .

Una vez que se inicia un bloque de hilos en un multiprocesador (SM), todos sus subprocesos (warps) permanecen en ejecución hasta que finaliza su procesamiento. Por lo tanto, no se inicia un nuevo bloque en un SM hasta que haya suficientes registros libres para todos los subprocesos del nuevo bloque y hasta que haya suficiente memoria compartida libre para dicho bloque.

Consideremos un warp de 32 hilos ejecutando una instrucción. Si uno o ambos operandos no están listos (por ejemplo, no se han obtenido aún de la memoria global), se produce un proceso llamado " cambio de contexto " que transfiere el control a otro warp. ^{[ 16 ]} Al cambiar de un warp en particular, todos los datos de ese warp permanecen en el archivo de registros 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, el warp correspondiente se considera listo para la ejecución. Si más de un warp es apto para la ejecución, el SM padre utiliza una política de planificación de warps para decidir qué warp obtiene la siguiente instrucción obtenida.

A continuación se analizan diferentes políticas para programar los warps que son elegibles para su ejecución: ^{[ 17 ]}

1. Round Robin (RR): Las instrucciones se obtienen de forma rotativa. RR garantiza que los SM se mantengan ocupados y que no se desperdicien ciclos de reloj en latencias de memoria.
2. Menos Recientemente Recuperado (LRF): en esta política, el warp para el que no se ha recuperado una instrucción durante más tiempo tiene prioridad en la recuperación de una instrucción.
3. Fair (FAIR) ^{[ 17 ]} - En esta política, el planificador se asegura de que todos los warps tengan una oportunidad "justa" en cuanto al número de instrucciones que se les solicitan. Solicita instrucciones a un warp para el cual se ha solicitado el número mínimo de instrucciones.
4. CAWS basado en bloques de hilos ^{[ 18 ]} (planificación de warps con conciencia de criticidad): el énfasis de esta política de planificación radica en mejorar el tiempo de ejecución de los bloques de hilos. Asigna más recursos de tiempo al warp que tardará más tiempo en ejecutarse. Al dar prioridad al warp más crítico, esta política permite que los bloques de hilos finalicen más rápido, de modo que los recursos estén disponibles con mayor rapidez.

El cambio de contexto de los hilos de la CPU tradicional requiere guardar y restaurar los valores de los registros asignados y el contador de programa en la memoria externa (o caché), por lo que es una operación mucho más compleja que el cambio de contexto de los warps. Todos los valores de los registros de un warp (incluido su contador de programa) permanecen en el archivo de registros, y la memoria compartida (y la caché) también se mantienen, ya que se comparten entre todos los warps del bloque de hilos.

Para aprovechar la arquitectura de subprocesos (warp), los lenguajes de programación y los desarrolladores deben comprender cómo agrupar los accesos a memoria y cómo gestionar la divergencia del flujo de control. Si cada subproceso en un subproceso sigue una ruta de ejecución diferente o si accede a memoria significativamente divergente, se pierden las ventajas de esta arquitectura y el rendimiento se degrada considerablemente.