- Mixture-of-Experts(MoE) 및 Expert Parallelism(EP)을 위한 고성능 통신 라이브러리
- GPU 기반 All-to-All 커널을 제공하여 MoE 디스패치 및 결합 연산을 고속으로 처리
-
FP8과 같은 저정밀 연산 지원
-
DeepSeek-V3 논문에서 제안한 그룹 제한 게이팅(group-limited gating) 알고리즘을 적용하여 비대칭 도메인 대역폭 포워딩을 최적화
- 예: NVLink → RDMA 데이터 전송 최적화
- 훈련 및 추론 프리필링(prefilling) 작업에 적합한 높은 처리량 제공
- 지연시간에 민감한 추론 디코딩을 위해 RDMA 전용 저지연 커널 포함
-
통신-연산 오버랩 기법 제공 (SM 리소스를 점유하지 않음)
성능
일반 커널 (NVLink 및 RDMA 전송)
- DeepEP는 H800 GPU와 CX7 InfiniBand 400Gb/s RDMA 네트워크 환경에서 성능을 테스트함
- DeepSeek-V3/R1 설정을 기반으로, 배치당 4096개 토큰, 7168개 히든 노드, top-4 그룹, top-8 전문가 구조를 적용하고, FP8 디스패칭 및 BF16 결합을 사용함
- 성능 테스트 결과, 노드 내부(NVLink 기반) 통신은 약 150GB/s 이상의 대역폭을 보였으며, 노드 간(RDMA 기반) 통신은 전문가 수에 따라 40~47GB/s 수준의 대역폭을 기록함
- 전문가 수가 많아질수록 RDMA 대역폭이 소폭 증가하는 경향을 보였음 (예: 16명 전문가일 때 43GB/s, 64명 전문가일 때 46GB/s)
저지연 커널 (순수 RDMA)
- 저지연 커널의 성능을 측정한 결과, 일반 커널보다 지연시간(latency)이 크게 단축됨
- 배치당 128개 토큰을 처리하는 환경에서, 전문가 수에 따라 지연시간이 증가했지만 RDMA 대역폭은 비교적 일정하게 유지됨
- 예를 들어, 전문가 8명일 때 163마이크로초(us), 전문가 256명일 때 194마이크로초(us) 수준으로 증가하였음
- 결합(combine) 연산에서는 디스패치보다 더 높은 지연시간이 발생했으며, 전문가 수가 증가할수록 RDMA 대역폭이 40GB/s 이하로 점진적으로 감소하는 경향을 보임
- 즉, 저지연 커널은 작은 규모의 전문가 그룹에서는 매우 빠르게 동작하지만, 전문가 수가 많아지면 지연시간이 늘어나므로 적절한 균형이 필요함
네트워크 설정
트래픽 격리 (Traffic Isolation)
- InfiniBand의 Virtual Lanes(VL) 을 활용하여 트래픽을 격리 가능
- 권장 분리 방식:
- 일반 커널 사용 작업
- 저지연 커널 사용 작업
- 기타 작업
-
NVSHMEM_IB_SL 환경 변수를 통해 VL 설정 가능
적응형 라우팅 (Adaptive Routing)
- InfiniBand 스위치의 적응형 라우팅을 지원
-
저지연 커널에서는 활성화 가능, 일반 커널에서는 비활성화 필요 (활성화 시 데이터 손상 위험)
- 설정 권장 사항:
-
네트워크 부하가 높은 경우: 적응형 라우팅 활성화
-
네트워크 부하가 낮은 경우: 정적 라우팅 유지
혼잡 제어 (Congestion Control)
- DeepEP는 혼잡 제어 기능을 비활성화한 상태로 운영
- 실제 환경에서 네트워크 혼잡이 심각하지 않음을 확인함
주요 기술적 고려 사항
-
비공식 PTX 명령어 사용:
ld.global.nc.L1::no_allocate.L2::256B를 활용하여 성능 최적화
- Hopper 아키텍처에서는 정상 작동하지만, 다른 플랫폼에서는
DISABLE_AGGRESSIVE_PTX_INSTRS=1을 설정하여 비활성화 가능
-
자동 튜닝 권장: 최적의 성능을 위해 클러스터별 성능 테스트 후 설정 적용 필요