Capítulo 7: Diseño de Streaming Multiprocessor en el Diseño de GPU

El streaming multiprocessor (SM) es el bloque de construcción fundamental de las arquitecturas de GPU de NVIDIA. Cada SM contiene un conjunto de núcleos CUDA que ejecutan instrucciones en un estilo SIMT (Single Instruction, Multiple Thread). El SM es responsable de administrar y programar warps, manejar la divergencia de ramificaciones y proporcionar un acceso rápido a la memoria compartida y las memorias caché. En este capítulo, exploraremos la microarquitectura del SM, incluyendo sus tuberías, mecanismos de programación de warps, diseño del archivo de registros y organización de la memoria compartida y la memoria caché L1.

Microarquitectura y Tuberías del SM

El SM es un procesador altamente paralelo y con pipeline diseñado para ejecutar eficientemente cientos de hilos simultáneamente. La Figura 7.1 muestra un diagrama de bloques simplificado de un SM en la arquitectura NVIDIA Volta.

                                 Caché de Instrucciones
                                         |
                                         v
                                    Programador de Warps
                                         |
                                         v
                               Unidad de Despacho (4 warps)
                                 |   |   |   |
                                 v   v   v   v
                               Núcleo CUDA (FP64/FP32/INT)
                               Núcleo CUDA (FP64/FP32/INT)
                               Núcleo CUDA (FP64/FP32/INT)
                               ...
                               Tensor Core
                               Tensor Core
                               ...
                               Unidad de Carga/Almacenamiento
                               Unidad de Carga/Almacenamiento
                               ...
                               Unidad de Funciones Especiales
                                         ^
                                         |
                                Archivo de Registros (64 KB)
                                         ^
