- ไลบรารีสื่อสารประสิทธิภาพสูง สำหรับ Mixture-of-Experts(MoE) และ Expert Parallelism(EP)
- มี เคอร์เนล All-to-All บน GPU เพื่อเร่งการประมวลผล MoE dispatch และ combine ให้ทำงานได้รวดเร็ว
- รองรับการคำนวณความละเอียดต่ำ เช่น FP8
- ใช้ อัลกอริทึม group-limited gating ที่เสนอในงานวิจัย DeepSeek-V3 เพื่อเพิ่มประสิทธิภาพ asymmetric domain bandwidth forwarding
- ตัวอย่าง: การปรับแต่งการส่งข้อมูล NVLink → RDMA
- ให้ throughput สูง เหมาะกับงานฝึกสอนและงาน inference prefilling
- มี เคอร์เนลหน่วงต่ำสำหรับ RDMA โดยเฉพาะ สำหรับ inference decoding ที่ไวต่อ latency
- มี เทคนิคการซ้อนทับการสื่อสาร-การคำนวณ (ไม่ใช้ทรัพยากร SM)
ประสิทธิภาพ
เคอร์เนลทั่วไป (การส่งผ่าน NVLink และ RDMA)
- DeepEP ทดสอบประสิทธิภาพในสภาพแวดล้อม H800 GPU และ เครือข่าย CX7 InfiniBand 400Gb/s RDMA
- อิงจากการตั้งค่า DeepSeek-V3/R1 โดยใช้โครงสร้าง 4096 โทเค็นต่อแบตช์, hidden nodes 7168 ตัว, top-4 groups, top-8 experts และใช้ FP8 dispatching กับ BF16 combine
- ผลการทดสอบพบว่า การสื่อสารภายในโหนด (บน NVLink) มีแบนด์วิดท์มากกว่า 150GB/s โดยประมาณ และ การสื่อสารระหว่างโหนด (บน RDMA) มีแบนด์วิดท์ระดับ 40~47GB/s ตามจำนวน experts
- เมื่อจำนวน experts มากขึ้น แบนด์วิดท์ RDMA มีแนวโน้มเพิ่มขึ้นเล็กน้อย (เช่น 43GB/s เมื่อมี experts 16 ตัว และ 46GB/s เมื่อมี experts 64 ตัว)
เคอร์เนลหน่วงต่ำ (RDMA ล้วน)
- ผลการวัดประสิทธิภาพของเคอร์เนลหน่วงต่ำพบว่า latency ลดลงอย่างมาก เมื่อเทียบกับเคอร์เนลทั่วไป
- ในสภาพแวดล้อมที่ประมวลผล 128 โทเค็นต่อแบตช์ latency เพิ่มขึ้นตามจำนวน experts แต่แบนด์วิดท์ RDMA ยังค่อนข้างคงที่
- ตัวอย่างเช่น อยู่ที่ 163 ไมโครวินาที(us) เมื่อมี experts 8 ตัว และเพิ่มเป็น 194 ไมโครวินาที(us) เมื่อมี experts 256 ตัว
- ในการทำงาน combine มี latency สูงกว่า dispatch และเมื่อจำนวน experts เพิ่มขึ้น แบนด์วิดท์ RDMA มีแนวโน้ม ค่อย ๆ ลดลงต่ำกว่า 40GB/s
- กล่าวคือ เคอร์เนลหน่วงต่ำทำงานได้เร็วมากกับกลุ่ม experts ขนาดเล็ก แต่เมื่อจำนวน experts มากขึ้น latency ก็จะเพิ่มขึ้น จึงต้องหาสมดุลที่เหมาะสม
การตั้งค่าเครือข่าย
การแยกทราฟฟิก (Traffic Isolation)
- สามารถแยกทราฟฟิกได้โดยใช้ Virtual Lanes(VL) ของ InfiniBand
- แนวทางการแยกที่แนะนำ:
- งานที่ใช้เคอร์เนลทั่วไป
- งานที่ใช้เคอร์เนลหน่วงต่ำ
- งานอื่น ๆ
- สามารถตั้งค่า VL ได้ผ่านตัวแปรแวดล้อม
NVSHMEM_IB_SL
การกำหนดเส้นทางแบบปรับตัว (Adaptive Routing)
- รองรับ adaptive routing ของสวิตช์ InfiniBand
- สามารถเปิดใช้กับเคอร์เนลหน่วงต่ำได้ แต่ต้องปิดกับเคอร์เนลทั่วไป (หากเปิดอาจเสี่ยงข้อมูลเสียหาย)
- คำแนะนำในการตั้งค่า:
- กรณีมีโหลดเครือข่ายสูง: เปิด adaptive routing
- กรณีมีโหลดเครือข่ายต่ำ: คงการ routing แบบคงที่
การควบคุมความแออัด (Congestion Control)
- DeepEP ทำงานโดยปิดฟังก์ชัน congestion control
- ยืนยันแล้วว่าในสภาพแวดล้อมจริง ความแออัดของเครือข่ายไม่ได้รุนแรง
ประเด็นทางเทคนิคสำคัญ
- ใช้คำสั่ง PTX ที่ไม่เป็นทางการ: ใช้
ld.global.nc.L1::no_allocate.L2::256B เพื่อปรับแต่งประสิทธิภาพ
- บนสถาปัตยกรรม Hopper ทำงานได้ตามปกติ แต่บนแพลตฟอร์มอื่นสามารถปิดได้ด้วยการตั้งค่า
DISABLE_AGGRESSIVE_PTX_INSTRS=1
- แนะนำให้ทำ auto-tuning: เพื่อให้ได้ประสิทธิภาพสูงสุด ควรทดสอบประสิทธิภาพในแต่ละคลัสเตอร์ก่อนนำค่าตั้งไปใช้
ยังไม่มีความคิดเห็น