DeepSeek publica DeepEP, una biblioteca open source para entrenamiento e inferencia MoE

 (github.com/deepseek-ai)
1 puntos por GN⁺ 2025-02-26 | 1 comentarios | Compartir por WhatsApp
  • DeepEP es una biblioteca de comunicación de alto rendimiento enfocada en la paralelización de expertos (EP) para el entrenamiento y la inferencia modernos de ML, que ofrece kernels GPU all-to-all para dispatch/combine de MoE y soporte de baja precisión, incluido FP8
  • La versión V2 refactoriza por completo EP para lograr un rendimiento igual o mejor que V1 con muchos menos recursos SM, y cambia el backend de NVSHMEM a NCCL Gin, más liviano
  • En pruebas con la configuración V3, se midió con 8K tokens por batch, hidden 7168, top 8 experts, dispatch FP8 y combine BF16; V2 registró hasta 1.3 veces más rendimiento pico y hasta 4 veces menos SM frente a V1
  • Todos los kernels se compilan en runtime como módulos JIT livianos, sin requerir compilación CUDA durante la instalación, y V2 unifica las API de alto throughput y baja latencia en una sola interfaz ElasticBuffer
  • Requiere GPU Hopper SM90, Python 3.8+, CUDA 12.3+, PyTorch 2.10+, NCCL 2.30.4+, NVLink y una red RDMA entre nodos; Engram, PP y CP son funciones experimentales

Alcance de DeepEP

  • DeepEP (DeepEveryParallel) es una biblioteca de comunicación de alto rendimiento para entrenamiento e inferencia modernos de machine learning
  • Su función central actual es la paralelización de expertos (Expert Parallelism, EP), y ofrece kernels GPU all-to-all de alto throughput y baja latencia para dispatch y combine de MoE
  • Soporta comunicación de baja precisión, incluido FP8
  • También incluye primitivas experimentales para paralelización de pipeline (PP), paralelización de contexto (CP) y acceso remoto a memoria (Engram)
  • Todos los kernels se compilan en runtime como módulos JIT (Just-In-Time) livianos, sin requerir compilación CUDA durante la instalación
  • A pesar de su diseño liviano, apunta a alcanzar o superar el límite de ancho de banda del hardware en varias configuraciones

Cambios clave de la versión V2

  • V2 es una versión que refactoriza por completo Expert Parallelism
    • Está diseñada para lograr rendimiento extremo usando varias veces menos recursos SM que V1
    • Soporta dominios de scale-up y scale-out más grandes
    • El backend cambia de NVSHMEM a un backend NCCL Gin más liviano
  • Las nuevas funciones son las siguientes
    • Compilación JIT completa
    • Backend NCCL Gin liviano y header-only
    • Posibilidad de reutilizar communicators NCCL existentes
    • En EPv2, unifica la API de alto throughput y la API de baja latencia en una sola interfaz ElasticBuffer
    • Proporciona un nuevo layout GEMM
    • Soporta dominios de scale-up y scale-out más grandes, hasta EP2048
    • Calcula analíticamente la cantidad de SM y QP, por lo que ya no hace falta autotuning
    • Sigue soportando tanto el modo híbrido como el modo directo
    • En entrenamiento legacy similar a V3, mantiene un rendimiento equivalente o mejor reduciendo el uso de SM de 24 a 4~6
    • Engram de 0 SM basado en RDMA
    • PP de 0 SM basado en RDMA
    • CP de 0 SM basado en Copy Engine

Limitaciones y funciones en desarrollo

  • V2 tiene un consumo de tamaño de buffer mayor que V1
  • EP de baja latencia RDMA de 0 SM ya no está soportado
  • Engram, PP y CP son funciones experimentales
  • Las funciones en desarrollo son las siguientes
    • Elastic GPU & CPU buffers, un espacio de direcciones virtuales continuo que mapea de forma mixta la memoria física de GPU y CPU
      • Apunta a habilitar Engram o EP desbalanceado de forma totalmente automática y transparente
    • Trabajo para manejar desbalance de carga mediante EP replay y reducir el tamaño de buffers intermedios
    • Actualizaciones de all-gather e implementación de reduce-scatter para DP y TP
  • La documentación de V1 basada en NVSHMEM está disponible en docs/legacy.md

