Accurate Models of NVIDIA Tensor Cores

1. El Problema

La multiplicación de matrices es una operación fundamental para el entrenamiento de redes neuronales y la inferencia. Para acelerar esto, las GPUs modernas (especialmente las de NVIDIA) utilizan unidades especializadas llamadas Tensor Cores (o núcleos de matriz). Sin embargo, existen problemas críticos de reproducibilidad y precisión:

• No conformidad con IEEE 754: Los Tensor Cores no siguen estrictamente el estándar IEEE 754 de aritmética de punto flotante. Diferentes fabricantes y generaciones de hardware implementan características numéricas distintas (modos de redondeo, anchos de acumuladores, puntos de normalización, bits de guardado adicionales, etc.).
• Falta de Reproducibilidad: Esto provoca que los resultados de las multiplicaciones de matrices varíen entre diferentes arquitecturas de GPU (ej. V100, A100, H100, B200) e incluso entre diferentes ejecuciones si no se controlan los detalles de implementación.
• Limitaciones de los Modelos Actuales: Los modelos existentes a menudo se basan en vectores de prueba manuales o teóricos que no distinguen suficientemente entre modelos de hardware o no han sido validados exhaustivamente en múltiples plataformas. Además, la disponibilidad de hardware para pruebas es limitada.

2. Metodología

Los autores desarrollaron un enfoque iterativo para crear modelos de software precisos que emulan el comportamiento de los Tensor Cores de NVIDIA. El proceso se divide en tres etapas principales:

1. Pruebas de Características Numéricas Generalizadas (GNFT):

   • Se utilizan vectores de prueba diseñados específicamente para activar casos especiales en la lógica de los multiplicadores (ej. subnormales, desbordamientos, alineación de significados).
   • El objetivo es determinar características clave como: el tamaño del bloque FMA ($N_{FMA}$), el número de bits de alineación extra ($n_{eab}$), los puntos de normalización y los modos de redondeo.

2. Búsqueda en el Espacio de Entrada (ISSM):

   • Una vez obtenida una aproximación inicial, se realiza una comparación semi-exhaustiva entre el modelo de software y el hardware real.
   • Se generan millones de vectores de entrada aleatorios (distribuciones normal y uniforme) para cubrir el espacio de entrada.
   • Si hay discrepancias entre el modelo y la GPU, se analizan los casos de fallo para refinar el modelo. Este ciclo se repite hasta lograr una precisión bit a bit.

3. Desarrollo del Modelo en MATLAB:

   • Se combinan aritmética de punto fijo con el simulador de formato de punto flotante personalizado CPFloat.
   • Se crean modelos específicos para las arquitecturas: V100 (Volta), A100/A2/A30 (Ampere), L40S/RTX 1000 (Ada Lovelace), H100/H200 (Hopper) y B200 (Blackwell).
   • Se soportan formatos de entrada de 8, 16 y 19 bits (fp8, fp16, bf16, tf19) y salidas en fp16/fp32.

3. Contribuciones Clave

• Determinación de Características Numéricas Precisas: Se han identificado y documentado las características numéricas exactas para 31 combinaciones de modelos de GPU y formatos de entrada/salida, incluyendo los nuevos formatos de 8 bits en las arquitecturas Hopper y Blackwell.
• Modelos de Software Validados: Se ha desarrollado una caja de herramientas en MATLAB (MATLAB Tensor Core v0.4.1) que emula el comportamiento interno de los Tensor Cores con precisión bit a bit.
• Descubrimiento de Comportamientos Ocultos:
  • Se identificó que en ciertas arquitecturas (H100/H200/B200 con entrada fp8), los productos se calculan en un patrón entrelazado (interleaved) cuando se usan instrucciones específicas (mma.sync), dividiendo los vectores de entrada entre dos invocaciones de Tensor Core.
  • Se descubrió que las instrucciones mma.sync en Hopper/Blackwell para fp8 a menudo se mapean internamente a Tensor Cores de fp16 en lugar de usar núcleos nativos de fp8, a diferencia de lo que sugiere la documentación de instrucciones QMMA.
  • Se caracterizó el comportamiento de los exponentes en rangos subnormales, donde el exponente máximo para la alineación se limita artificialmente (ej. a -133) en ciertas condiciones.
• Modelo Personalizable: La herramienta permite a los usuarios definir sus propias configuraciones (tamaño FMA, bits de alineación, modo de redondeo) para investigar el impacto de diferentes parámetros numéricos.

4. Resultados

• Precisión Bit a Bit: Tras la validación mediante millones de pruebas aleatorias, los modelos de software producen resultados idénticos a los del hardware real en todos los casos probados (incluyendo casos extremos como Inf y NaN).
• Diferencias entre Arquitecturas:
  • V100: Usa $N_{FMA}=4$ y 0 bits de alineación extra.
  • A100: Usa $N_{FMA}=8$ (fp16/bf16) y 1 bit de alineación extra.
  • H100/H200/B200: Usan $N_{FMA}=16$ (fp16/bf16) y 2 bits de alineación extra.
  • L40S/Ada: Comparten características con A100 pero soportan fp8 nativo con $N_{FMA}=16$ y 13 bits fraccionarios en el acumulador.
• Validación de Algoritmos Multi-Palabra: Se demostró el uso de los modelos para emular multiplicaciones de matrices de alta precisión utilizando aritmética de baja precisión (multi-word). Los resultados mostraron que las diferencias en los modelos (ej. bits de alineación extra o modos de redondeo) afectan significativamente el error de norma, especialmente en tamaños de problema grandes ($n \approx 10^6$).

5. Significado e Impacto

• Reproducibilidad Científica: Proporciona a los científicos computacionales y desarrolladores de algoritmos de precisión mixta una herramienta confiable para verificar la reproducibilidad de sus resultados sin necesidad de acceso inmediato a hardware específico.
• Análisis de Errores: Facilita el análisis preciso de errores de redondeo en algoritmos que dependen de Tensor Cores, algo que antes era imposible debido a la opacidad de la implementación de hardware.
• Estándares Futuros: Al documentar estas diferencias, el trabajo apoya los esfuerzos de estandarización (como IEEE P3109 y Open Compute Project) para definir comportamientos aritméticos más consistentes en hardware de IA.
• Accesibilidad: Al estar disponible en MATLAB y GitHub, democratiza el acceso a modelos de hardware de vanguardia para la investigación numérica, permitiendo simular arquitecturas futuras (como B200) antes de que estén ampliamente disponibles.

En resumen, este trabajo cierra la brecha entre la teoría de la aritmética de punto flotante y la implementación real en hardware de GPU, ofreciendo la primera suite de modelos de software que replican con exactitud el comportamiento interno de los Tensor Cores de NVIDIA a través de múltiples generaciones de hardware.