Accurate Models of NVIDIA Tensor Cores

Título: Modelos Precisos de Tensor Cores da NVIDIA

Autores: Faizan A. Khattak e Mantas Mikaitis (Universidade de Leeds, Reino Unido)

---

1. O Problema

A multiplicação de matrizes é uma operação fundamental para o treinamento de redes neurais e inferência, além de ser amplamente utilizada em computação científica de alto desempenho (HPC). As GPUs modernas utilizam unidades especializadas, conhecidas como Tensor Cores (ou Matrix Cores), para acelerar essas operações.

No entanto, existem desafios críticos:

• Não Conformidade com IEEE 754: Os Tensor Cores não seguem estritamente o padrão IEEE 754 de aritmética de ponto flutuante. Diferentes fabricantes e gerações de hardware implementam comportamentos numéricos distintos (modos de arredondamento, largura do acumulador, pontos de normalização, bits extras de alinhamento, etc.).
• Falta de Reprodutibilidade: Essas diferenças levam a resultados não reprodutíveis entre diferentes arquiteturas de GPU (ex: V100 vs. A100 vs. H100) e entre diferentes fornecedores.
• Limitações dos Modelos Atuais: Os modelos existentes são frequentemente aproximados ou baseados em vetores de teste manuais que podem não distinguir entre diferentes modelos de hardware, especialmente devido à escassez de acesso a múltiplas plataformas físicas para validação.

2. Metodologia

Os autores desenvolveram uma abordagem iterativa para criar modelos de software precisos (bit-accurate) que emulam o comportamento interno dos Tensor Cores da NVIDIA. O processo envolve duas etapas principais:

1. Teste Generalizado de Recursos Numéricos (GNFT - Generalised Numerical Feature Testing):

   • Utiliza vetores de teste cuidadosamente projetados para disparar casos especiais e revelar características numéricas específicas (ex: suporte a números subnormais, comportamento de arredondamento, tamanho do bloco FMA).
   • Este método fornece uma aproximação inicial do modelo, identificando parâmetros como o número de bits de alinhamento extra ($n_{eab}$) e o tamanho do bloco de multiplicação-acumulação ($N_{FMA}$).

2. Busca no Espaço de Entrada (ISSM - Input Space Search Method):

   • Para refinar o modelo inicial, o ISSM gera milhões de vetores de entrada aleatórios (distribuições normal e uniforme) e compara os resultados do modelo de software com os resultados reais obtidos no hardware da GPU.
   • Quando há discrepâncias (mismatches), os autores analisam os casos de falha para identificar recursos numéricos não detectados inicialmente (como comportamentos de subnormais ou padrões de intercalação de dados).
   • O modelo é ajustado manualmente (por especialistas) e o processo se repete até que o modelo seja bit-accurate (produza resultados idênticos ao hardware em todos os testes).

3. Principais Contribuições

• Modelos MATLAB Precisos: Desenvolvimento de uma toolbox MATLAB (versão v0.4.1) que contém modelos de Tensor Cores para as GPUs V100, A100, A2, A30, H100, H200, B200, L40S e RTX 1000 (Ada).
• Cobertura de Formatos: Os modelos cobrem formatos de entrada de baixa precisão (8-bit, 16-bit, 19-bit/Tf19) e mistos, incluindo os novos formatos de ponto flutuante de 8 bits (fp8-E4M3 e fp8-E5M2) e 6/4 bits (embora estes últimos não tenham sido detalhados em profundidade no estudo).
• Descoberta de Novos Comportamentos:
  • Identificação de que os produtos parciais permanecem desnormalizados durante a alinhamento e acumulação em várias arquiteturas.
  • Detecção de um limite de expoente máximo fixo (-133) para alinhamento de mantissas quando se somam produtos subnormais pequenos.
  • Revelação de um padrão de intercalação (interleaved) no processamento de entradas fp8 nas arquiteturas Hopper (H100/H200) e Blackwell (B200) quando acessadas via instruções mma.sync, onde os dados são distribuídos entre duas chamadas de tensor core.
  • Confirmação de que, em certas instruções de fp8 na arquitetura Hopper, o hardware converte internamente para fp16 antes de processar.
• Modelo Personalizável: Fornecimento de um modelo genérico que permite aos pesquisadores simular configurações arbitrárias (alterando tamanho FMA, bits de alinhamento, modos de arredondamento), facilitando a exploração de futuros padrões de hardware.

4. Resultados

• Validação Estatística: Os modelos foram validados através de testes randomizados massivos (até $10^8$ amostras para algumas configurações). Em todos os casos, os modelos de software produziram resultados idênticos ao nível de bits em comparação com o hardware real.
• Detalhamento de Arquiteturas:
  • V100: $N_{FMA}=4$, alinhamento de 25 bits (2 inteiros, 23 fracionários), sem bits extras de alinhamento ($n_{eab}=0$).
  • A100/A2/A30: $N_{FMA}=8$ (fp16/bf16), $n_{eab}=1$, alinhamento de 26 bits.
  • H100/H200/B200: $N_{FMA}=16$ (fp16/bf16), $n_{eab}=2$, alinhamento de 27 bits. Para fp8, a arquitetura Hopper usa um padrão de intercalação, enquanto a Blackwell (B200) usa um modelo de 5ª geração com $N_{FMA}=32$ e 23 bits fracionários para fp8 nativo.
• Aplicação em Algoritmos Multi-Word: Os autores demonstraram o uso dos modelos em algoritmos de multiplicação de matrizes de alta precisão (multi-word) usando Tensor Cores de baixa precisão. Os resultados mostraram que as diferenças nos modelos (como o número de bits extras de alinhamento e o modo de arredondamento) impactam significativamente a precisão final, especialmente em problemas de grande escala.

5. Significância

• Reprodutibilidade Científica: Permite que cientistas computacionais e desenvolvedores de algoritmos de precisão mista verifiquem e validem seus testes e algoritmos em software antes de executá-los em hardware físico, garantindo consistência entre diferentes gerações de GPUs.
• Análise de Erros: Facilita a análise rigorosa de erros de arredondamento em algoritmos que dependem de Tensor Cores, algo que era anteriormente difícil devido à "caixa preta" do hardware.
• Padronização Futura: Ao documentar e modelar as diferenças entre fabricantes e gerações, este trabalho fornece dados cruciais para futuras iniciativas de padronização (como o IEEE P3109 e o Open Compute Project), ajudando a definir comportamentos esperados para aritmética de baixa precisão.
• Acesso Democratizado: A toolbox MATLAB torna acessível a emulação precisa de hardware de ponta (como B200 e H100) para pesquisadores que não possuem acesso físico a essas máquinas.

O código fonte e os modelos estão disponíveis publicamente no repositório GitHub dos autores, integrados à biblioteca MATLAB-tensor-core.