Resultados de medición de rendimiento

  • Se probó con las siguientes condiciones, alineadas con la configuración V3
    • 8K tokens por batch
    • hidden dimension 7168
    • top 8 experts
    • dispatch FP8
    • combine BF16
  • Los resultados principales son los siguientes
    • 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
  • Los valores medidos son ancho de banda lógico; por ejemplo, los 90 GB/s de EP 8 x 2 incluyen el tráfico de local rank
  • V2 logra hasta 1.3 veces el rendimiento pico frente a V1 y reduce hasta 4 veces la cantidad de SM
  • Los resultados de configuraciones EP más grandes se omiten por ahora, y se recomienda que los usuarios hagan sus propios benchmarks
  • Según la experiencia interna, se espera que los kernels sigan saturando el ancho de banda del hardware incluso a mayor escala
  • Los datos de rendimiento de V1 están en docs/legacy.md

Instalación y requisitos

  • Los requisitos son los siguientes
    • GPU Hopper SM90 o arquitectura compatible con SM90 PTX ISA
    • Python 3.8 o superior
    • CUDA 12.3 o superior para GPU SM90
    • PyTorch 2.10 o superior
    • NCCL 2.30.4 o superior
    • NVLink para comunicación dentro del nodo
    • Red RDMA para comunicación entre nodos
  • Se recomienda instalar NCCL con pip para que DeepEP pueda detectarlo automáticamente en el entorno Python
pip install "nvidia-nccl-cu13>=2.30.4" --no-deps
  • Para soportar métodos legacy, también depende de NVSHMEM; consultar las instrucciones de instalación en NVSHMEM Installation Guide
  • Ejemplos para ejecutar build y pruebas durante el desarrollo
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
  • La instalación se realiza con el siguiente comando
python setup.py install
  • Después de instalar, se puede usar importando deep_ep en un proyecto Python

Interfaz centrada en ElasticBuffer

  • En V2, todas las operaciones EP se integran bajo una única interfaz ElasticBuffer
    • Maneja las API de alto throughput y baja latencia con la misma interfaz
    • El buffer puede inicializarse especificando directamente la configuración MoE
    • Calcula analíticamente la cantidad óptima de SM y QP
  • Un ejemplo de inicialización del buffer calcula el tamaño necesario con ElasticBuffer.get_buffer_size_hint() y verifica si puede reutilizarse un buffer existente
  • Al crear un buffer nuevo, se especifican parámetros como num_max_tokens_per_rank, hidden, num_topk y use_fp8_dispatch
  • Con _buffer.get_theoretical_num_sms(num_experts, num_topk) se obtiene la cantidad teórica de SM que usará el kernel de comunicación
  • Si se especifica directamente num_sms en las llamadas a dispatch y combine, se puede sobrescribir el valor calculado

Patrones de uso en entrenamiento, prefill y decoding

  • En entrenamiento o prefill de inferencia, MoE dispatch enruta tokens hacia el expert correspondiente de todos los ranks
    • Soporta entradas BF16 y FP8
    • handle contiene los metadatos de routing necesarios para la llamada posterior a combine
    • handle.num_recv_tokens_per_expert_list proporciona la cantidad de tokens por expert necesaria para GEMM
  • El backward pass de MoE dispatch en realidad se procesa con combine
  • MoE combine reduce las salidas de los experts hacia el rank original
  • El backward pass de MoE combine en realidad se procesa con dispatch
  • La superposición de comunicación y cómputo se gestiona con la interfaz EventOverlap
    • Permite realizar cómputo independiente mientras la comunicación está en curso
    • Antes de usar el resultado, se sincroniza el compute stream con event.current_stream_wait()
  • En decoding de inferencia también se usa el mismo ElasticBuffer
    • Cuando la gating decision no cambia, se reutilizan los metadatos de routing con cached_handle
    • Este patrón evita recalcular el layout y la sincronización con CPU