```Memoria compartida / Caché L1 (96 KB)

Figura 7.1: Diagrama de bloques simplificado de un SM en la arquitectura NVIDIA Volta.

Los principales componentes del SM incluyen:

Caché de instrucciones: Almacena instrucciones de acceso frecuente para reducir la latencia y mejorar el rendimiento.
Planificador de warps: Selecciona los warps que están listos para ejecutarse y los envía a las unidades de ejecución disponibles.
Unidad de despacho: Recupera y decodifica instrucciones para hasta 4 warps por ciclo y las envía a las unidades de ejecución apropiadas.
Núcleos CUDA: Unidades de ejecución programables que admiten una amplia gama de operaciones de enteros y coma flotante. Cada SM en Volta contiene 64 núcleos CUDA.
Núcleos de tensor: Unidades de ejecución especializadas diseñadas para acelerar las cargas de trabajo de aprendizaje profundo y IA. Cada SM en Volta contiene 8 núcleos de tensor.
Unidades de carga/almacenamiento: Manejan las operaciones de memoria, incluidas las cargas y almacenamientos en la memoria global, memoria compartida y cachés.
Unidades de funciones especiales: Ejecutan operaciones trascendentales y otras operaciones matemáticas complejas.
Registro de archivo: Proporciona un acceso rápido a los registros privados de los hilos. Cada SM en Volta tiene un registro de archivo de 64 KB.
Memoria compartida / Caché L1: Un espacio de memoria configurable que se puede usar como una caché administrada por software (memoria compartida) o como una caché de datos L1 administrada por hardware.

La tubería SM está diseñada para maximizar el rendimiento al permitir que múltiples warps se ejecuten concurrentemente y oculten la latencia de la memoria. La Figura 7.2 ilustra una vista simplificada de la tubería SM.

    Búsqueda de instrucciones
            |
            v
    Decodificación de instrucciones
            |
            v
    Recolección de operandos
            |
            v
    Ejecución (Núcleos CUDA, Núcleos de tensor, Unidades de carga/almacenamiento, Unidades de funciones especiales)
            |
            v
    Escritura de resultados

Figura 7.2: Tubería SM simplificada.

Las etapas de la tubería son las siguientes:

Búsqueda de instrucciones: El planificador de warps selecciona un warp que está listo para ejeAquí está la traducción al español del archivo con comentarios traducidos, pero sin traducir el código:

Rápido y recupera la siguiente instrucción para esa cuadrícula desde la caché de instrucciones.

Decodificación de instrucciones: La instrucción recuperada se decodifica para determinar el tipo de operación, los operandos y los registros de destino.
Recopilación de operandos: Los operandos requeridos para la instrucción se recopilan del archivo de registros o de la memoria compartida.
Ejecución: La instrucción se ejecuta en la unidad de ejecución apropiada (CUDA Core, Tensor Core, Unidad de Carga/Almacenamiento o Unidad de Función Especial).
Escritura de vuelta: El resultado de la ejecución se escribe de vuelta al archivo de registros o a la memoria compartida.

Para lograr un alto rendimiento, el SM emplea varias técnicas para maximizar la utilización de los recursos y ocultar la latencia:

Doble emisión: El SM puede emitir dos instrucciones independientes por cuadrícula en un solo ciclo, lo que permite una mayor paralelismo a nivel de instrucción.
Unidades de ejecución con pipeline: Las unidades de ejecución tienen pipeline, lo que permite al SM iniciar una nueva operación en una unidad antes de que la operación anterior haya finalizado.
Ocultación de latencia: El SM puede cambiar entre cuadrículas de forma cíclica, lo que le permite ocultar la latencia de los accesos a memoria y las operaciones de alta latencia ejecutando instrucciones de otras cuadrículas.

El Ejemplo 7.1 muestra un kernel CUDA simple que realiza la suma elemento a elemento de dos vectores.

__global__ void vectorAdd(int *a, int *b, int *c, int n) {
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    if (tid < n) {
        c[tid] = a[tid] + b[tid];
    }
}

Ejemplo 7.1: Kernel CUDA para la suma de vectores.

En este ejemplo, cada hilo en el kernel calcula la suma de los elementos correspondientes de los vectores de entrada a y b y almacena el resultado en el vector de salida c. El SM ejecuta este kernel asignando cada hilo a un CUDA core y programando cuadrículas de hilos para ejecutarse en los cores disponibles. Las unidades de carga/almacenamiento se utilizan para recuperar los datos de entrada de la memoria global y escribir los resultados de vuelta.

Programación de cuadrículas y manejo de divergencia

EfAquí está la traducción al español del archivo Markdown, con los comentarios traducidos pero sin traducir el código:

Una programación eficiente de warps es crucial para maximizar el rendimiento del SM. El planificador de warps es responsable de seleccionar los warps que están listos para ejecutarse y enviarlos a las unidades de ejecución disponibles. El objetivo principal del planificador de warps es mantener ocupadas las unidades de ejecución, asegurándose de que siempre haya warps disponibles para ejecutarse.

El SM emplea un mecanismo de planificación de warps de dos niveles:

Planificación de Warps: El planificador de warps selecciona los warps que están listos para ejecutarse en función de una política de planificación, como round-robin o first-in-first-out. Los warps seleccionados se envían entonces a las unidades de ejecución disponibles.
Planificación de Instrucciones: Dentro de cada warp, el SM planifica las instrucciones en función de sus dependencias y la disponibilidad de las unidades de ejecución. El SM puede emitir múltiples instrucciones independientes del mismo warp en un solo ciclo para maximizar el paralelismo a nivel de instrucción.

La Figura 7.3 ilustra el mecanismo de planificación de warps de dos niveles.

    Grupo de Warps
    Warp 1 (Listo)
    Warp 2 (Esperando)
    Warp 3 (Listo)
    ...
    Warp N (Listo)
        |
        v
    Planificador de Warps
        |
        v
    Unidad de Despacho
        |
        v
    Unidades de Ejecución

Figura 7.3: Mecanismo de planificación de warps de dos niveles.

Uno de los desafíos clave en la planificación de warps es manejar la divergencia de ramificación. En el modelo de ejecución SIMT, todos los hilos de un warp ejecutan la misma instrucción de manera sincronizada. Sin embargo, cuando un warp se encuentra con una instrucción de ramificación (por ejemplo, una declaración if-else), algunos hilos pueden tomar la ruta if mientras que otros toman la ruta else. Esta situación se conoce como divergencia de ramificación.

Para manejar la divergencia de ramificación, el SM emplea una técnica llamada predicación. Cuando un warp se encuentra con una ramificación divergente, el SM ejecuta ambas rutas de la ramificación de manera secuencial, enmascarando los hilos que no toman cada ruta. Los resultados se combinan entonces utilizando registros de predicado para asegurar que cada hilo reciba el resultado correcto.

El Ejemplo 7.2 muestra un kernel CUDA con una ramificación divergente.Aquí está la traducción al español del archivo Markdown, con los comentarios traducidos, pero sin traducir el código:

__global__ void divergentKernel(int *data, int *result) {
    // ID de hilo
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    // Si el dato es mayor que 0
    if (data[tid] > 0) {
        // Multiplica el dato por 2
        result[tid] = data[tid] * 2;
    } else {
        // Multiplica el dato por 3
        result[tid] = data[tid] * 3;
    }
}

Ejemplo 7.2: Kernel CUDA con una rama divergente.

En este ejemplo, la condición de la rama data[tid] > 0 puede hacer que algunos hilos en una warp tomen el camino del if mientras otros toman el camino del else. El SM maneja esta divergencia ejecutando ambos caminos secuencialmente y enmascarando los hilos inactivos en cada camino.

La Figura 7.4 ilustra el proceso de predicación para una warp con hilos divergentes.

    Warp (32 hilos)
    Hilo 1: data[1] = 5, result[1] = 10
    Hilo 2: data[2] = -3, result[2] = -9
    ...
    Hilo 32: data[32] = 7, result[32] = 14

    Rama divergente:
    if (data[tid] > 0) {
        result[tid] = data[tid] * 2;
    } else {
        result[tid] = data[tid] * 3;
    }

    Predicación:
    Paso 1: Ejecutar el camino del if con máscara
        Hilo 1: result[1] = 10
        Hilo 2: (enmascarado)
        ...
        Hilo 32: result[32] = 14

    Paso 2: Ejecutar el camino del else con máscara
        Hilo 1: (enmascarado)
        Hilo 2: result[2] = -9
        ...
        Hilo 32: (enmascarado)

    Resultado final:
    Hilo 1: result[1] = 10
    Hilo 2: result[2] = -9
    ...
    Hilo 32: result[32] = 14

Figura 7.4: Proceso de predicación para una warp con hilos divergentes.

Mediante el uso de la predicación, el SM puede manejar la divergencia de ramas sin la necesidad de instrucciones de rama explícitas o divergencia de flujo de control. Sin embargo, las ramas divergentes aún pueden afectar el rendimiento, ya que el SM debe ejecutar ambos caminos secuencialmente, reduciendo el paralelismo efectivo.

Registro de archivo y recolectores de operandos

El registro de archivo es un componente crítico del SM, que proporciona un acceso rápido a los registros privados de los hilos. Cada SM tiene un registro de archivo grande para admitir los muchos hilos activos y permitir un cambio de contexto eficiente entre warps.Aquí está la traducción al español del archivo markdown, con los comentarios de código traducidos:

En la arquitectura NVIDIA Volta, cada SM tiene un registro de archivo de 64 KB, organizado en 32 bancos de 2 KB cada uno. El registro de archivo está diseñado para proporcionar un alto ancho de banda y un acceso de baja latencia para admitir el gran número de hilos concurrentes.

Para minimizar los conflictos de banco y mejorar el rendimiento, el SM emplea una técnica llamada recolección de operandos. Los recolectores de operandos son unidades especializadas que reúnen los operandos de los bancos de registros y los entregan a las unidades de ejecución. Mediante el uso de recolectores de operandos, el SM puede reducir el impacto de los conflictos de banco y mejorar la utilización de las unidades de ejecución.

La Figura 7.5 muestra un diagrama simplificado del registro de archivo y los recolectores de operandos en un SM.

    Registro de Archivo (64 KB)
    Banco 1 (2 KB)
    Banco 2 (2 KB)
    ...
    Banco 32 (2 KB)
        |
        v
    Recolectores de Operandos
        |
        v
    Unidades de Ejecución

Figura 7.5: Registro de archivo y recolectores de operandos en un SM.

Los recolectores de operandos funcionan reuniendo operandos de múltiples instrucciones y múltiples warps, permitiendo al SM emitir instrucciones de diferentes warps a las unidades de ejecución en un solo ciclo. Esto ayuda a ocultar la latencia de los accesos al registro de archivo y mejora el rendimiento general del SM.

El Ejemplo 7.3 muestra un kernel CUDA que realiza un producto punto de dos vectores.

__global__ void dotProduct(float *a, float *b, float *result, int n) {
    // Cada hilo crea una suma parcial en la memoria compartida
    __shared__ float partialSum[256];
    int tid = threadIdx.x;
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    partialSum[tid] = 0;
 
    // Cada hilo calcula una parte del producto punto
    while (i < n) {
        partialSum[tid] += a[i] * b[i];
        i += blockDim.x * gridDim.x;
    }
 
    // Sincroniza los hilos del bloque
    __syncthreads();
 
    // Realiza una reducción en árbol para calcular la suma final
    for (int s = blockDim.x / 2; s > 0; s >>= 1) {
        if (tid < s) {
            partialSum[tid] += partialSum[tid + s];
        }
        __syncthreads();
    }
 
    // El hilo 0 de cada bloque almacena el resultado final
    if (tid == 0) {
        result[blockIdx.x] = partialSum[0];
    }
}

En este ejemplo, cada hilo calcula una suma parcial del producto punto utilizando su asignaciónAquí está la traducción al español del archivo Markdown proporcionado, con los comentarios del código traducidos al español:

Elementos de los vectores de entrada. Las sumas parciales se almacenan en la matriz de memoria compartida partialSum. Después de que todos los hilos hayan calculado sus sumas parciales, se realiza una reducción en paralelo para sumar las sumas parciales y obtener el resultado final del producto punto.

El recolector de operandos juega un papel crucial en este ejemplo al recoger eficientemente los operandos para los accesos a la memoria compartida y las operaciones aritméticas. Ayuda a evitar conflictos de banco y mejora la utilización de las unidades de ejecución.

Conclusión

El multiprocesador de flujo (streaming multiprocessor) es la unidad de cómputo central en las arquitecturas de GPU modernas. Su diseño se enfoca en maximizar el rendimiento y ocultar la latencia de memoria a través de una combinación de multihilos de grano fino, ejecución SIMT y recolección de operandos eficiente.

Los componentes clave del SM incluyen el planificador de warps, que selecciona los warps para su ejecución; la pila SIMT, que maneja la divergencia y convergencia de ramas; el archivo de registros y los recolectores de operandos, que proporcionan un acceso rápido a los registros privados de los hilos; y la memoria compartida y la caché L1, que permiten un intercambio y reutilización de datos de baja latencia.

A medida que las arquitecturas de GPU continúan evolucionando, la investigación en áreas como el manejo de la divergencia de ramas, la programación de warps y el diseño del archivo de registros será crucial para mejorar el rendimiento y la eficiencia de las futuras GPU. Técnicas novedosas como la formación dinámica de warps, la compactación de bloques de hilos y las memorias caché de reutilización de operandos tienen el potencial de mejorar significativamente las capacidades del SM y permitir nuevos niveles de rendimiento en cargas de trabajo de computación paralela.

Chapter 6 Gpu Performance Metrics and Analysis Chapter 8 Interconnect and on Chip Networks