DeepSeek publica la biblioteca DeepEP de código abierto para entrenamiento e inferencia de MoE
(github.com/deepseek-ai)DeepEP
DeepEP es una biblioteca de comunicación para Mixture-of-Experts (MoE) y expert parallelism (EP). Proporciona kernels GPU all-to-all de alta velocidad y baja latencia, conocidos como dispatch y combine de MoE. También admite operaciones de baja precisión, incluyendo FP8. En línea con el algoritmo de gating con grupos limitados propuesto en el paper de DeepSeek-V3, ofrece kernels optimizados para transferencias de ancho de banda asimétrico entre dominios, enviando datos desde el dominio NVLink al dominio RDMA. Estos kernels ofrecen un alto throughput, por lo que son adecuados para entrenamiento y tareas de prefill en inferencia. Además, admite control sobre la cantidad de SM (Streaming Multiprocessors). Para decoding de inferencia sensible a la latencia, DeepEP incluye kernels de baja latencia que usan RDMA puro para minimizar la latencia. Esta biblioteca también introduce un método de superposición entre comunicación y cómputo basado en hooks que no ocupa recursos de SM.
Rendimiento
Kernels generales usando transferencia por NVLink y RDMA
- En H800 se probó el kernel general con el ancho de banda máximo de NVLink de aproximadamente 160 GB/s, conectado a una tarjeta de red RDMA CX7 InfiniBand de 400 Gb/s (~50 GB/s de ancho de banda máximo).
- Sigue la configuración de preentrenamiento de DeepSeek-V3/R1 (4096 tokens por batch, 7168 hidden, top 4 grupos, top 8 expertos, dispatch en FP8 y combine en BF16).
Kernels de baja latencia usando RDMA puro
- En H800 se probó el kernel de baja latencia conectado a una tarjeta de red RDMA CX7 InfiniBand de 400 Gb/s (~50 GB/s de ancho de banda máximo).
- Sigue la configuración típica de producción de DeepSeek-V3/R1 (128 tokens por batch, 7168 hidden, top 8 expertos, dispatch en FP8 y combine en BF16).
Inicio rápido
Requisitos
- GPU Hopper (más adelante podría admitir más arquitecturas o dispositivos)
- Python 3.8 o superior
- CUDA 12.3 o superior
- PyTorch 2.1 o superior
- NVLink para comunicación dentro del nodo
- Red RDMA para comunicación entre nodos
Descarga e instalación de la dependencia NVSHMEM
DeepEP depende de una versión modificada de NVSHMEM. Debe instalarse consultando la guía de instalación.
Configuración de red
DeepEP ha sido probado por completo en redes InfiniBand y, en teoría, también es compatible con RDMA over Converged Ethernet (RoCE).
Aislamiento de tráfico
InfiniBand admite aislamiento de tráfico mediante Virtual Lanes (VL). Para evitar interferencias entre distintos tipos de tráfico, se recomienda separar las cargas de trabajo en carriles virtuales de la siguiente manera:
- Cargas de trabajo que usan kernels generales
- Cargas de trabajo que usan kernels de baja latencia
- Otras cargas de trabajo
En DeepEP, se puede controlar la asignación de carriles virtuales configurando la variable de entorno NVSHMEM_IB_SL.
Enrutamiento adaptativo
El enrutamiento adaptativo es una función avanzada de enrutamiento que ofrecen los switches InfiniBand y permite distribuir el tráfico de forma uniforme entre múltiples rutas. Actualmente, los kernels de baja latencia admiten enrutamiento adaptativo, pero los kernels generales no (aunque podrían admitirlo pronto). Si se habilita el enrutamiento adaptativo para kernels generales entre nodos, pueden producirse bloqueos mutuos o corrupción de datos. En el caso de los kernels de baja latencia, habilitar el enrutamiento adaptativo puede eliminar por completo la congestión de red causada por conflictos de enrutamiento, pero agrega latencia adicional. Para un rendimiento óptimo, se recomienda la siguiente configuración:
- Habilitar enrutamiento adaptativo en entornos con alta carga de red
- Usar enrutamiento estático en entornos con baja carga de red
Control de congestión
Dado que no se observó congestión significativa en entornos de producción, el control de congestión está desactivado.
Interfaces y ejemplos
Uso del ejemplo en entrenamiento del modelo o prefill de inferencia
Los kernels generales pueden usarse en la etapa de entrenamiento del modelo o en la etapa de prefill de inferencia (sin la parte backward).
Uso del ejemplo en decoding de inferencia
Los kernels de baja latencia pueden usarse en la etapa de decoding de inferencia.
Precauciones
- Para lograr rendimiento extremo, se descubrió y utilizó la instrucción PTX no documentada
ld.global.nc.L1::no_allocate.L2::256B. Esta instrucción provoca comportamiento indefinido al acceder a memoria volátil de GPU usando un modificador PTX de lectura no coherente y de solo lectura. Sin embargo, al probarla con.L1::no_allocateen la arquitectura Hopper, el rendimiento mejoró considerablemente. Si el kernel no funciona en otras plataformas, puede desactivarse agregandoDISABLE_AGGRESSIVE_PTX_INSTRS=1ensetup.py, o bien se puede reportar un issue. - Para obtener mejor rendimiento en el clúster, se recomienda ejecutar todas las pruebas y usar la configuración óptima de autotuning. La configuración predeterminada está optimizada para el clúster interno de DeepSeek.
Licencia
Este repositorio de código se publica bajo la licencia MIT, y el código que hace referencia a NVSHMEM (incluyendo csrc/kernels/ibgda_device.cuh y third-party/nvshmem.patch) está sujeto a la NVSHMEM SLA.
1 comentarios
Opiniones de Hacker News
.L1::no_allocateen la arquitectura Hopper debería estar garantizada y el rendimiento será mucho mejor