DeepSeek revela a biblioteca open source DeepEP para treino e inferência de MoE

 (github.com/deepseek-ai)
1 pontos por GN⁺ 2025-02-26 | 1 comentários | Compartilhar no WhatsApp
  • DeepEP é uma biblioteca de comunicação de alto desempenho focada em paralelismo de especialistas (EP) no treino e na inferência modernos de ML, oferecendo kernels GPU all-to-all para dispatch/combine de MoE e suporte a baixa precisão, como FP8
  • A versão V2 refatora totalmente o EP e entrega desempenho equivalente ou melhor que a V1 usando muito menos recursos de SM, além de trocar o backend de NVSHMEM para o mais leve NCCL Gin
  • Nos testes com configuração baseada na V3, medidos com 8K tokens por batch, hidden 7168, top 8 experts, dispatch em FP8 e combine em BF16, a V2 registrou até 1,3x mais desempenho de pico e até 4x menos uso de SM em relação à V1
  • Todos os kernels são compilados em tempo de execução como módulos JIT leves, sem exigir compilação CUDA na instalação, e a V2 unifica APIs de alta vazão e baixa latência em uma única interface ElasticBuffer
  • São necessários GPU Hopper SM90, Python 3.8+, CUDA 12.3+, PyTorch 2.10+, NCCL 2.30.4+, NVLink e rede RDMA entre nós; Engram, PP e CP são recursos experimentais

O que o DeepEP oferece

  • DeepEP (DeepEveryParallel) é uma biblioteca de comunicação de alto desempenho para treino e inferência de machine learning modernos
  • Atualmente, sua principal funcionalidade é o paralelismo de especialistas (Expert Parallelism, EP), com kernels GPU all-to-all de alta vazão e baixa latência para dispatch e combine de MoE
  • Suporta comunicação de baixa precisão, incluindo FP8
  • Também inclui primitivas experimentais para paralelismo de pipeline (PP), paralelismo de contexto (CP) e acesso remoto à memória (Engram)
  • Todos os kernels são compilados em tempo de execução como módulos JIT (Just-In-Time) leves, sem necessidade de compilação CUDA durante a instalação
  • Apesar do design leve, mira desempenho próximo ou superior ao limite de largura de banda do hardware em várias configurações

Principais mudanças da versão V2

  • A V2 é uma versão com refatoração completa do Expert Parallelism
    • Foi projetada para alcançar desempenho extremo com várias vezes menos recursos de SM do que a V1
    • Suporta domínios maiores de scale-up e scale-out
    • O backend foi trocado de NVSHMEM para o mais leve backend NCCL Gin
  • Os novos recursos incluem:
    • Compilação JIT completa
    • Backend NCCL Gin leve e apenas com headers
    • Reutilização de communicators NCCL existentes
    • Unificação das APIs de alta vazão e baixa latência do EPv2 em uma única interface ElasticBuffer
    • Novo layout de GEMM
    • Suporte a domínios maiores de scale-up e scale-out, até EP2048
    • Cálculo analítico automático do número de SMs e QPs, eliminando a necessidade de auto-tuning
    • Continuidade do suporte aos modos híbrido e direto
    • Em treino legado semelhante ao da V3, redução do uso de SM de 24 para 4~6, mantendo desempenho equivalente ou melhor
    • 0 SM Engram baseado em RDMA
    • 0 SM PP baseado em RDMA
    • 0 SM CP baseado em Copy Engine

Limitações e recursos em desenvolvimento

  • A V2 consome mais tamanho de buffer do que a V1
  • O EP de baixa latência com RDMA e 0 SM não é mais suportado
  • Engram, PP e CP são recursos experimentais
  • Os recursos em desenvolvimento incluem:
    • Elastic GPU & CPU buffers, um espaço contínuo de endereços virtuais com mapeamento misto de memória física de GPU e CPU
      • A direção é permitir Engram totalmente automático e transparente ou EP desbalanceado
    • Uso de EP replay para tratar desequilíbrio de carga e reduzir o tamanho de buffers intermediários
    • Implementação de atualizações all-gather e reduce-scatter para DP e TP
  • A documentação da V1 baseada em NVSHMEM está disponível em docs/legacy.md

