- Biblioteca de comunicación de alto rendimiento para Mixture-of-Experts (MoE) y Expert Parallelism (EP)
- Proporciona un kernel All-to-All basado en GPU para procesar rápidamente las operaciones de despacho y combinación de MoE
- Soporta operaciones de baja precisión como FP8
- Aplica el algoritmo de compuerta con grupos limitados (group-limited gating) propuesto en el paper de DeepSeek-V3 para optimizar el reenvío de ancho de banda de dominio asimétrico
- Ejemplo: optimización de transferencia de datos NVLink → RDMA
- Ofrece alto rendimiento, adecuado para entrenamiento y tareas de prefilling de inferencia
- Incluye un kernel de baja latencia exclusivo para RDMA para decodificación de inferencia sensible a la latencia
- Proporciona una técnica de solapamiento entre comunicación y cómputo (sin ocupar recursos de SM)
Rendimiento
Kernel general (transferencia por NVLink y RDMA)
- DeepEP probó su rendimiento en un entorno con GPU H800 y una red RDMA InfiniBand CX7 de 400Gb/s
- Con base en la configuración de DeepSeek-V3/R1, se aplicó una estructura con 4096 tokens por lote, 7168 nodos ocultos, grupo top-4 y expertos top-8, usando despacho en FP8 y combinación en BF16
- Como resultado de las pruebas, la comunicación dentro del nodo (basada en NVLink) mostró un ancho de banda superior a ~150GB/s, mientras que la comunicación entre nodos (basada en RDMA) registró entre 40 y 47GB/s según la cantidad de expertos
- Se observó una ligera tendencia al aumento del ancho de banda RDMA a medida que crecía el número de expertos (por ejemplo, 43GB/s con 16 expertos y 46GB/s con 64 expertos)
Kernel de baja latencia (RDMA puro)
- Al medir el rendimiento del kernel de baja latencia, la latencia se redujo considerablemente frente al kernel general
- En un entorno que procesaba 128 tokens por lote, la latencia aumentó según la cantidad de expertos, pero el ancho de banda RDMA se mantuvo relativamente estable
- Por ejemplo, fue de 163 microsegundos (us) con 8 expertos y aumentó a 194 microsegundos (us) con 256 expertos
- En la operación de combinación (combine), la latencia fue mayor que en el despacho, y el ancho de banda RDMA tendió a disminuir gradualmente por debajo de 40GB/s a medida que aumentaba el número de expertos
- En otras palabras, el kernel de baja latencia funciona muy rápido con grupos pequeños de expertos, pero al aumentar la cantidad de expertos también aumenta la latencia, por lo que se necesita un equilibrio adecuado
Configuración de red
Aislamiento de tráfico (Traffic Isolation)
- Es posible aislar tráfico usando Virtual Lanes (VL) de InfiniBand
- Método de separación recomendado:
- Trabajos que usan el kernel general
- Trabajos que usan el kernel de baja latencia
- Otros trabajos
- Es posible configurar VL mediante la variable de entorno
NVSHMEM_IB_SL
Enrutamiento adaptativo (Adaptive Routing)
- Soporta enrutamiento adaptativo en switches InfiniBand
- Puede activarse en el kernel de baja latencia, pero debe desactivarse en el kernel general (si se activa, hay riesgo de corrupción de datos)
- Recomendaciones de configuración:
- Cuando la carga de red es alta: activar enrutamiento adaptativo
- Cuando la carga de red es baja: mantener enrutamiento estático
Control de congestión (Congestion Control)
- DeepEP opera con la función de control de congestión desactivada
- Se confirmó que en entornos reales la congestión de red no era grave
Consideraciones técnicas principales
- Uso de instrucciones PTX no oficiales: se utiliza
ld.global.nc.L1::no_allocate.L2::256B para optimizar el rendimiento
- En la arquitectura Hopper funciona normalmente, pero en otras plataformas puede desactivarse con
DISABLE_AGGRESSIVE_PTX_INSTRS=1
- Se recomienda autotuning: para lograr el mejor rendimiento, es necesario aplicar la configuración después de probar el desempeño en cada clúster
Aún no hay comentarios.