DeepSeek công bố thư viện mã nguồn mở DeepEP cho huấn luyện và suy luận MoE

 (github.com/deepseek-ai)
1 điểm bởi GN⁺ 2025-02-26 | 1 bình luận | Chia sẻ qua WhatsApp
  • DeepEP là thư viện truyền thông hiệu năng cao tập trung vào Expert Parallelism (EP) trong huấn luyện và suy luận ML hiện đại, cung cấp các kernel GPU all-to-all cho MoE dispatch/combine và hỗ trợ độ chính xác thấp như FP8
  • Bản phát hành V2 tái cấu trúc toàn diện EP, đạt hiệu năng tương đương hoặc tốt hơn V1 với ít tài nguyên SM hơn rất nhiều, đồng thời chuyển backend từ NVSHMEM sang NCCL Gin gọn nhẹ hơn
  • Trong các bài kiểm thử theo cấu hình V3 với 8K tokens mỗi batch, hidden 7168, top 8 experts, FP8 dispatch và BF16 combine, V2 ghi nhận hiệu năng đỉnh cao hơn tới 1,3 lần và giảm tới 4 lần số SM so với V1
  • Tất cả kernel đều được biên dịch lúc chạy bằng mô-đun JIT gọn nhẹ, không cần biên dịch CUDA khi cài đặt; V2 cũng hợp nhất API thông lượng cao và độ trễ thấp vào giao diện ElasticBuffer duy nhất
  • Yêu cầu Hopper SM90 GPU, Python 3.8+, CUDA 12.3+, PyTorch 2.10+, NCCL 2.30.4+, NVLink và mạng RDMA liên nút; Engram, PP và CP là tính năng thử nghiệm

Phạm vi mà DeepEP cung cấp

  • DeepEP (DeepEveryParallel) là thư viện truyền thông hiệu năng cao cho huấn luyện và suy luận máy học hiện đại
  • Tính năng trọng tâm hiện tại là Expert Parallelism (EP), cung cấp các kernel GPU all-to-all thông lượng cao, độ trễ thấp cho MoE dispatch và combine
  • Hỗ trợ truyền thông độ chính xác thấp bao gồm FP8
  • Cũng bao gồm các primitive thử nghiệm cho pipeline parallelism (PP), context parallelism (CP) và truy cập bộ nhớ từ xa (Engram)
  • Mọi kernel đều được biên dịch lúc chạy bằng mô-đun JIT (Just-In-Time) gọn nhẹ, không cần biên dịch CUDA trong quá trình cài đặt
  • Dù thiết kế gọn nhẹ, dự án vẫn hướng tới hiệu năng tiệm cận hoặc vượt giới hạn băng thông phần cứng trong nhiều cấu hình

Những thay đổi chính của bản phát hành V2

  • V2 là bản phát hành tái cấu trúc hoàn toàn Expert Parallelism
    • Được thiết kế để đạt hiệu năng cực cao với lượng tài nguyên SM chỉ bằng vài phần so với V1
    • Hỗ trợ miền scale-up và scale-out lớn hơn
    • Backend được chuyển từ NVSHMEM sang NCCL Gin backend gọn nhẹ hơn
  • Các tính năng mới gồm có
    • Biên dịch JIT toàn phần
    • NCCL Gin backend dạng header-only, gọn nhẹ
    • Có thể tái sử dụng NCCL communicator hiện có
    • Trong EPv2, API thông lượng cao và API độ trễ thấp được hợp nhất vào một giao diện ElasticBuffer
    • Cung cấp bố cục GEMM mới
    • Hỗ trợ miền scale-up và scale-out lớn hơn, tới EP2048
    • Tự động tính toán số lượng SM và QP theo phương pháp phân tích, không còn cần auto-tuning
    • Tiếp tục hỗ trợ cả hybrid mode lẫn direct mode
    • Trong huấn luyện legacy tương tự V3, giảm mức dùng SM từ 24 xuống 4~6 mà vẫn giữ hiệu năng tương đương hoặc tốt hơn
    • 0 SM Engram dựa trên RDMA
    • 0 SM PP dựa trên RDMA
    • 0 SM CP dựa trên Copy Engine

Hạn chế và các tính năng đang được phát triển

  • V2 tiêu tốn dung lượng bộ đệm lớn hơn V1
  • EP độ trễ thấp RDMA 0 SM không còn được hỗ trợ
  • Engram, PP, CP là các tính năng thử nghiệm
  • Các tính năng đang được phát triển gồm
    • Elastic GPU & CPU buffers, không gian địa chỉ ảo liên tục ánh xạ trộn bộ nhớ vật lý GPU và CPU
      • Hướng tới việc cho phép Engram hoàn toàn tự động, minh bạch hoặc EP mất cân bằng
    • Xử lý mất cân bằng tải và giảm kích thước bộ đệm trung gian bằng cách tận dụng EP replay
    • Triển khai cập nhật all-gather và reduce-scatter cho DP và TP
  • Tài liệu V1 dựa trên NVSHMEM được cung cấp tại docs/legacy.md

