Accelerating GPU Inference of Large Language Models with Moderately Unstructured Sparse Weight Matrices

Resumo Técnico: Acelerando a Inferência de LLMs em GPU com Matrizes de Pesos Esparsas Moderadamente Não Estruturadas

Declaração do Problema
A implantação de Grandes Modelos de Linguagem (LLMs) enfrenta um gargalo crítico nos custos de inferência. Embora técnicas de poda de modelos (ex: Wanda, RIA) possam reduzir a carga computacional ao introduzir esparsidade em matrizes de pesos, elas são limitadas por requisitos de qualidade do modelo. Diferente de modelos de pequena escala, os LLMs são altamente sensíveis a níveis elevados de esparsidade; manter o desempenho tipicamente limita a poda a uma esparsidade não estruturada moderada (em torno de 50%).

Neste nível específico de esparsidade, os kernels de GPU existentes para Multiplicação de Matrizes Esparsas (SpMM) falham em superar as bibliotecas de multiplicação de matrizes densas como a cuBLAS. Essa lacuna de desempenho decorre de três desafios fundamentais:

1. Incompatibilidade com Tensor Cores: GPUs modernas oferecem Tensor Cores esparsos de alta eficiência, mas eles suportam estritamente a esparsidade estruturada 2:4. Eles não podem acelerar diretamente padrões de poda irregulares e não estruturados.
2. Sobrecarga de Metadados: Formatos de armazenamento tradicionais (ex: CSR) exigem o armazenamento de metadados posicionais completos para cada elemento não nulo. Com 50% de esparsidade, esse metadado pode ocupar tanto espaço quanto os próprios valores, anulando os benefícios da largura de banda de memória.
3. Gargalos de Decodificação On-Chip: Técnicas de compressão que reduzem metadados (ex: codificação de bitmap) deslocam a carga de trabalho para a decodificação on-chip. Em GPUs modernas com Memória de Alta Largura de Banda (HBM), o acesso à memória global não consegue sobrepor totalmente o trabalho de decodificação realizado pelos núcleos CUDA mais lentos, criando um travamento de desempenho (performance stall).

Metodologia
Os autores propõem uma solução holística composta por um novo formato de armazenamento de matriz de três camadas e um kernel SpMM co-otimizado projetado para esparsidade não estruturada moderada.

1. Formato de Armazenamento de Três Camadas
O formato proposto decompõe a matriz esparsa em três camadas distintas para equilibrar compatibilidade de hardware e eficiência de armazenamento:

• Camada Sparse-TC: Esta camada extrai padrões de esparsidade estruturada 2:4 da distribuição irregular. Ela retém dois valores não nulos de cada quatro pesos consecutivos. Como o padrão é fixo, requer apenas 2 bits por valor para metadados, garantindo compatibilidade direta com os sparse tensor cores sem overhead de decodificação.
• Camada de Preenchimento de Slots (Slot-Filling): Elementos não nulos excedentes que não se ajustam ao padrão 2:4 são redistribuídos nos slots vazios deixados pela camada Sparse-TC. Para minimizar os metadados para esses ~17,7% dos elementos, os autores introduzem a codificação de Distância Diferencial Paralela (PDD). Em vez de armazenar índices absolutos, a PDD codifica as distâncias relativas entre os elementos redistribuídos. A maioria das distâncias é pequena o suficiente para ser codificada em 4 bits. Para lacunas maiores, pontos fictícios (dummy points) são inseridos para quebrar o intervalo, garantindo que nenhuma distância individual exceda o limite de codificação.
• Camada Residual: Os elementos não nulos restantes (tipicamente <1% do total) que não podem ser tratados eficientemente pelas duas primeiras camadas são armazenados em um formato de Linha de Matriz Esparsa Comprimida (CSR) tradicional. Devido ao volume mínimo de dados, o overhead de metadados aqui é desprezível.

2. Kernel Co-Otimizado
Com base neste formato, os autores projetam um kernel de GPU que utiliza conjuntamente sparse tensor cores e núcleos CUDA dentro de um modelo de execução em pipeline:

