▲

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

 (github.com/deepseek-ai)
1P by neo 3일전 | ★ favorite | 댓글 1개

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의 적용을 받음.

▲
neo 3일전  [-]
Hacker News 의견
  • 극한의 성능을 위해 문서에 없는 PTX 명령어를 발견하고 사용함. 이 명령어는 비일관적인 읽기 전용 PTX 수정자를 사용하여 휘발성 GPU 메모리에 접근하는 것으로 정의되지 않은 동작을 초래할 수 있음. 그러나 Hopper 아키텍처에서 .L1::no_allocate로 테스트된 정확성이 보장되며 성능이 훨씬 나아질 것임
  • Zuckerberg는 Meta가 AI를 오픈 소싱한다고 주장하는 것을 멈춰야 함. 그들은 코드가 아닌 가중치만 공개하고 있음. 진정한 오픈 소스 AI는 DeepSeek뿐임
  • 마치 사탕 가게에 있는 아이처럼 느껴짐. 이러한 트릭 중 일부는 논문을 기반으로 올바르게 역공학하는 데 너무 오래 걸릴 것임. 이번 주의 발표가 MoE를 기본 학술 모델로 사용하는 르네상스를 시작하길 바람
  • 이 사람들을 사랑하지 않을 수 없음. 그들은 우리 모두를 위해 오픈 소스의 경계를 정말로 밀어붙이고 있음. 공유해줘서 고마움
    • 효율적이고 최적화된 all-to-all 통신
    • NVLink와 RDMA를 통한 노드 내 및 노드 간 지원
    • 훈련 및 추론 사전 채우기를 위한 고처리량 커널
    • 추론 디코딩을 위한 저지연 커널
    • 네이티브 FP8 디스패치 지원
    • 계산-통신 중첩을 위한 유연한 GPU 자원 제어
  • DeepSeek의 작업 뒤에 있는 동기는 잘못된 것일 수 있음 (예: AI에서 미국의 선도적 이점을 없애려는 국가 후원 시도). 그러나 전 세계적으로 그 결과는 단순히 환상적임
    • 최악의 경우 (잘못된 이유로 이 작업을 수행하는 경우)에도 DeepSeek에게 감사함. 그들은 OpenAI가 수년간 전 세계에 거짓말한 것을 실제로 하고 있음
    • 정말 대단함
  • 모두가 기대했던 PTX가 이번에 포함되었는지 궁금함
  • 기술 보고서에서 언급한 PTX 명령어가 여기 코드로 연결되어야 함
  • 미국이 싱가포르에서 GPU 영수증을 추적하는 동안 DeepSeek이 H800만 사용했는지 확인하기 위해, 나머지 세계는 전체 H100에서 이러한 최적화를 실행할 수 있음
    • 미국의 제재와 그들의 명령이 전 세계를 덮고 있다고 믿는 오만함 때문에 H100을 얻거나 접근하기 어려웠다고 가장하는 것인지 궁금함
  • 실제 "Open AI™" 회사의 두 번째 오픈 소스 릴리스이며 MIT 라이선스 하에 있음
    • DeepSeek은 $157B+를 주장하는 회사보다 더 개방적임
    • 거의 아무도 Meta의 Llama에 대해 이야기하지 않으며 모두가 Llama 4를 이유와 함께 출시할 것으로 예상해야 함
    • 목표는 제로로의 경주에서 중간에 끼이지 않는 것임
답변달기