Kết quả đo hiệu năng

  • Thử nghiệm được thực hiện theo cấu hình V3 với các điều kiện sau
    • 8K tokens mỗi batch
    • hidden dimension 7168
    • top 8 experts
    • FP8 dispatch
    • BF16 combine
  • Các kết quả chính gồm
    • 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
  • Các giá trị đo là băng thông logic; ví dụ, 90 GB/s của EP 8 x 2 bao gồm cả local rank traffic
  • V2 đạt hiệu năng đỉnh cao hơn tới 1,3 lần và giảm tới 4 lần số lượng SM so với V1
  • Kết quả cho cấu hình EP lớn hơn hiện chưa được đưa vào; người dùng được khuyến nghị tự benchmark
  • Theo kinh nghiệm nội bộ, các kernel được kỳ vọng vẫn bão hòa băng thông phần cứng khi mở rộng quy mô
  • Dữ liệu hiệu năng V1 có tại docs/legacy.md

Cài đặt và yêu cầu

  • Các yêu cầu gồm có
    • GPU Hopper SM90 hoặc kiến trúc hỗ trợ SM90 PTX ISA
    • Python 3.8 trở lên
    • CUDA 12.3 trở lên cho GPU SM90
    • PyTorch 2.10 trở lên
    • NCCL 2.30.4 trở lên
    • NVLink cho giao tiếp nội nút
    • Mạng RDMA cho giao tiếp liên nút
  • Khuyến nghị cài NCCL bằng pip để DeepEP có thể tự động phát hiện trong môi trường Python
pip install "nvidia-nccl-cu13>=2.30.4" --no-deps
  • Để hỗ trợ các phương thức legacy, dự án cũng phụ thuộc vào NVSHMEM; xem hướng dẫn cài đặt tại NVSHMEM Installation Guide
  • Ví dụ chạy build và test trong quá trình phát triển như sau
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
  • Việc cài đặt được thực hiện bằng lệnh sau
python setup.py install
  • Sau khi cài đặt, có thể import deep_ep để sử dụng trong dự án Python

Giao diện xoay quanh ElasticBuffer

  • Trong V2, mọi phép toán EP đều được hợp nhất dưới một giao diện ElasticBuffer duy nhất
    • Xử lý cả API thông lượng cao và độ trễ thấp trong cùng một giao diện
    • Bộ đệm có thể được khởi tạo bằng cách chỉ định trực tiếp cấu hình MoE
    • Tính toán số SM và QP tối ưu theo phương pháp phân tích
  • Ví dụ khởi tạo bộ đệm sử dụng ElasticBuffer.get_buffer_size_hint() để tính kích thước cần thiết và kiểm tra khả năng tái sử dụng bộ đệm hiện có
  • Khi tạo bộ đệm mới, có thể chỉ định num_max_tokens_per_rank, hidden, num_topk, use_fp8_dispatch v.v.
  • Dùng _buffer.get_theoretical_num_sms(num_experts, num_topk) để lấy số SM lý thuyết dùng cho kernel truyền thông
  • Nếu chỉ định trực tiếp num_sms trong lời gọi dispatch và combine, giá trị tính toán này sẽ bị ghi đè

Mẫu sử dụng cho huấn luyện, prefill và decoding

  • Trong huấn luyện hoặc prefill suy luận, MoE dispatch định tuyến token tới expert tương ứng trên mọi rank
    • Hỗ trợ đầu vào BF16 và FP8
    • handle chứa metadata định tuyến cần cho lần gọi combine sau đó
    • handle.num_recv_tokens_per_expert_list cung cấp số lượng token theo từng expert cần cho GEMM
  • Backward pass của MoE dispatch thực tế được xử lý bằng combine
  • MoE combine reduce đầu ra của expert về rank ban đầu
  • Backward pass của MoE combine thực tế được xử lý bằng dispatch
  • Việc chồng lấp truyền thông và tính toán được quản lý qua giao diện EventOverlap
    • Có thể thực hiện tính toán độc lập trong khi truyền thông đang diễn ra
    • Đồng bộ compute stream bằng event.current_stream_wait() trước khi dùng kết quả
  • Trong suy luận decoding cũng dùng cùng ElasticBuffer
    • Có thể tái sử dụng metadata định tuyến bằng cached_handle khi quyết định gating không thay đổi
    • Mẫu này tránh việc tính lại layout và đồng bộ CPU