• Decodificação Paralela: Os núcleos CUDA são encarregados de decodificar os metadados PDD da camada de Slot-Filling em paralelo. Como esta camada representa uma pequena fração do total de dados, o overhead de decodificação é gerenciado de forma eficaz.
• Aceleração Estruturada: Os sparse tensor cores processam a camada Sparse-TC usando instruções wgmma (arquitetura Hopper), aproveitando a estrutura 2:4 para alto rendimento (throughput).
• Tratamento de Resíduos: Os núcleos CUDA também processam a camada residual esparsa, escrevendo os resultados em um buffer on-chip separado para serem reduzidos com as saídas dos tensor cores ao final do pipeline.
• Sobreposição de Pipeline: O kernel implementa um pipeline de grão fino que sobrepõe três operações críticas:
  • Carregamento de matrizes de entrada densas (em cache na L2) e tiles de pesos esparsos (da Memória Global).
  • Decodificação de metadados PDD nos núcleos CUDA.
  • Execução de multiplicação de matrizes nos sparse tensor cores.
    Esta estratégia de "carregar-como-esparso, computar-como-esparso" garante que a computação on-chip se sobreponha ao acesso à memória global, maximizando a utilização da HBM.

Principais Contribuições

• Novo Formato de Armazenamento: Um formato de três camadas adaptado para esparsidade não estruturada moderada que permite a compatibilidade com sparse tensor cores, reduz drasticamente o overhead de metadados via codificação PDD e minimiza os custos de decodificação on-chip.
• Kernel Co-Otimizado: Um kernel SpMM que coordena eficientemente sparse tensor cores e núcleos CUDA, utilizando pipelining para sobrepor o acesso à memória com a computação.
• Primeiro a Superar o Processamento Denso: A implementação é a primeira a demonstrar que o SpMM esparso moderadamente não estruturado pode superar a multiplicação de matrizes densas em GPUs modernas equipadas com HBM.

Resultados
As avaliações foram conduzidas em GPUs NVIDIA H100 SXM5 usando os modelos OPT-30B e OPT-66B com níveis de esparsidade de 50%, 60% e 70%.

• Desempenho ao Nível de Kernel: Com 50% de esparsidade, o método proposto alcança até 1,64× de aceleração sobre o SpInfer (Best Paper do EuroSys'25) e supera significativamente outros baselines como SparTA e FlashLLM. Também supera o desempenho da cuBLAS densa.
• Inferência de Ponta a Ponta: O método entrega até 1,41× de aceleração de ponta a ponta sobre o FlashLLM (VLDB'24) e até 41,4% de maior throughput comparado ao FlashLLM no OPT-30B com um tamanho de lote (batch size) de 16.
• Eficiência de Memória: A abordagem reduz o uso de pico de memória global em aproximadamente 21,4% em comparação com os baselines densos com 50% de esparsidade, embora incorra em um uso de memória ligeiramente superior ao SpInfer devido à ausência de codificação de bitmap para metadados.
• Escalabilidade: O método mantém vantagens através dos níveis de esparsidade de 50% a 70%.

Significância e Alegações
O artigo afirma abordar uma lacuna crítica na otimização de inferência de LLMs: a incapacidade de kernels esparsos existentes em aproveitar efetivamente a esparsidade não estruturada moderada. Ao introduzir um formato de armazenamento que faz a ponte entre padrões de poda irregulares e aceleração de hardware estruturada, os autores demonstram que a poda pode gerar acelerações reais de inferência sem sacrificar a qualidade do modelo.

O trabalho destaca que, enquanto a esparsidade alta (>70%) degrada a qualidade do modelo, a esparsidade moderada (~50%) é o "ponto ideal" (sweet spot) para LLMs. A solução proposta desbloqueia o potencial de desempenho deste ponto ideal, que era anteriormente inacessível devido às ineficiências dos kernels esparsos existentes. Os autores observam que sua otimização é especificamente direcionada para a fase de decodificação da inferência de LLM (onde os tamanhos de lote são pequenos e limitados por memória), tornando-a altamente adequada para implantação escalável em ambientes de produção onde as fases de prefill e decode são frequentemente desacopladas. O código fonte é disponibilizado para facilitar a pesquisa e adoção futuras.