GitHub - deepseek-ai/DeepEP: DeepEP: an efficient expert-parallel communication library

DeepEP는 Mixture-of-Experts (MoE) 및 expert parallelism을 위해 특별히 설계된 고성능 통신 라이브러리입니다. 이 라이브러리는 MoE dispatch 및 combine이라고도 알려진 All-to-All GPU 커널에서 높은 처리량(high-throughput)과 낮은 지연 시간(low-latency)을 제공하며, FP8과 같은 저정밀도(low-precision) 연산을 지원합니다.

핵심 방법론 및 기능:

MoE 통신 최적화:

* Normal Kernels (학습 및 추론 Pre-filling): 높은 처리량에 중점을 둡니다.
* Asymmetric-domain bandwidth forwarding: DeepSeek-V3 논문에서 제안된 group-limited gating 알고리즘에 맞춰, NVLink 도메인에서 RDMA 도메인으로 데이터를 전달하는 것과 같은 비대칭 도메인 대역폭 포워딩을 위한 커널을 제공합니다. 이는 학습 및 추론 pre-filling 작업에 적합합니다.
* SM (Streaming Multiprocessors) number control: GPU의 SM 사용량을 제어할 수 있습니다.
* deep_ep.Buffer 클래스를 통해 통신 버퍼를 관리하며, get_dispatch_layout, dispatch, combine 함수를 사용합니다. dispatch 커널은 현재 랭크가 수신할 토큰 수를 알기 위해 GPU의 신호를 기다리는 CPU 대기(CPU wait for GPU signal)를 수반하며, 이는 CUDA Graph와 호환되지 않을 수 있습니다.
* 성능: H800 GPU에서 Intranode의 경우 NVLink를 통해 Dispatch 153 GB/s, Combine 158 GB/s를, Internode의 경우 RDMA를 통해 #EP 64에서 Dispatch 51 GB/s, Combine 50 GB/s를 달성합니다.
* Low-Latency Kernels (추론 Decoding): 낮은 지연 시간에 중점을 둡니다.
* Pure RDMA: 지연 시간을 최소화하기 위해 순수 RDMA를 사용합니다.
* Hook-based communication-computation overlapping: 통신-연산 중첩(communication-computation overlapping)을 위한 독점적인 후크(hook) 기반 방식을 도입하여 GPU SM 자원을 전혀 점유하지 않습니다. low_latency_dispatch 및 low_latency_combine 함수는 $return_recv_hook=True$ 옵션을 통해 후크를 반환하며, 실제 데이터 수신은 이 후크를 호출할 때 이루어져, 특히 double-batch overlapping 시나리오에서 효율적인 자원 활용을 가능하게 합니다.
* CUDA Graph 호환성: CUDA Graph와 호환됩니다.
* 성능: H800 GPU에서 #EP 8일 때 Dispatch 77 us (98 GB/s), Combine 114 us (127 GB/s)를, #EP 256일 때 Dispatch 194 us (39 GB/s), Combine 360 us (40 GB/s)를 달성합니다.

하드웨어 및 네트워크 최적화:

* GPU 호환성: Ampere (SM80), Hopper (SM90) GPU 또는 SM90 PTX ISA를 지원하는 기타 아키텍처를 지원합니다.
* 통신 기술: Intranode 통신을 위해 NVLink를, Internode 통신을 위해 RDMA 네트워크를 사용합니다. InfiniBand 네트워크에서 완전히 테스트되었으며, 이론적으로 RoCE(RDMA over Converged Ethernet)와도 호환됩니다.
* NVSHMEM 의존성: DeepEP는 NVSHMEM에 의존하며, NVSHMEM_DIR 환경 변수를 통해 경로를 지정해야 합니다.
* 네트워크 설정:
* Traffic isolation: InfiniBand의 Virtual Lanes (VL)를 통해 트래픽 격리를 지원하여 서로 다른 유형의 워크로드(normal kernels, low-latency kernels, other workloads) 간의 간섭을 방지합니다. NVSHMEM_IB_SL 환경 변수로 VL 할당을 제어합니다.
* Adaptive routing: InfiniBand 스위치의 고급 라우팅 기능으로, 트래픽을 여러 경로에 고르게 분산하여 라우팅 충돌로 인한 네트워크 혼잡을 제거할 수 있습니다. 무거운 네트워크 부하 환경에서는 활성화하고, 가벼운 부하 환경에서는 정적 라우팅을 권장합니다.
* Congestion control: 현재 프로덕션 환경에서 심각한 혼잡이 관찰되지 않아 비활성화되어 있습니다.

성능 최적화 기법:

* Undefined-behavior PTX usage: 극대화된 성능을 위해 ld.global.nc.L1::no_allocate.L2::256B와 같은 정의되지 않은 PTX 명령어를 활용합니다. 이는 Hopper 아키텍처에서 L1 캐시와 비일관성 캐시(non-coherent cache)가 통합되어 발생하는 특정 동작을 활용하는 것으로, 성능 향상에 기여합니다. 호환성 문제가 발생할 경우 $DISABLE_AGGRESSIVE_PTX_INSTRS=1$을 설정하여 비활성화할 수 있습니다.
* Auto-tuning: 최적의 성능을 위해 사용자 클러스터 환경에서 테스트를 실행하여 자동 튜닝된 설정을 사용하는 것을 권장합니다. 기본 설정은 DeepSeek의 내부 클러스터에 최적화되어 있습니다.

개발 및 향후 계획 (Roadmap):
AR(Adaptive Routing) 지원, low-latency 모드 코드 리팩토링, A100 intranode 지원, low-latency dispatch 커널의 BF16 지원, NVLink 프로토콜을 이용한 intranode low-latency 커널 지원, TMA(Tensor Memory Accelerator) 복사, SM-free 커널 및 리팩토링, PCIe 커널 지원, NVFP4 데이터 타입 지원 등이 계획되어 있습니다.

주요 참고 사항:
* 현재 DeepEP 구현은 통신 버퍼로 큐(queues)를 사용하는데, 이는 메모리를 절약하지만 복잡성과 잠재적 데드락을 유발할 수 있습니다. 단순성과 성능 향상을 위해 고정 크기(fixed-size) 버퍼를 사용하는 대안적 접근 방식이 제안됩니다.
* NVSHMEM을 참조하는 코드(예: csrc/kernels/ibgda_device.cuh 및 third-party/nvshmem.patch)는 NVSHMEM SLA에 따르며, 그 외 코드는 MIT License로 배포됩니다.