DeepSeek công bố mã nguồn mở DeepEP (2/5)

(github.com/deepseek-ai)

3 điểm bởi xguru 2025-02-25 | Chưa có bình luận nào. | Chia sẻ qua WhatsApp

Thư viện giao tiếp hiệu năng cao dành cho Mixture-of-Experts (MoE) và Expert Parallelism (EP)
Cung cấp kernel All-to-All dựa trên GPU để xử lý tốc độ cao các phép toán dispatch và combine của MoE
Hỗ trợ tính toán độ chính xác thấp như FP8
Áp dụng thuật toán group-limited gating được đề xuất trong bài báo DeepSeek-V3 để tối ưu forwarding băng thông miền bất đối xứng
- Ví dụ: tối ưu truyền dữ liệu NVLink → RDMA
- Cung cấp thông lượng cao, phù hợp cho huấn luyện và tác vụ prefilling suy luận
Bao gồm kernel độ trễ thấp chuyên dụng cho RDMA để phục vụ giải mã suy luận nhạy cảm với độ trễ
Cung cấp kỹ thuật chồng lấp giao tiếp-tính toán (không chiếm tài nguyên SM)

Hiệu năng

DeepEP đã kiểm thử hiệu năng trong môi trường GPU H800 và mạng RDMA CX7 InfiniBand 400Gb/s
Dựa trên cấu hình DeepSeek-V3/R1, áp dụng kiến trúc 4096 token mỗi batch, 7168 hidden node, top-4 group, top-8 expert, đồng thời sử dụng dispatch FP8 và combine BF16
Kết quả kiểm thử cho thấy giao tiếp trong node (dựa trên NVLink) đạt băng thông khoảng trên 150GB/s, còn giao tiếp liên node (dựa trên RDMA) đạt mức băng thông 40~47GB/s tùy theo số lượng expert
Khi số lượng expert tăng lên, băng thông RDMA có xu hướng tăng nhẹ (ví dụ: 43GB/s với 16 expert, 46GB/s với 64 expert)

Kết quả đo hiệu năng kernel độ trễ thấp cho thấy độ trễ (latency) giảm đáng kể so với kernel thông thường
Trong môi trường xử lý 128 token mỗi batch, độ trễ tăng theo số lượng expert nhưng băng thông RDMA được duy trì tương đối ổn định
Ví dụ, tăng từ 163 micro giây (us) với 8 expert lên 194 micro giây (us) với 256 expert
Ở phép toán combine, độ trễ cao hơn dispatch, đồng thời khi số lượng expert tăng thì băng thông RDMA có xu hướng giảm dần xuống dưới 40GB/s
Nói cách khác, kernel độ trễ thấp hoạt động rất nhanh với nhóm expert quy mô nhỏ, nhưng khi số lượng expert tăng thì độ trễ cũng tăng, nên cần một sự cân bằng phù hợp

Có thể cô lập lưu lượng bằng cách tận dụng Virtual Lanes (VL) của InfiniBand
Cách phân tách được khuyến nghị:
- Tác vụ dùng kernel thông thường
- Tác vụ dùng kernel độ trễ thấp
- Các tác vụ khác
Có thể thiết lập VL thông qua biến môi trường NVSHMEM_IB_SL

Hỗ trợ định tuyến thích ứng của switch InfiniBand
Có thể bật với kernel độ trễ thấp, nhưng cần tắt với kernel thông thường (nếu bật có nguy cơ làm hỏng dữ liệu)
Khuyến nghị cấu hình:
- Khi tải mạng cao: bật định tuyến thích ứng
- Khi tải mạng thấp: giữ định tuyến tĩnh

DeepEP được vận hành trong trạng thái tắt chức năng kiểm soát tắc nghẽn
Đã xác nhận rằng trong môi trường thực tế, tình trạng nghẽn mạng không nghiêm trọng

Sử dụng lệnh PTX không chính thức: dùng ld.global.nc.L1::no_allocate.L2::256B để tối ưu hiệu năng
- Trên kiến trúc Hopper, lệnh này hoạt động bình thường; trên nền tảng khác có thể vô hiệu hóa bằng cách đặt DISABLE_AGGRESSIVE_PTX_INSTRS=1
Khuyến nghị autotuning: để đạt hiệu năng tối ưu, cần kiểm thử hiệu năng theo từng cụm rồi mới áp dụng cấu hình