Biến môi trường và giá trị cố định khi build

  • Cấu hình chung
    • EP_BUFFER_DEBUG: in thông tin debug về khởi tạo bộ đệm, xấp xỉ SM và backend
    • EP_SUPPRESS_NCCL_CHECK: bỏ kiểm tra sai lệch phiên bản NCCL
    • EP_AVOID_RECORD_STREAM: tránh record_stream cho tensor đầu ra
    • EP_NUM_TOPK_IDX_BITS: ghi đè số bit mã hóa top-k index
  • Cấu hình mạng
    • EP_NIC_NAME: tên NIC mặc định dùng để truy vấn thuộc tính NIC, mặc định mlx5_0
    • EP_OVERRIDE_RDMA_SL: ghi đè RDMA service level index
    • EP_DISABLE_GIN: vô hiệu hóa NCCL Gin backend
  • Cấu hình JIT
    • EP_JIT_CACHE_DIR: thư mục cache kernel đã biên dịch, mặc định $HOME/.deep_ep
    • EP_JIT_NVCC_COMPILER: đường dẫn trình biên dịch NVCC
    • EP_JIT_CPP_STANDARD: phiên bản chuẩn C++, mặc định 20
    • EP_JIT_DUMP_PTX, EP_JIT_DUMP_SASS, EP_JIT_DUMP_ASM: thiết lập dump đầu ra PTX·SASS
  • Một số biến môi trường hoạt động theo kiểu persistent
    • Được ghi nhận tại thời điểm build và đưa vào làm giá trị mặc định của gói cài đặt
    • Nếu không bị ghi đè bởi biến môi trường hiện tại khi import, các giá trị mặc định này sẽ tự động được áp dụng
    • Các biến mục tiêu là EP_JIT_CACHE_DIR, EP_JIT_PRINT_COMPILER_COMMAND, EP_NUM_TOPK_IDX_BITS, EP_NCCL_ROOT_DIR
  • Xem thêm chi tiết tại test_ep.py hoặc tài liệu Python

Khuyến nghị cấu hình mạng

  • DeepEP đã được kiểm thử đầy đủ trên mạng InfiniBand
  • Về lý thuyết cũng tương thích với RDMA over Converged Ethernet, tức RoCE
  • Cách ly lưu lượng
    • Được hỗ trợ qua Virtual Lanes của InfiniBand
    • Khuyến nghị tách workload expert-parallel và các workload khác sang các virtual lane khác nhau
    • Trong V2, có thể điều khiển phân bổ virtual lane bằng tham số sl_idx hoặc biến môi trường EP_OVERRIDE_RDMA_SL
  • Adaptive routing
    • Là tính năng định tuyến nâng cao giúp switch InfiniBand phân phối lưu lượng đồng đều qua nhiều đường đi
    • Khuyến nghị bật trong mọi điều kiện tải mạng, ngay cả khi có phát sinh thêm độ trễ
  • Congestion control
    • Nên tắt vì gây hại cho băng thông tối đa
    • Nếu không thể tránh nghẽn mạng, nên gán workload đó vào virtual lane ưu tiên thấp
  • PCI atomic mode
    • Nếu phần cứng hỗ trợ, nên cấu hình PCI_ATOMIC_MODE của NIC để cải thiện hiệu năng RDMA atomic operation
sudo mlxconfig -y -d mlx5_$i set PCI_ATOMIC_MODE=4

Nhánh thử nghiệm và các fork cộng đồng

  • Nhánh thử nghiệm
    • Zero-copy: loại bỏ sao chép giữa tensor PyTorch và bộ đệm truyền thông để giảm mạnh mức dùng SM của kernel thông thường
    • Eager: dùng giao thức độ trễ thấp để loại bỏ extra RTT latency do RDMA atomic OP bổ sung
    • Hybrid-EP: triển khai backend mới dùng TMA instructions, hỗ trợ mức dùng SM tối thiểu, miền NVLink lớn hơn, chồng lấp truyền thông-tính toán tinh hạt ở single-batch, kernel PCIe và hỗ trợ NVFP4
    • AntGroup-Opt: chuỗi tối ưu hóa do AntGroup Network Platform Department thực hiện
    • Mori-EP: hỗ trợ chế độ độ trễ thấp ROCm/AMD GPU dựa trên backend MORI
    • nvDev: nhánh dựa trên V2 bao gồm các tính năng CUDA mới nhất như Compute Fabric Transport
  • Fork cộng đồng
    • uccl/uccl-ep: hỗ trợ chạy DeepEP trên GPU dị chủng như Nvidia và AMD, cùng NIC như EFA, Broadcom và CX7
    • Infrawaves/DeepEP_ibrc_dual-ports_multiQP: bổ sung giải pháp multi-QP và hỗ trợ NIC dual-port cho IBRC transport
    • antgroup/DeepXTrace: công cụ chẩn đoán giúp tìm slow rank hiệu quả và chính xác
    • ROCm/mori: thư viện truyền thông thế hệ mới của AMD cho các workload AI nhạy cảm hiệu năng như Wide EP, KVCache transfer và Collectives