Resultados de desempenho

  • Os testes seguiram uma configuração alinhada à V3 com as seguintes condições:
    • 8K tokens por batch
    • hidden dimension 7168
    • top 8 experts
    • dispatch em FP8
    • combine em BF16
  • Os principais resultados foram:
    • SM90, CX7, EP 8 x 2: dispatch 90 GB/s RDMA, combine 81 GB/s RDMA, 12 SM
    • SM90, CX7, EP 8 x 4: dispatch 61 GB/s RDMA, combine 61 GB/s RDMA, 6 SM
    • SM100, CX7, EP 8 x 2: dispatch 90 GB/s RDMA, combine 91 GB/s RDMA, 12 SM
    • SM100, EP 8: dispatch 726 GB/s NVLink, combine 740 GB/s NVLink, 64 SM
    • SM100, EP 8: dispatch 643 GB/s NVLink, combine 675 GB/s NVLink, 24 SM
  • As medições são de largura de banda lógica; por exemplo, os 90 GB/s de EP 8 x 2 incluem tráfego de local rank
  • A V2 alcança até 1,3x mais desempenho de pico e reduz em até 4x o número de SMs em comparação com a V1
  • Os resultados para configurações de EP maiores foram omitidos por enquanto, e os usuários são incentivados a executar seus próprios benchmarks
  • Pela experiência interna, espera-se que os kernels continuem saturando a largura de banda do hardware mesmo em escalas maiores
  • Os dados de desempenho da V1 estão em docs/legacy.md

Instalação e requisitos

  • Os requisitos são:
    • GPU Hopper SM90 ou arquitetura com suporte a SM90 PTX ISA
    • Python 3.8 ou superior
    • CUDA 12.3 ou superior para GPU SM90
    • PyTorch 2.10 ou superior
    • NCCL 2.30.4 ou superior
    • NVLink para comunicação dentro do nó
    • Rede RDMA para comunicação entre nós
  • Recomenda-se instalar o NCCL via pip para que o DeepEP possa detectá-lo automaticamente no ambiente Python
pip install "nvidia-nccl-cu13>=2.30.4" --no-deps
  • Para suporte aos métodos legados, também há dependência de NVSHMEM; as instruções de instalação estão no NVSHMEM Installation Guide
  • Exemplos de build e execução de testes em desenvolvimento:
python setup.py build
ln -s build/lib.linux-x86_64-cpython-38/deep_ep_cpp.cpython-38-x86_64-linux-gnu.so

python tests/elastic/test_ep.py
python tests/elastic/test_agrs.py
python tests/elastic/test_engram.py
python tests/elastic/test_pp.py
  • A instalação pode ser feita com o seguinte comando:
python setup.py install
  • Depois da instalação, é possível importar deep_ep em projetos Python

Interface centrada em ElasticBuffer

  • Na V2, todas as operações de EP foram unificadas sob uma única interface ElasticBuffer
    • Ela trata APIs de alta vazão e baixa latência na mesma interface
    • O buffer pode ser inicializado especificando diretamente a configuração do MoE
    • O número ideal de SMs e QPs é calculado analiticamente
  • O exemplo de inicialização do buffer usa ElasticBuffer.get_buffer_size_hint() para calcular o tamanho necessário e verificar se um buffer existente pode ser reutilizado
  • Ao criar um novo buffer, são especificados num_max_tokens_per_rank, hidden, num_topk, use_fp8_dispatch etc.
  • _buffer.get_theoretical_num_sms(num_experts, num_topk) retorna o número teórico de SMs a ser usado pelos kernels de comunicação
  • Se num_sms for informado diretamente nas chamadas dispatch e combine, esse valor sobrescreve o cálculo automático

