DeepSeek divulga a biblioteca open source DeepEP para treinamento e inferência com MoE

DeepEP

DeepEP é uma biblioteca de comunicação para Mixture-of-Experts (MoE) e expert parallelism (EP). Ela fornece kernels GPU all-to-all de alta velocidade e baixa latência, conhecidos como dispatch e combine de MoE. Também oferece suporte a operações de baixa precisão, incluindo FP8. De acordo com o algoritmo de gating com grupos restritos proposto no artigo do DeepSeek-V3, ela fornece kernels otimizados para encaminhamento com largura de banda assimétrica entre domínios, transferindo dados do domínio NVLink para o domínio RDMA. Esses kernels oferecem alta taxa de transferência, sendo adequados para treinamento e para o estágio de prefill na inferência. Também há suporte ao controle da quantidade de SMs (Streaming Multiprocessors). Para decoding de inferência sensível à latência, o DeepEP inclui kernels de baixa latência que usam RDMA puro para minimizar a latência. A biblioteca também introduz um método de sobreposição entre comunicação e computação baseado em hooks, que não ocupa recursos de SM.

Desempenho

Kernels gerais com encaminhamento via NVLink e RDMA

• No H800, os kernels gerais foram testados com largura de banda máxima de NVLink de cerca de 160 GB/s, conectados a uma placa de rede RDMA CX7 InfiniBand de 400 Gb/s (~50 GB/s de largura de banda máxima).
• Segue a configuração de pré-treinamento do DeepSeek-V3/R1 (4096 tokens por batch, hidden size 7168, top 4 grupos, top 8 experts, dispatch em FP8 e combine em BF16).

Kernels de baixa latência com RDMA puro

• No H800, os kernels de baixa latência foram testados com uma placa de rede RDMA CX7 InfiniBand de 400 Gb/s (~50 GB/s de largura de banda máxima).
• Segue a configuração típica de produção do DeepSeek-V3/R1 (128 tokens por batch, hidden size 7168, top 8 experts, dispatch em FP8 e combine em BF16).

Início rápido

Requisitos

• GPU Hopper (mais arquiteturas ou dispositivos poderão ser suportados no futuro)
• Python 3.8 ou superior
• CUDA 12.3 ou superior
• PyTorch 2.1 ou superior
• NVLink para comunicação dentro do nó
• Rede RDMA para comunicação entre nós

Baixar e instalar a dependência NVSHMEM

O DeepEP depende de uma versão modificada do NVSHMEM. É preciso instalá-la consultando o guia de instalação.

Configuração de rede

O DeepEP foi totalmente testado em rede InfiniBand e, em teoria, também é compatível com RDMA over Converged Ethernet (RoCE).

Isolamento de tráfego

O InfiniBand oferece suporte a isolamento de tráfego por meio de Virtual Lanes (VL). Para evitar interferência entre diferentes tipos de tráfego, recomenda-se separar as cargas nas virtual lanes da seguinte forma:

• Cargas que usam kernels gerais
• Cargas que usam kernels de baixa latência
• Outras cargas

No DeepEP, é possível controlar a atribuição de virtual lanes definindo a variável de ambiente NVSHMEM_IB_SL.

Roteamento adaptativo

O roteamento adaptativo é um recurso avançado de roteamento fornecido por switches InfiniBand, capaz de distribuir o tráfego de forma equilibrada entre múltiplos caminhos. Atualmente, os kernels de baixa latência suportam roteamento adaptativo, mas os kernels gerais ainda não (isso pode ser suportado em breve). Ativar roteamento adaptativo para kernels gerais entre nós pode causar deadlock ou corrupção de dados. No caso dos kernels de baixa latência, ativar o roteamento adaptativo pode eliminar completamente a congestão de rede causada por conflitos de rota, mas introduz latência adicional. Para melhor desempenho, recomenda-se a seguinte configuração:

• Ativar roteamento adaptativo em ambientes com alta carga de rede
• Usar roteamento estático em ambientes com baixa carga de rede

Controle de congestionamento

Como não foi observada congestão significativa em ambiente de produção, o controle de congestionamento fica desativado.