Giấy phép và trích dẫn

  • DeepEP V2 được xây dựng trên NCCL Gin backend của NCCL
  • Mã nguồn trong kho được phát hành theo MIT License
  • Mục trích dẫn là DeepEP: an efficient expert-parallel communication library, với năm ghi là 2025

Bài viết liên quan

1 bình luận

 
GN⁺ 2025-02-26
Ý kiến trên Hacker News

  • Tìm ra và dùng lệnh PTX không được tài liệu hóa ld.global.nc.L1::no_allocate.L2::256B để đạt hiệu năng cực hạn
    Lệnh này truy cập bộ nhớ GPU volatile bằng bộ sửa đổi PTX chỉ đọc không nhất quán .nc, nên có thể gây ra hành vi không xác định
    Tuy nhiên, trên kiến trúc Hopper, khi dùng cùng .L1::no_allocate thì qua thử nghiệm được cho là bảo đảm tính đúng đắn, còn hiệu năng thì tốt hơn rất nhiều

    • Trên thực tế, liệu NVIDIA sau này có thể âm thầm thay đổi hành vi của lệnh ngoài tài liệu này trên kiến trúc mới, dù cố ý hay không, và lật ngược cuộc chơi không?

  • Cảm giác như một đứa trẻ bước vào cửa hàng kẹo vậy
    Có quá nhiều thủ thuật mà nếu chỉ nhìn bài báo để đảo ngược kỹ thuật cho đúng thì sẽ mất rất lâu, và hy vọng các công bố tuần này sẽ mở ra một thời kỳ phục hưng trong đó MoE trở thành mô hình học thuật mặc định

    • Nhìn từ góc độ này thì tôi không hiểu chuyện gì đang xảy ra giữa thực hành mô hình tối tân ngoài đời và các mô hình học thuật
      Từ GPT-4 trở đi, nhóm trước về cơ bản đã toàn là MoE, nhưng các mô hình công khai, trừ DeepSeek V3 và Mixtral, thường vẫn luôn là mô hình dense

  • Không thể không thích đội này được
    Họ đang đẩy xa ranh giới của open source vì tất cả mọi người

    • Là kiểu viết cách ra như Open AI™
    • Thực ra thì không hẳn là open source
      Nếu muốn xem một mô hình open source thật sự, hãy xem OLMo 2 của AI2: https://allenai.org/blog/olmo2
      Họ thực sự chia sẻ mọi thứ cần thiết để tái tạo mô hình, kể cả chính dữ liệu
      Trong liên kết trên họ cũng nói: “vì khoa học mở hoàn toàn đòi hỏi nhiều hơn các trọng số công khai, chúng tôi vui mừng chia sẻ với cộng đồng mô hình hóa ngôn ngữ rộng hơn bản cập nhật OLMo mới, bao gồm trọng số, dữ liệu, mã, recipe, checkpoint trung gian và mô hình tinh chỉnh theo chỉ dẫn”

  • Zuckerberg nên ngừng tuyên bố rằng Meta công bố AI dưới dạng open source
    Họ thậm chí còn chạy quảng cáo TV, nhưng thực tế chỉ công khai trọng số chứ không có mã
    AI open source thật sự chỉ có DeepSeek

    • Nói chặt chẽ thì DeepSeek cũng không open source bằng OLMo hay Open Euro
      Vì họ không công khai dữ liệu
    • DeepSeek rõ ràng không phải open source thật sự
      Muốn là open source thì phải dùng giấy phép open source thật sự như OSI liệt kê, và phải chia sẻ mã tiền huấn luyện/hậu huấn luyện, mã liên quan đến tuning, mã đánh giá, mọi thứ liên quan đến an toàn/kiểm duyệt, và có lẽ cả toàn bộ dữ liệu huấn luyện
      Nếu không thì không thể tái tạo trọng số, và việc chia sẻ trọng số cũng giống như chia sẻ một chương trình đã biên dịch
      Theo tôi biết, mô hình cạnh tranh mà thật sự open source chỉ có OLMo 2 của AI2: https://allenai.org/blog/olmo2
      Gần đây họ còn công bố một ứng dụng chạy suy luận ngay trên thiết bị, cái này cũng open source: https://allenai.org/blog/olmoe-app
      Họ còn có một mô hình khác tên Tülu 3, được cho là có hiệu năng tốt hơn DeepSeek V3: https://allenai.org/blog/tulu-3-405B
    • Meta đã trau chuốt PyTorch hơn 10 năm rồi
      Gần như mọi thứ cần để huấn luyện LLM, kể cả công nghệ mới nhất, đều nằm trong đó
      Còn cần gì nữa? Phần mã đặc thù cho hạ tầng của Meta à?
    • PyTorch có tính vào không?
    • Trọng số công khai = một cục binary
      Vậy là quay lại mô hình FREEWARE / SHAREWARE
      Nên dùng cách gọi như thế cho các trọng số “công khai”

  • Cung cấp giao tiếp all-to-all hiệu quả và được tối ưu hóa, hỗ trợ trong nút và giữa các nút qua NVLink và RDMA, kernel thông lượng cao cho huấn luyện và prefill trong suy luận, kernel độ trễ thấp cho decoding trong suy luận, hỗ trợ dispatch FP8 native, và điều khiển tài nguyên GPU linh hoạt để chồng lấp tính toán-giao tiếp
    X: https://x.com/deepseek_ai/status/1894211757604049133

  • Động cơ của công việc DeepSeek có thể là sai trái
    Chẳng hạn có thể là một nỗ lực được nhà nước hậu thuẫn nhằm đưa lợi thế đi trước của Mỹ trong AI về 0, nhưng hiệu ứng ròng đối với toàn thế giới thì đơn giản là tuyệt vời
    Trong trường hợp tệ nhất, tức là ngay cả khi họ làm vì lý do sai, tôi vẫn biết ơn DeepSeek; họ đang thực sự làm điều mà OpenAI đã nói dối với cả thế giới suốt nhiều năm rằng họ sẽ làm

    • Trong lĩnh vực quan hệ quốc tế, đúng sai không áp dụng mạnh đến vậy
      Việc công bố cái này dưới dạng open source có “sai” hơn lệnh cấm xuất khẩu GPU Nvidia cao cấp không?
      Việc DeepSeek công bố open source có lẽ chỉ là, với sự đồng thuận của Đảng Cộng sản Trung Quốc, một điều đồng thời có lợi cho cả Đảng Cộng sản Trung Quốc lẫn cộng đồng AI open source rộng lớn hơn, chứ không nên hiểu là một lập trường nguyên tắc nào đó
      Tìm cách xóa bỏ lợi thế cạnh tranh của quốc gia khác là hoạt động chính của mọi chính phủ, lớn cũng như nhỏ

  • Đây là vòng công bố open source thứ 2 dưới giấy phép MIT từ công ty Open AI™ thật sự
    Một lần nữa, DeepSeek cởi mở hơn công ty trị giá 157 tỷ USD tự nhận là “Open” kia
    Hầu như chẳng ai nói về Llama của Meta, nhưng mọi người nên kỳ vọng Llama 4 sẽ ra mắt cùng khả năng suy luận
    Mục tiêu là không bị nghiền nát ở giữa cuộc đua tiến về 0

  • Trong lúc Mỹ lục hóa đơn GPU ở Singapore để kiểm tra DeepSeek có chỉ dùng H800 hay không, phần còn lại của thế giới có thể chạy các tối ưu hóa này trên những H100 nguyên vẹn phải không?
    Vì lệnh trừng phạt của Mỹ khiến H100 khó mua hoặc khó tiếp cận, trong khi họ vẫn tiếp tục giả vờ, do sự ngạo mạn của Mỹ, rằng mệnh lệnh của họ bao trùm cả thế giới?
    Hiểu như vậy có đúng không?

  • Lần này đã bao gồm PTX mà mọi người chờ đợi chưa?

    • Đúng, có một phần trong thư mục csrc/kernels
      Tìm asm là sẽ thấy các chỗ dùng
    • Cần giải thích cho những người còn lại vì sao PTX mà mọi người chờ đợi lại quan trọng đến thế

  • Gợi nhớ thập niên 80–90, khi người ta hack assembly hoặc tìm các lệnh không được tài liệu hóa để vắt kiệt hiệu năng CPU
    Có lẽ một ngày nào đó compiler sẽ tối ưu đủ tốt, hoặc GPU sẽ mạnh đến mức các thủ thuật như thế này không còn tạo ra khác biệt lớn nữa, giống CPU ngày nay