DeepSeek, MoE 훈련 및 추론을 위한 오픈 소스 DeepEP 라이브러리 공개

DeepEP

DeepEP는 Mixture-of-Experts (MoE)와 expert parallelism (EP)을 위한 통신 라이브러리임. 고속의 저지연 all-to-all GPU 커널을 제공하며, MoE dispatch와 combine으로 알려져 있음. 또한 FP8을 포함한 저정밀도 연산을 지원함. DeepSeek-V3 논문에서 제안된 그룹 제한 게이팅 알고리듬에 맞춰, 비대칭 도메인 대역폭 전달을 최적화한 커널을 제공하며, NVLink 도메인에서 RDMA 도메인으로 데이터를 전달함. 이러한 커널은 높은 처리량을 제공하여 학습과 추론 사전 채우기 작업에 적합함. 또한 SM(Streaming Multiprocessors) 수 제어를 지원함. 지연에 민감한 추론 디코딩을 위해, DeepEP는 순수 RDMA를 사용하여 지연을 최소화하는 저지연 커널을 포함함. 이 라이브러리는 SM 리소스를 점유하지 않는 hook 기반의 통신-계산 중첩 방법을 도입함.

성능

NVLink 및 RDMA 전달을 사용하는 일반 커널

• H800에서 NVLink 최대 대역폭 약 160 GB/s, CX7 InfiniBand 400 Gb/s RDMA 네트워크 카드(~50 GB/s 최대 대역폭)와 연결하여 일반 커널을 테스트함.
• DeepSeek-V3/R1 사전 학습 설정(배치당 4096 토큰, 7168 히든, 상위 4개 그룹, 상위 8개 전문가, FP8 디스패칭 및 BF16 결합)을 따름.

순수 RDMA를 사용하는 저지연 커널

• H800에서 CX7 InfiniBand 400 Gb/s RDMA 네트워크 카드(~50 GB/s 최대 대역폭)와 연결하여 저지연 커널을 테스트함.
• 일반적인 DeepSeek-V3/R1 프로덕션 설정(배치당 128 토큰, 7168 히든, 상위 8개 전문가, FP8 디스패칭 및 BF16 결합)을 따름.

빠른 시작

요구 사항

• Hopper GPU (추후 더 많은 아키텍처나 장치를 지원할 수 있음)
• Python 3.8 이상
• CUDA 12.3 이상
• PyTorch 2.1 이상
• 노드 내 통신을 위한 NVLink
• 노드 간 통신을 위한 RDMA 네트워크

NVSHMEM 종속성 다운로드 및 설치

DeepEP는 수정된 NVSHMEM에 의존함. 설치 가이드를 참조하여 설치해야 함.

네트워크 구성

DeepEP는 InfiniBand 네트워크에서 완전히 테스트되었으며, 이론적으로 RDMA over Converged Ethernet (RoCE)와도 호환 가능함.

트래픽 격리

InfiniBand는 Virtual Lanes (VL)을 통해 트래픽 격리를 지원함. 서로 다른 유형의 트래픽 간 간섭을 방지하기 위해, 다음과 같이 가상 레인에 작업을 분리할 것을 권장함:

• 일반 커널을 사용하는 작업
• 저지연 커널을 사용하는 작업
• 기타 작업

DeepEP에서는 NVSHMEM_IB_SL 환경 변수를 설정하여 가상 레인 할당을 제어할 수 있음.

적응형 라우팅

적응형 라우팅은 InfiniBand 스위치에서 제공하는 고급 라우팅 기능으로, 여러 경로에 트래픽을 고르게 분산할 수 있음. 현재 저지연 커널은 적응형 라우팅을 지원하지만, 일반 커널은 지원하지 않음(곧 지원될 수 있음). 일반 노드 간 커널에 적응형 라우팅을 활성화하면 교착 상태나 데이터 손상 문제가 발생할 수 있음. 저지연 커널의 경우, 적응형 라우팅을 활성화하면 라우팅 충돌로 인한 네트워크 혼잡을 완전히 제거할 수 있지만, 추가적인 지연이 발생함. 최적의 성능을 위해 다음 구성을 권장함:

• 네트워크 부하가 큰 환경에서는 적응형 라우팅을 활성화
• 네트워크 부하가 적은 환경에서는 정적 라우팅 사용

혼잡 제어

생산 환경에서 유의미한 혼잡을 관찰하지 못했기 때문에 혼잡 제어는 비활성화됨.

인터페이스 및 예제

모델 학습 또는 추론 사전 채우기에서의 예제 사용

일반 커널은 모델 학습 또는 추론 사전 채우기 단계(역방향 부분 없이)에서 사용할 수 있음.

추론 디코딩에서의 예제 사용

저지연 커널은 추론 디코딩 단계에서 사용할 수 있음.

주의 사항

• 극한의 성능을 위해, 문서에 없는 PTX 명령어 ld.global.nc.L1::no_allocate.L2::256B를 발견하고 사용함. 이 명령어는 비일관성 읽기 전용 PTX 수정자를 사용하여 휘발성 GPU 메모리에 접근하는 정의되지 않은 동작을 유발함. 그러나 Hopper 아키텍처에서 .L1::no_allocate로 테스트한 결과, 성능이 훨씬 향상됨. 다른 플랫폼에서 커널이 작동하지 않으면 DISABLE_AGGRESSIVE_PTX_INSTRS=1을 setup.py에 추가하여 이를 비활성화하거나 이슈를 제기할 수 있음.
• 클러스터에서 더 나은 성능을 위해 모든 테스트를 실행하고 최적의 자동 튜닝 구성을 사용하는 것을 권장함. 기본 구성은 DeepSeek의 내부 클러스터에서 최적화됨.

라이선스

이 코드 저장소는 MIT 라이선스 하에 공개되며, NVSHMEM을 참조하는 코드(csrc/kernels/ibgda_device.cuh 및 third-party/nvshmem.patch 포함)는 NVSHMEM SLA의 적용을 받음.