Interface e exemplos

Uso dos exemplos em treinamento de modelo ou prefill de inferência

Os kernels gerais podem ser usados no treinamento do modelo ou na etapa de prefill da inferência (sem a parte de backward).

Uso dos exemplos em decoding de inferência

Os kernels de baixa latência podem ser usados na etapa de decoding da inferência.

Observações

• Para desempenho extremo, foi identificado e usado o comando PTX não documentado ld.global.nc.L1::no_allocate.L2::256B. Esse comando provoca comportamento indefinido ao acessar memória volátil da GPU usando um modificador PTX de leitura somente não coerente. No entanto, ao testar com .L1::no_allocate na arquitetura Hopper, o desempenho melhorou bastante. Se os kernels não funcionarem em outras plataformas, isso pode ser desativado adicionando DISABLE_AGGRESSIVE_PTX_INSTRS=1 em setup.py, ou você pode abrir uma issue.
• Para melhor desempenho no cluster, recomenda-se executar todos os testes e usar a configuração ideal de autotuning. A configuração padrão foi otimizada no cluster interno da DeepSeek.

Licença

Este repositório de código é disponibilizado sob a licença MIT, e o código que faz referência ao NVSHMEM (incluindo csrc/kernels/ibgda_device.cuh e third-party/nvshmem.patch) está sujeito ao NVSHMEM SLA.

[GN⁺]

Comentários do Hacker News

• Descobriram e usaram instruções PTX não documentadas em busca de desempenho extremo. Essa instrução pode causar comportamento indefinido por acessar memória volátil da GPU usando um modificador PTX de somente leitura inconsistente. Porém, a corretude testada com .L1::no_allocate na arquitetura Hopper é garantida, e o desempenho deve ser muito melhor
• Zuckerberg deveria parar de afirmar que a Meta faz open source de IA. Eles estão liberando apenas os pesos, não o código. A única IA verdadeiramente open source é a da DeepSeek
• Me sinto como uma criança em uma loja de doces. Alguns desses truques levariam tempo demais para serem corretamente revertidos com base apenas em artigos. Espero que os anúncios desta semana iniciem um renascimento do uso de MoE como modelo acadêmico padrão
• Não tem como não gostar dessas pessoas. Elas estão realmente levando adiante as fronteiras do open source por todos nós. Obrigado por compartilhar
  • comunicação all-to-all eficiente e otimizada
  • suporte intra-nó e entre nós via NVLink e RDMA
  • kernels de alta vazão para treinamento e preenchimento prévio de inferência
  • kernels de baixa latência para decodificação de inferência
  • suporte nativo a despacho FP8
  • controle flexível de recursos de GPU para sobreposição entre computação e comunicação
• A motivação por trás do trabalho da DeepSeek pode estar errada (por exemplo, uma tentativa patrocinada por um Estado de eliminar a vantagem de liderança dos EUA em IA). Mas, globalmente, os resultados são simplesmente fantásticos
  • Mesmo no pior caso (se estiverem fazendo isso pelos motivos errados), ainda assim agradeço à DeepSeek. Eles estão de fato fazendo aquilo sobre o qual a OpenAI mentiu para o mundo durante anos
  • Absolutamente incrível
• Estou curioso se o PTX que todo mundo estava esperando foi incluído desta vez
• A instrução PTX mencionada no relatório técnico deveria estar ligada ao código aqui
• Enquanto os EUA rastreiam recibos de GPU em Singapura para verificar se a DeepSeek usou apenas H800, o resto do mundo pode executar essas otimizações em H100 completos
  • Fico me perguntando se isso é fingir que era difícil obter ou ter acesso a H100 por causa da arrogância de acreditar que as sanções dos EUA e suas ordens cobrem o mundo inteiro
• Este é o segundo lançamento open source da verdadeira empresa "Open AI™", e está sob licença MIT
  • A DeepSeek é mais aberta do que a empresa que afirma valer mais de $157B
  • Quase ninguém fala sobre a Llama da Meta, e todos deveriam esperar que lancem a Llama 4 com razão
  • O objetivo é não ficar preso no meio de uma corrida até o zero