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)DeepEP
DeepEP là một thư viện truyền thông dành cho Mixture-of-Experts (MoE) và expert parallelism (EP). Thư viện cung cấp các kernel all-to-all GPU tốc độ cao, độ trễ thấp, thường được dùng cho MoE dispatch và combine. Ngoài ra, thư viện còn hỗ trợ tính toán độ chính xác thấp, bao gồm cả FP8. Phù hợp với thuật toán gating giới hạn theo nhóm được đề xuất trong bài báo DeepSeek-V3, DeepEP cung cấp các kernel tối ưu hóa việc truyền băng thông miền bất đối xứng, chuyển dữ liệu từ miền NVLink sang miền RDMA. Các kernel này mang lại thông lượng cao, phù hợp cho huấn luyện và các tác vụ prefill trong suy luận. Ngoài ra, thư viện còn hỗ trợ điều khiển số lượng SM (Streaming Multiprocessors). Đối với tác vụ decoding suy luận nhạy cảm với độ trễ, DeepEP bao gồm các kernel độ trễ thấp sử dụng RDMA thuần để giảm thiểu độ trễ. Thư viện này cũng giới thiệu phương pháp chồng lấp giao tiếp - tính toán dựa trên hook, không chiếm dụng tài nguyên SM.
Hiệu năng
Kernel thông thường sử dụng truyền NVLink và RDMA
- Trên H800, kernel thông thường được kiểm thử với băng thông NVLink tối đa khoảng 160 GB/s, kết nối với card mạng RDMA CX7 InfiniBand 400 Gb/s (~50 GB/s băng thông tối đa).
- Tuân theo cấu hình pre-training DeepSeek-V3/R1 (4096 token mỗi batch, hidden 7168, top 4 group, top 8 expert, dispatch FP8 và combine BF16).
Kernel độ trễ thấp sử dụng RDMA thuần
- Trên H800, kernel độ trễ thấp được kiểm thử với card mạng RDMA CX7 InfiniBand 400 Gb/s (~50 GB/s băng thông tối đa).
- Tuân theo cấu hình production DeepSeek-V3/R1 điển hình (128 token mỗi batch, hidden 7168, top 8 expert, dispatch FP8 và combine BF16).
Bắt đầu nhanh
Yêu cầu
- GPU Hopper (trong tương lai có thể hỗ trợ thêm nhiều kiến trúc hoặc thiết bị hơn)
- Python 3.8 trở lên
- CUDA 12.3 trở lên
- PyTorch 2.1 trở lên
- NVLink cho giao tiếp trong node
- Mạng RDMA cho giao tiếp liên node
Tải xuống và cài đặt phụ thuộc NVSHMEM
DeepEP phụ thuộc vào một phiên bản NVSHMEM đã được chỉnh sửa. Cần tham khảo hướng dẫn cài đặt để tiến hành cài đặt.
Cấu hình mạng
DeepEP đã được kiểm thử đầy đủ trên mạng InfiniBand và về mặt lý thuyết cũng tương thích với RDMA over Converged Ethernet (RoCE).
Cô lập lưu lượng
InfiniBand hỗ trợ cô lập lưu lượng thông qua Virtual Lanes (VL). Để tránh nhiễu giữa các loại lưu lượng khác nhau, nên tách workload theo các làn ảo như sau:
- workload sử dụng kernel thông thường
- workload sử dụng kernel độ trễ thấp
- các workload khác
Trong DeepEP, có thể kiểm soát việc gán làn ảo bằng cách thiết lập biến môi trường NVSHMEM_IB_SL.
Định tuyến thích ứng
Định tuyến thích ứng là một tính năng định tuyến nâng cao do switch InfiniBand cung cấp, cho phép phân phối lưu lượng đồng đều qua nhiều đường đi. Hiện tại, kernel độ trễ thấp hỗ trợ định tuyến thích ứng, còn kernel thông thường thì chưa hỗ trợ (có thể sẽ sớm có). Nếu bật định tuyến thích ứng cho kernel liên node thông thường, có thể phát sinh tình trạng deadlock hoặc lỗi hỏng dữ liệu. Đối với kernel độ trễ thấp, việc bật định tuyến thích ứng có thể loại bỏ hoàn toàn tắc nghẽn mạng do xung đột định tuyến, nhưng sẽ gây thêm độ trễ. Để đạt hiệu năng tối ưu, khuyến nghị cấu hình như sau:
- Bật định tuyến thích ứng trong môi trường có tải mạng cao
- Dùng định tuyến tĩnh trong môi trường có tải mạng thấp
Kiểm soát tắc nghẽn
Vì không quan sát thấy tắc nghẽn đáng kể trong môi trường production, nên kiểm soát tắc nghẽn được tắt.
Giao diện và ví dụ
Sử dụng ví dụ trong huấn luyện mô hình hoặc prefill suy luận
Kernel thông thường có thể được sử dụng trong giai đoạn huấn luyện mô hình hoặc prefill suy luận (không bao gồm phần backward).
Sử dụng ví dụ trong decoding suy luận
Kernel độ trễ thấp có thể được sử dụng trong giai đoạn decoding suy luận.
Lưu ý
- Để đạt hiệu năng cực hạn, nhóm phát triển đã phát hiện và sử dụng lệnh PTX không được ghi nhận trong tài liệu là
ld.global.nc.L1::no_allocate.L2::256B. Lệnh này gây ra hành vi không xác định khi truy cập bộ nhớ GPU volatile bằng một PTX modifier đọc chỉ đọc không nhất quán. Tuy nhiên, khi kiểm thử với.L1::no_allocatetrên kiến trúc Hopper, hiệu năng được cải thiện đáng kể. Nếu kernel không hoạt động trên nền tảng khác, có thể vô hiệu hóa bằng cách thêmDISABLE_AGGRESSIVE_PTX_INSTRS=1vàosetup.pyhoặc tạo issue. - Để có hiệu năng tốt hơn trên cluster, nên chạy toàn bộ các bài kiểm thử và sử dụng cấu hình auto-tuning tối ưu. Cấu hình mặc định được tối ưu trên cluster nội bộ của DeepSeek.
Giấy phép
Kho mã này được phát hành theo giấy phép MIT, và phần mã tham chiếu tới NVSHMEM (bao gồm csrc/kernels/ibgda_device.cuh và third-party/nvshmem.patch) chịu sự điều chỉnh của NVSHMEM SLA.
1 bình luận
Ý kiến Hacker News
.L1::no_allocatetrên kiến trúc Hopper, và hiệu năng sẽ tốt hơn nhiều