Variables de entorno y valores fijos en build

  • Configuración general
    • EP_BUFFER_DEBUG: muestra información de depuración sobre inicialización de buffers, aproximación de SM y backend
    • EP_SUPPRESS_NCCL_CHECK: suprime la verificación de incompatibilidad de versión de NCCL
    • EP_AVOID_RECORD_STREAM: evita record_stream en los tensores de salida
    • EP_NUM_TOPK_IDX_BITS: sobrescribe la cantidad de bits para codificar índices top-k
  • Configuración de red
    • EP_NIC_NAME: nombre de la NIC predeterminada que se usará para consultar atributos de NIC; valor por defecto mlx5_0
    • EP_OVERRIDE_RDMA_SL: sobrescribe el índice de service level RDMA
    • EP_DISABLE_GIN: desactiva el backend NCCL Gin
  • Configuración JIT
    • EP_JIT_CACHE_DIR: directorio de caché para kernels compilados; valor por defecto $HOME/.deep_ep
    • EP_JIT_NVCC_COMPILER: ruta del compilador NVCC
    • EP_JIT_CPP_STANDARD: versión del estándar C++; valor por defecto 20
    • EP_JIT_DUMP_PTX, EP_JIT_DUMP_SASS, EP_JIT_DUMP_ASM: configuración relacionada con dumps de salida PTX/SASS/ASM
  • Algunas variables de entorno funcionan como persistentes
    • Se capturan durante el build y se incluyen como valores por defecto en el paquete instalado
    • Si no se sobrescriben con las variables de entorno actuales al momento de importar, estos valores por defecto se aplican automáticamente
    • Las variables objetivo son EP_JIT_CACHE_DIR, EP_JIT_PRINT_COMPILER_COMMAND, EP_NUM_TOPK_IDX_BITS y EP_NCCL_ROOT_DIR
  • Para más detalles, consultar test_ep.py o la documentación de Python

Recomendaciones de configuración de red

  • DeepEP fue probado completamente en redes InfiniBand
  • En teoría, también es compatible con RDMA over Converged Ethernet, es decir, RoCE
  • Aislamiento de tráfico
    • Está soportado mediante Virtual Lanes de InfiniBand
    • Se recomienda separar las cargas expert-parallel y otras cargas en virtual lanes distintas
    • En V2, la asignación de virtual lanes puede controlarse con el argumento sl_idx o con la variable de entorno EP_OVERRIDE_RDMA_SL
  • Adaptive routing
    • Es una función avanzada de routing en la que los switches InfiniBand distribuyen el tráfico de forma uniforme entre múltiples rutas
    • Se recomienda activarla en todas las condiciones de carga de red, aunque agregue latencia adicional
  • Congestion control
    • Se desactiva porque perjudica el ancho de banda máximo
    • Si la congestión es inevitable, se recomienda asignar esa workload a una virtual lane de menor prioridad
  • PCI atomic mode
    • Si el hardware lo soporta, se recomienda configurar PCI_ATOMIC_MODE en la NIC para mejorar el rendimiento de las RDMA atomic operations
sudo mlxconfig -y -d mlx5_$i set PCI_ATOMIC_MODE=4