Padrões de uso em treino, prefill e decoding

  • Em treino ou prefill de inferência, o MoE dispatch roteia tokens para o expert correspondente em todos os ranks
    • Suporta entradas em BF16 e FP8
    • O handle contém os metadados de roteamento necessários para a chamada posterior de combine
    • handle.num_recv_tokens_per_expert_list fornece a quantidade de tokens por expert necessária para a GEMM
  • O backward pass de MoE dispatch é tratado, na prática, por combine
  • O MoE combine reduz a saída dos experts de volta ao rank original
  • O backward pass de MoE combine é tratado, na prática, por dispatch
  • A sobreposição entre comunicação e computação é gerenciada pela interface EventOverlap
    • É possível executar computação independente enquanto a comunicação está em andamento
    • Antes de usar o resultado, sincroniza-se a compute stream com event.current_stream_wait()
  • O mesmo ElasticBuffer também é usado no decoding de inferência
    • Quando a decisão de gating não muda, é possível reutilizar os metadados de roteamento com cached_handle
    • Esse padrão evita recálculo de layout e sincronização com a CPU

Variáveis de ambiente e valores fixados no build

  • Configurações gerais
    • EP_BUFFER_DEBUG: exibe informações de depuração sobre inicialização do buffer, estimativa de SM e backend
    • EP_SUPPRESS_NCCL_CHECK: suprime a verificação de incompatibilidade de versão do NCCL
    • EP_AVOID_RECORD_STREAM: evita record_stream nos tensores de saída
    • EP_NUM_TOPK_IDX_BITS: sobrescreve o número de bits da codificação do índice top-k
  • Configurações de rede
    • EP_NIC_NAME: nome da NIC padrão usada para consultar propriedades, com valor padrão mlx5_0
    • EP_OVERRIDE_RDMA_SL: sobrescreve o índice de service level do RDMA
    • EP_DISABLE_GIN: desabilita o backend NCCL Gin
  • Configurações de JIT
    • EP_JIT_CACHE_DIR: diretório de cache dos kernels compilados, padrão $HOME/.deep_ep
    • EP_JIT_NVCC_COMPILER: caminho do compilador NVCC
    • EP_JIT_CPP_STANDARD: versão do padrão C++, padrão 20
    • EP_JIT_DUMP_PTX, EP_JIT_DUMP_SASS, EP_JIT_DUMP_ASM: configurações relacionadas a dump de saídas PTX e SASS
  • Algumas variáveis de ambiente se comportam de forma persistente
    • São capturadas no momento do build e incluídas como valores padrão do pacote instalado
    • Se não forem sobrescritas pelas variáveis de ambiente atuais no momento do import, esses valores padrão são aplicados automaticamente
    • As variáveis afetadas são EP_JIT_CACHE_DIR, EP_JIT_PRINT_COMPILER_COMMAND, EP_NUM_TOPK_IDX_BITS, EP_NCCL_ROOT_DIR
  • Mais detalhes estão em test_ep.py ou na documentação Python

Recomendações de configuração de rede

  • O DeepEP foi totalmente testado em redes InfiniBand
  • Em teoria, também é compatível com RDMA over Converged Ethernet, ou RoCE
  • Isolamento de tráfego
    • É suportado por Virtual Lanes do InfiniBand
    • Recomenda-se separar workloads de expert-parallel e outros workloads em virtual lanes diferentes
    • Na V2, a alocação de virtual lane pode ser controlada pelo argumento sl_idx ou pela variável de ambiente EP_OVERRIDE_RDMA_SL
  • Adaptive routing
    • É um recurso avançado de roteamento em que switches InfiniBand distribuem o tráfego de forma equilibrada entre múltiplos caminhos
    • Recomenda-se mantê-lo habilitado em todas as condições de carga de rede, mesmo com alguma latência adicional
  • Congestion control
    • Deve ser desabilitado por prejudicar a largura de banda máxima
    • Quando a congestão for inevitável, recomenda-se atribuir o workload a uma virtual lane de prioridade mais baixa
  • PCI atomic mode
    • Se o hardware suportar, recomenda-se configurar PCI_ATOMIC_MODE na NIC para melhorar o desempenho de operações atômicas RDMA
sudo mlxconfig -y -d mlx5_$i set PCI_ATOMIC_MODE=4

