- Biblioteca de comunicação de alto desempenho para Mixture-of-Experts (MoE) e Expert Parallelism (EP)
- Fornece kernel All-to-All baseado em GPU para processar rapidamente operações de despacho e combinação de MoE
- Suporte a operações de baixa precisão como FP8
- Aplica o algoritmo de group-limited gating proposto no artigo do DeepSeek-V3 para otimizar o encaminhamento assimétrico da largura de banda entre domínios
- Ex.: otimização de transferência de dados NVLink → RDMA
- Oferece alta taxa de transferência adequada para tarefas de treinamento e prefilling de inferência
- Inclui kernel dedicado de baixa latência para RDMA voltado à decodificação de inferência sensível à latência
- Fornece técnica de sobreposição entre comunicação e computação (sem ocupar recursos de SM)
Desempenho
Kernel geral (transferências via NVLink e RDMA)
- O DeepEP testou o desempenho em um ambiente com GPU H800 e rede RDMA InfiniBand CX7 de 400Gb/s
- Com base na configuração do DeepSeek-V3/R1, foi aplicada uma estrutura com 4096 tokens por batch, 7168 nós ocultos, top-4 grupos e top-8 especialistas, usando despacho em FP8 e combinação em BF16
- Nos testes de desempenho, a comunicação intra-nó (baseada em NVLink) apresentou largura de banda acima de cerca de 150GB/s, enquanto a comunicação entre nós (baseada em RDMA) registrou largura de banda de 40~47GB/s, dependendo do número de especialistas
- À medida que o número de especialistas aumentava, houve tendência de leve aumento na largura de banda RDMA (ex.: 43GB/s com 16 especialistas, 46GB/s com 64 especialistas)
Kernel de baixa latência (RDMA puro)
- A medição de desempenho do kernel de baixa latência mostrou que a latência foi significativamente reduzida em comparação com o kernel geral
- Em um ambiente que processa 128 tokens por batch, a latência aumentou conforme o número de especialistas crescia, mas a largura de banda RDMA permaneceu relativamente estável
- Por exemplo, foi de 163 microssegundos (us) com 8 especialistas para 194 microssegundos (us) com 256 especialistas
- Na operação de combinação (
combine), ocorreu latência maior do que no despacho, e a largura de banda RDMA tendeu a cair gradualmente para abaixo de 40GB/s à medida que o número de especialistas aumentava
- Ou seja, o kernel de baixa latência funciona muito rapidamente em grupos pequenos de especialistas, mas a latência aumenta quando o número de especialistas cresce, exigindo um equilíbrio adequado
Configuração de rede
Isolamento de tráfego (Traffic Isolation)
- É possível isolar o tráfego usando Virtual Lanes (VL) do InfiniBand
- Forma de separação recomendada:
- tarefas que usam o kernel geral
- tarefas que usam o kernel de baixa latência
- outras tarefas
- É possível configurar VL por meio da variável de ambiente
NVSHMEM_IB_SL
Roteamento adaptativo (Adaptive Routing)
- Suporta roteamento adaptativo em switches InfiniBand
- Pode ser ativado no kernel de baixa latência, mas precisa permanecer desativado no kernel geral (há risco de corrupção de dados se ativado)
- Recomendações de configuração:
- Quando a carga de rede for alta: ativar roteamento adaptativo
- Quando a carga de rede for baixa: manter roteamento estático
Controle de congestionamento (Congestion Control)
- O DeepEP opera com o controle de congestionamento desativado
- Foi confirmado que, no ambiente real, o congestionamento de rede não era grave
Principais considerações técnicas
- Uso de instrução PTX não oficial: utiliza
ld.global.nc.L1::no_allocate.L2::256B para otimização de desempenho
- Na arquitetura Hopper, funciona normalmente, mas em outras plataformas pode ser desativada com
DISABLE_AGGRESSIVE_PTX_INSTRS=1
- Autotuning recomendado: para obter o melhor desempenho, é necessário aplicar as configurações após testes de desempenho em cada cluster
Ainda não há comentários.