Ramas experimentales y forks de la comunidad

  • Ramas experimentales
    • Zero-copy: elimina copias entre tensores de PyTorch y buffers de comunicación, reduciendo mucho el uso de SM de los kernels generales
    • Eager: usa un protocolo de baja latencia para eliminar la latencia extra de RTT que agregan las RDMA atomic OP
    • Hybrid-EP: implementación de un nuevo backend con instrucciones TMA, uso mínimo de SM, soporte para dominios NVLink más grandes, superposición de comunicación y cómputo fine-grained en single-batch, kernels PCIe y soporte NVFP4
    • AntGroup-Opt: serie de optimizaciones escrita por AntGroup Network Platform Department
    • Mori-EP: soporte de modo de baja latencia para ROCm/GPU AMD basado en el backend MORI
    • nvDev: rama basada en V2 que incluye funciones recientes de CUDA, como Compute Fabric Transport
  • Forks de la comunidad
    • uccl/uccl-ep: soporte para ejecutar DeepEP en GPU heterogéneas como Nvidia y AMD, y en NIC como EFA, Broadcom y CX7
    • Infrawaves/DeepEP_ibrc_dual-ports_multiQP: agrega una solución multi-QP y soporte para NIC dual-port al transporte IBRC
    • antgroup/DeepXTrace: analizador de diagnóstico para encontrar slow ranks de forma eficiente y precisa
    • ROCm/mori: biblioteca de comunicación AMD de próxima generación para workloads de IA sensibles al rendimiento, como Wide EP, transferencia de KVCache y Collectives

Licencia y cita

  • DeepEP V2 está construido sobre el backend Gin de NCCL
  • El código del repositorio se publica bajo la MIT License
  • El ítem de cita es DeepEP: an efficient expert-parallel communication library, y el año se indica como 2025

Lecturas relacionadas