Branches experimentais e forks da comunidade

  • Branches experimentais
    • Zero-copy: remove cópias entre tensores PyTorch e buffers de comunicação, reduzindo significativamente o uso de SM em kernels comuns
    • Eager: usa um protocolo de baixa latência para eliminar a extra RTT latency adicionada por RDMA atomic OP
    • Hybrid-EP: nova implementação de backend com instruções TMA, suporte a uso mínimo de SM, domínios NVLink maiores, sobreposição refinada de comunicação e computação em single-batch, kernel PCIe e suporte a NVFP4
    • AntGroup-Opt: série de otimizações escrita pelo AntGroup Network Platform Department
    • Mori-EP: suporte ao modo de baixa latência em ROCm/GPUs AMD com backend baseado em MORI
    • nvDev: branch baseada na V2 com recursos CUDA mais recentes, incluindo Compute Fabric Transport
  • Forks da comunidade
    • uccl/uccl-ep: suporte para executar o DeepEP em GPUs heterogêneas como Nvidia e AMD, e NICs como EFA, Broadcom e CX7
    • Infrawaves/DeepEP_ibrc_dual-ports_multiQP: adiciona solução multi-QP e suporte a NIC dual-port ao transporte IBRC
    • antgroup/DeepXTrace: analisador de diagnóstico para encontrar slow ranks com eficiência e precisão
    • ROCm/mori: biblioteca de comunicação de próxima geração da AMD para workloads de IA sensíveis a desempenho, incluindo Wide EP, transferência de KVCache e Collectives

Licença e citação

  • O DeepEP V2 é construído sobre o backend Gin do NCCL
  • O código do repositório é distribuído sob a MIT License
  • O item de citação é DeepEP: an efficient expert-parallel communication library, com o ano indicado como 2025

Leituras relacionadas