1 comentarios

 
GN⁺ 2025-02-26
Opiniones en Hacker News

  • Consiguen un rendimiento extremo al encontrar y usar la instrucción PTX no documentada ld.global.nc.L1::no_allocate.L2::256B
    Esta instrucción puede provocar comportamiento indefinido porque accede a memoria GPU volátil con el modificador PTX de solo lectura no coherente .nc
    Sin embargo, dicen que en la arquitectura Hopper se probó que garantiza la corrección cuando se usa junto con .L1::no_allocate, y que el rendimiento mejora mucho

    • En la práctica, ¿podría NVIDIA más adelante, en una nueva arquitectura, cambiar sutilmente el comportamiento de esta instrucción fuera de la documentación y dar vuelta la situación, ya sea intencionalmente o no?

  • Se siente como un niño entrando a una dulcería
    Hay muchos trucos que tomarían demasiado tiempo de ingeniería inversa solo a partir de los papers, y ojalá las publicaciones de esta semana inauguren un renacimiento en el que MoE se use como modelo académico base

    • Visto desde esta perspectiva, no entiendo qué está pasando entre la práctica real de modelos de vanguardia y los modelos académicos
      Los primeros ya son todos MoE desde GPT-4, pero los modelos abiertos, salvo DeepSeek V3 y Mixtral, suelen ser modelos densos

  • Es imposible no querer a este equipo
    Están empujando los límites del open source para todos

    • Como Open AI™ escrito separado
    • En realidad, no es open source
      Si quieres ver un modelo verdaderamente open source, revisa OLMo 2 de AI2: https://allenai.org/blog/olmo2
      Comparten de verdad todo lo necesario para reproducir el modelo, incluidos los datos mismos
      En el enlace de arriba también dicen: “Como la ciencia abierta completa requiere más que pesos públicos, nos complace compartir con la comunidad más amplia de modelado de lenguaje la nueva actualización de OLMo, que incluye pesos, datos, código, recetas, checkpoints intermedios y modelos ajustados con instrucciones”

  • Zuckerberg debería dejar de afirmar que Meta publica AI como open source
    Hasta hacen anuncios de TV, pero en realidad solo publican los pesos y no el código
    La única AI verdaderamente open source es DeepSeek

    • Estrictamente hablando, DeepSeek tampoco es tan open source como OLMo u Open Euro
      Porque no publicaron los datos
    • DeepSeek claramente no es verdadero open source
      Para ser open source tendría que usar una licencia realmente open source, como las enumeradas por la OSI, y compartir el código de preentrenamiento y postentrenamiento, el código relacionado con tuning, el código de evaluación, todo lo relacionado con seguridad y censura, y probablemente todo el conjunto de datos de entrenamiento
      De lo contrario no se pueden reproducir los pesos, y compartir pesos es parecido a compartir un programa compilado
      Hasta donde sé, el único modelo competitivo realmente open source es OLMo 2 de AI2: https://allenai.org/blog/olmo2
      Recientemente también publicaron una app que hace inferencia en el dispositivo, y también es open source: https://allenai.org/blog/olmoe-app
      También tienen otro modelo llamado Tülu 3, que dicen que rinde mejor que DeepSeek V3: https://allenai.org/blog/tulu-3-405B
    • Meta lleva más de 10 años puliendo PyTorch
      Casi todo lo necesario para entrenar LLM, incluida la tecnología más reciente, está ahí
      ¿Qué más hace falta? ¿Partes de código específicas de la infraestructura de Meta?
    • ¿PyTorch también cuenta?
    • Pesos públicos = un montón de binarios
      Es como volver al modelo FREEWARE / SHAREWARE
      Deberíamos usar esa expresión para los pesos “públicos”

  • Proporciona comunicación all-to-all eficiente y optimizada, soporte intra-nodo e inter-nodo mediante NVLink y RDMA, kernels de alto throughput para entrenamiento y prefill de inferencia, kernels de baja latencia para decodificación de inferencia, soporte nativo de despacho FP8 y control flexible de recursos GPU para superponer cómputo y comunicación
    X: https://x.com/deepseek_ai/status/1894211757604049133

  • La motivación del trabajo de DeepSeek podría estar equivocada
    Por ejemplo, podría ser un intento patrocinado por el Estado para reducir a cero la ventaja de primacía de Estados Unidos en AI, pero el efecto neto para todos en el mundo simplemente es excelente
    En el peor de los casos, aunque lo hagan por las razones equivocadas, gracias a DeepSeek: están haciendo realmente lo que OpenAI le mintió al mundo durante años diciendo que haría

    • En el ámbito de las relaciones internacionales, lo correcto y lo incorrecto no aplican tanto
      ¿Publicar esto como open source es más “incorrecto” que la prohibición de exportar GPUs Nvidia avanzadas?
      La publicación open source de DeepSeek probablemente solo sea algo que, con el consentimiento del Partido Comunista Chino, termina siendo bueno tanto para el Partido Comunista Chino como para la comunidad open source de AI en general, y no debería tomarse como una postura de principios
      Encontrar maneras de eliminar la ventaja competitiva de otros países es una actividad principal de todos los gobiernos, grandes y pequeños

  • Es la segunda ronda de publicaciones open source bajo licencia MIT por parte de la verdadera empresa Open AI™
    Una vez más, DeepSeek es más abierta que una empresa de 157.000 millones de dólares que dice ser “Open”
    Casi nadie habla de Llama de Meta, pero todos deberían esperar que Llama 4 salga con capacidades de razonamiento
    El objetivo es no morir aplastado en medio de la carrera hacia cero

  • Mientras Estados Unidos revisa los recibos de GPUs de Singapur para comprobar si DeepSeek solo usó H800, ¿el resto del mundo puede correr estas optimizaciones en H100 completas?
    ¿Mientras, por las sanciones de EE. UU., les costaba conseguir o acceder a H100, y mientras EE. UU. seguía fingiendo, por arrogancia, que sus órdenes cubren a todo el mundo?
    ¿Lo estoy entendiendo bien?

  • ¿Esta vez sí incluyeron el PTX que todos esperaban?

    • Sí, hay algo en el directorio csrc/kernels
      Si buscas asm, puedes encontrar dónde se usa
    • Hace falta explicarle al resto por qué ese PTX que todos esperaban es tan importante

  • Me recuerda a los años 80 y 90, cuando la gente hackeaba assembly o buscaba instrucciones no documentadas para exprimir el rendimiento de la CPU
    Algún día los compiladores optimizarán lo suficiente, o las GPUs serán tan potentes, que estos trucos ya no marcarán tanta diferencia, como ocurre hoy con las CPUs