1 comentários

 
GN⁺ 2025-02-26
Comentários do Hacker News

  • Encontraram e usaram a instrução PTX não documentada ld.global.nc.L1::no_allocate.L2::256B para obter desempenho extremo
    Essa instrução acessa memória volátil da GPU com o modificador PTX somente leitura não coerente .nc, portanto pode causar comportamento indefinido
    Porém, dizem que na arquitetura Hopper foi testado que, quando usada junto com .L1::no_allocate, a correção é garantida e o desempenho fica muito melhor

    • Na prática, a NVIDIA poderia mais tarde, em uma nova arquitetura, mudar sutilmente o comportamento dessa instrução fora da documentação e virar o jogo, de propósito ou não?

  • Parece uma criança entrando numa loja de doces
    Há muitos truques que levariam tempo demais para fazer engenharia reversa direito só a partir do artigo, e espero que os lançamentos desta semana abram um renascimento em que MoE seja usado como modelo acadêmico padrão

    • Vendo por esse ângulo, não entendo o que está acontecendo entre a prática real dos modelos de ponta e os modelos acadêmicos
      O primeiro grupo já é todo MoE desde o GPT-4, mas os modelos abertos, com exceção do DeepSeek V3 e do Mixtral, muitas vezes continuam sendo modelos densos

  • É impossível não gostar desta equipe
    Eles estão expandindo os limites do open source para todos

    • Do tipo escrito separado, como Open AI™
    • Na verdade, não é open source
      Para ver um modelo realmente open source, confira o OLMo 2 da AI2: https://allenai.org/blog/olmo2
      Eles de fato compartilham tudo o que é necessário para reproduzir o modelo, até os próprios dados
      No link acima, eles também dizem: “como a ciência aberta completa exige mais do que pesos públicos, temos o prazer de compartilhar com a comunidade mais ampla de modelagem de linguagem uma nova atualização do OLMo que inclui pesos, dados, código, receitas, checkpoints intermediários e modelos ajustados por instruções”

  • Zuckerberg deveria parar de afirmar que a Meta publica IA como open source
    Eles estão até fazendo comerciais de TV, mas na prática só liberam os pesos, sem código
    A única IA realmente open source é a DeepSeek

    • Estritamente falando, a DeepSeek também não é tão open source quanto o OLMo ou o Open Euro
      Porque não divulgou os dados
    • A DeepSeek claramente não é open source de verdade
      Para ser open source, teria que usar uma licença open source de verdade, como as listadas pela OSI, e compartilhar o código de pré e pós-treinamento, o código relacionado a tuning, o código de avaliação, tudo relacionado a segurança e censura e, provavelmente, todo o conjunto de dados de treinamento
      Caso contrário, não é possível reproduzir os pesos, e compartilhar pesos é parecido com compartilhar um programa compilado
      Pelo que sei, o único modelo competitivo realmente open source é o OLMo 2 da AI2: https://allenai.org/blog/olmo2
      Recentemente, eles também lançaram um app que faz inferência no próprio dispositivo, e isso também é open source: https://allenai.org/blog/olmoe-app
      Há também outro modelo chamado Tülu 3, que dizem ter desempenho melhor que o DeepSeek V3: https://allenai.org/blog/tulu-3-405B
    • A Meta vem refinando o PyTorch há mais de 10 anos
      Quase tudo que é necessário para treinar LLMs, incluindo tecnologia de ponta, está ali
      O que mais falta? As partes do código específicas da infraestrutura da Meta?
    • O PyTorch conta também?
    • Pesos públicos = um blob binário
      Voltamos ao modelo FREEWARE / SHAREWARE
      É assim que deveríamos chamar esses pesos “públicos”

  • Oferece 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 prefill de inferência, kernels de baixa latência para decoding de inferência, suporte nativo a dispatch em FP8 e controle flexível de recursos de GPU para sobreposição de computação e comunicação
    X: https://x.com/deepseek_ai/status/1894211757604049133

  • A motivação por trás do trabalho da DeepSeek pode estar errada
    Por exemplo, pode ser uma tentativa patrocinada pelo Estado de reduzir a zero a vantagem dos EUA em IA, mas o efeito líquido para todos no mundo é simplesmente excelente
    No pior caso, mesmo que estejam fazendo isso pelos motivos errados, sou grato à DeepSeek, pois ela está fazendo de fato o que a OpenAI mentiu durante anos dizendo que faria pelo mundo

    • No campo das relações internacionais, certo e errado não se aplicam tanto assim
      Tornar isso open source é mais “errado” do que a proibição de exportação de GPUs Nvidia avançadas?
      O lançamento open source da DeepSeek provavelmente é apenas algo que, com a concordância do Partido Comunista Chinês, acabou sendo bom ao mesmo tempo para o Partido Comunista Chinês e para a comunidade open source de IA mais ampla, e não deve ser tomado como uma posição de princípio
      Encontrar formas de eliminar a vantagem competitiva de outros países é uma atividade central de todos os governos, grandes e pequenos

  • É a segunda rodada de lançamentos open source sob licença MIT feita pela verdadeira empresa Open AI™
    Mais uma vez, a DeepSeek é mais aberta do que uma empresa de US$ 157 bilhões que se diz “Open”
    Quase ninguém fala sobre o Llama da Meta, mas todos deveriam esperar que o Llama 4 venha com capacidade de raciocínio
    O objetivo é não ser esmagado no meio de uma corrida rumo a zero

  • Enquanto os EUA vasculham recibos de GPUs em Singapura para confirmar se a DeepSeek usou apenas H800, o resto do mundo pode rodar essas otimizações em H100 completas?
    Porque, por causa das sanções dos EUA, era difícil obter ou acessar H100, e por causa da arrogância americana eles continuam fingindo acreditar que suas ordens valem para o mundo inteiro?
    Entendi certo?

  • O PTX que todos estavam esperando foi incluído desta vez?

    • Sim, há alguns no diretório csrc/kernels
      Pesquisando por asm, dá para encontrar onde foi usado
    • O resto das pessoas precisa de uma explicação sobre por que o PTX que todos estavam esperando é tão importante

  • Isso me lembra os anos 80 e 90, quando as pessoas hackeavam assembly ou procuravam instruções não documentadas para extrair desempenho da CPU
    Imagino que um dia os compiladores vão otimizar o suficiente, ou as GPUs ficarão tão poderosas, que esses truques não farão mais tanta diferença, como acontece hoje com CPUs