DeepSeek เปิดตัวไลบรารีโอเพนซอร์ส DeepEP สำหรับการฝึกและอนุมาน MoE

 (github.com/deepseek-ai)
1 คะแนน โดย GN⁺ 2025-02-26 | 1 ความคิดเห็น | แชร์ทาง WhatsApp

DeepEP

DeepEP เป็นไลบรารีสื่อสารสำหรับ Mixture-of-Experts (MoE) และ expert parallelism (EP) โดยมีเคอร์เนล all-to-all สำหรับ GPU ที่มีความเร็วสูงและหน่วงต่ำ ซึ่งเป็นที่รู้จักในชื่อ MoE dispatch และ combine นอกจากนี้ยังรองรับการคำนวณความแม่นยำต่ำรวมถึง FP8 อีกด้วย โดยสอดคล้องกับอัลกอริทึม group-limited gating ที่นำเสนอในงานวิจัย DeepSeek-V3 จึงมีเคอร์เนลที่ปรับแต่งการส่งผ่านแบนด์วิดท์โดเมนแบบไม่สมมาตรให้เหมาะสม โดยส่งข้อมูลจากโดเมน NVLink ไปยังโดเมน RDMA เคอร์เนลเหล่านี้ให้ throughput สูง จึงเหมาะกับงานฝึกและงาน inference prefill อีกทั้งยังรองรับการควบคุมจำนวน SM (Streaming Multiprocessors) สำหรับ inference decoding ที่ไวต่อความหน่วง DeepEP มีเคอร์เนลหน่วงต่ำที่ใช้ RDMA ล้วนเพื่อลด latency ให้น้อยที่สุด ไลบรารียังนำเสนอวิธีซ้อนทับการสื่อสารและการคำนวณแบบอิง hook ซึ่งไม่ยึดครองทรัพยากร SM

ประสิทธิภาพ

เคอร์เนลทั่วไปที่ใช้การส่งผ่าน NVLink และ RDMA

  • ทดสอบเคอร์เนลทั่วไปบน H800 โดยใช้แบนด์วิดท์สูงสุดของ NVLink ราว 160 GB/s และเชื่อมต่อกับการ์ดเครือข่าย RDMA CX7 InfiniBand 400 Gb/s (~50 GB/s แบนด์วิดท์สูงสุด)
  • ใช้การตั้งค่าการพรีเทรนของ DeepSeek-V3/R1 (4096 โทเค็นต่อแบตช์, hidden 7168, top 4 groups, top 8 experts, dispatch แบบ FP8 และ combine แบบ BF16)

เคอร์เนลหน่วงต่ำที่ใช้ RDMA ล้วน

  • ทดสอบเคอร์เนลหน่วงต่ำบน H800 โดยเชื่อมต่อกับการ์ดเครือข่าย RDMA CX7 InfiniBand 400 Gb/s (~50 GB/s แบนด์วิดท์สูงสุด)
  • ใช้การตั้งค่าโปรดักชันทั่วไปของ DeepSeek-V3/R1 (128 โทเค็นต่อแบตช์, hidden 7168, top 8 experts, dispatch แบบ FP8 และ combine แบบ BF16)

เริ่มต้นอย่างรวดเร็ว

ข้อกำหนด

  • GPU Hopper (อาจรองรับสถาปัตยกรรมหรืออุปกรณ์เพิ่มเติมในอนาคต)
  • Python 3.8 ขึ้นไป
  • CUDA 12.3 ขึ้นไป
  • PyTorch 2.1 ขึ้นไป
  • NVLink สำหรับการสื่อสารภายในโหนด
  • เครือข่าย RDMA สำหรับการสื่อสารระหว่างโหนด

ดาวน์โหลดและติดตั้ง dependency ของ NVSHMEM

DeepEP พึ่งพา NVSHMEM ที่มีการปรับแก้ จึงต้องติดตั้งตามคู่มือการติดตั้ง

การกำหนดค่าเครือข่าย

DeepEP ได้รับการทดสอบอย่างครบถ้วนบนเครือข่าย InfiniBand และในทางทฤษฎีก็สามารถเข้ากันได้กับ RDMA over Converged Ethernet (RoCE) ด้วย

การแยกทราฟฟิก

InfiniBand รองรับการแยกทราฟฟิกผ่าน Virtual Lanes (VL) เพื่อป้องกันการรบกวนระหว่างทราฟฟิกแต่ละประเภท แนะนำให้แยกงานไปยัง virtual lane ดังนี้:

  • งานที่ใช้เคอร์เนลทั่วไป
  • งานที่ใช้เคอร์เนลหน่วงต่ำ
  • งานประเภทอื่น

ใน DeepEP สามารถควบคุมการจัดสรร virtual lane ได้โดยตั้งค่าตัวแปรสภาพแวดล้อม NVSHMEM_IB_SL

Adaptive routing

Adaptive routing เป็นความสามารถด้านการกำหนดเส้นทางขั้นสูงที่มีในสวิตช์ InfiniBand ซึ่งช่วยกระจายทราฟฟิกให้สมดุล across หลายเส้นทาง ปัจจุบันเคอร์เนลหน่วงต่ำรองรับ adaptive routing แต่เคอร์เนลทั่วไปยังไม่รองรับ (แต่อาจรองรับในเร็ว ๆ นี้) หากเปิด adaptive routing กับเคอร์เนลทั่วไประหว่างโหนด อาจเกิด deadlock หรือปัญหาข้อมูลเสียหายได้ สำหรับเคอร์เนลหน่วงต่ำ การเปิด adaptive routing สามารถขจัดความแออัดของเครือข่ายที่เกิดจาก routing collision ได้ทั้งหมด แต่จะเพิ่มความหน่วงเพิ่มเติม เพื่อให้ได้ประสิทธิภาพที่เหมาะสมที่สุด แนะนำการตั้งค่าดังนี้:

  • เปิด adaptive routing ในสภาพแวดล้อมที่มีภาระเครือข่ายสูง
  • ใช้ static routing ในสภาพแวดล้อมที่มีภาระเครือข่ายต่ำ

การควบคุมความแออัด

เนื่องจากไม่พบความแออัดที่มีนัยสำคัญในสภาพแวดล้อมโปรดักชัน จึงปิดการควบคุมความแออัดไว้

อินเทอร์เฟซและตัวอย่าง

การใช้งานตัวอย่างในงานฝึกโมเดลหรือ inference prefill

เคอร์เนลทั่วไปสามารถใช้ในขั้นตอนการฝึกโมเดล หรือในขั้น inference prefill (โดยไม่มีส่วน backward)

การใช้งานตัวอย่างใน inference decoding

เคอร์เนลหน่วงต่ำสามารถใช้ในขั้นตอน inference decoding

ข้อควรระวัง

  • เพื่อประสิทธิภาพสูงสุด ได้มีการค้นพบและใช้งานคำสั่ง PTX ที่ไม่มีในเอกสารคือ ld.global.nc.L1::no_allocate.L2::256B คำสั่งนี้ก่อให้เกิดพฤติกรรมที่ไม่ถูกกำหนดไว้ โดยใช้ PTX modifier สำหรับการอ่านแบบไม่สอดคล้องและอ่านอย่างเดียวเพื่อเข้าถึงหน่วยความจำ GPU แบบ volatile อย่างไรก็ตาม เมื่อทดสอบบนสถาปัตยกรรม Hopper ด้วย .L1::no_allocate พบว่าประสิทธิภาพดีขึ้นอย่างมาก หากเคอร์เนลไม่ทำงานบนแพลตฟอร์มอื่น สามารถปิดการใช้งานนี้ได้โดยเพิ่ม DISABLE_AGGRESSIVE_PTX_INSTRS=1 ใน setup.py หรือเปิด issue
  • เพื่อประสิทธิภาพที่ดียิ่งขึ้นบนคลัสเตอร์ แนะนำให้รันการทดสอบทั้งหมดและใช้การตั้งค่า auto-tuning ที่เหมาะสมที่สุด โดยค่าตั้งต้นได้รับการปรับแต่งบนคลัสเตอร์ภายในของ DeepSeek

ใบอนุญาต

โค้ดรีโพซิทอรีนี้เผยแพร่ภายใต้สัญญาอนุญาต MIT และโค้ดที่อ้างอิง NVSHMEM (รวมถึง csrc/kernels/ibgda_device.cuh และ third-party/nvshmem.patch) อยู่ภายใต้ NVSHMEM SLA

บทความที่เกี่ยวข้อง

1 ความคิดเห็น

 
GN⁺ 2025-02-26
ความคิดเห็นจาก Hacker News
  • เพื่อรีดประสิทธิภาพให้ถึงขีดสุด พวกเขาค้นพบและใช้งานคำสั่ง PTX ที่ไม่มีอยู่ในเอกสาร คำสั่งนี้ใช้ตัวปรับแต่ง PTX แบบอ่านอย่างเดียวที่ไม่สอดคล้องกันเพื่อเข้าถึงหน่วยความจำ GPU แบบ volatile ซึ่งอาจนำไปสู่พฤติกรรมที่ไม่ถูกกำหนดไว้ อย่างไรก็ตาม มีการยืนยันความถูกต้องเมื่อทดสอบด้วย .L1::no_allocate บนสถาปัตยกรรม Hopper และประสิทธิภาพก็ดีขึ้นมาก
  • Zuckerberg ควรหยุดอ้างว่า Meta ทำโอเพนซอร์ส AI ได้แล้ว พวกเขาเปิดเผยแค่น้ำหนักโมเดล ไม่ใช่โค้ด AI โอเพนซอร์สที่แท้จริงมีแค่ DeepSeek เท่านั้น
  • รู้สึกเหมือนเด็กที่อยู่ในร้านขายลูกกวาด บางเทคนิคเหล่านี้คงใช้เวลานานมากกว่าจะทำ reverse engineer จากตัวบทความวิจัยได้อย่างถูกต้อง หวังว่าการประกาศในสัปดาห์นี้จะเป็นจุดเริ่มต้นของยุคฟื้นฟูที่ใช้ MoE เป็นโมเดลวิชาการพื้นฐาน
  • ไม่รักคนกลุ่มนี้คงไม่ได้จริงๆ พวกเขากำลังผลักขอบเขตของโอเพนซอร์สเพื่อพวกเราทุกคนอย่างแท้จริง ขอบคุณที่แบ่งปัน
    • การสื่อสาร all-to-all ที่มีประสิทธิภาพและผ่านการปรับแต่งมาอย่างดี
    • รองรับทั้งภายในโหนดและข้ามโหนดผ่าน NVLink และ RDMA
    • เคอร์เนล throughput สูงสำหรับการเทรนและการ prefill ในงานอนุมาน
    • เคอร์เนล latency ต่ำสำหรับการ decode ในงานอนุมาน
    • รองรับ native FP8 dispatch
    • การควบคุมทรัพยากร GPU ที่ยืดหยุ่นสำหรับการซ้อนทับการคำนวณกับการสื่อสาร
  • แรงจูงใจเบื้องหลังงานของ DeepSeek อาจไม่ถูกต้องนัก (เช่น ความพยายามที่ได้รับการสนับสนุนจากรัฐเพื่อทำลายความได้เปรียบด้านการเป็นผู้นำของสหรัฐใน AI) แต่ผลลัพธ์ที่เกิดขึ้นกับคนทั่วโลกนั้นยอดเยี่ยมอย่างเรียบง่าย
    • แม้ในกรณีเลวร้ายที่สุด (ถ้าพวกเขาทำสิ่งนี้ด้วยเหตุผลที่ผิด) ก็ยังต้องขอบคุณ DeepSeek พวกเขากำลังทำสิ่งที่ OpenAI โกหกคนทั้งโลกมาหลายปีว่าตัวเองทำอยู่จริงๆ
    • น่าทึ่งจริงๆ
  • สงสัยว่า PTX ที่ทุกคนรอคอยจะถูกรวมมาในครั้งนี้หรือไม่
  • คำสั่ง PTX ที่กล่าวถึงในรายงานทางเทคนิคควรถูกเชื่อมโยงกับโค้ดที่นี่
  • ขณะที่สหรัฐกำลังไล่ตามตรวจใบเสร็จ GPU ในสิงคโปร์เพื่อยืนยันว่า DeepSeek ใช้แค่ H800 หรือไม่ คนที่เหลือของโลกสามารถรันการปรับแต่งเหล่านี้บน H100 เต็มรูปแบบได้
    • สงสัยว่าพวกเขากำลังแกล้งทำเป็นว่าการได้มาหรือเข้าถึง H100 เป็นเรื่องยาก เพราะความหยิ่งผยองที่เชื่อว่ามาตรการคว่ำบาตรและคำสั่งของสหรัฐครอบคลุมไปทั้งโลกหรือเปล่า
  • นี่คือการปล่อยโอเพนซอร์สครั้งที่สองจากบริษัท "Open AI™" ตัวจริง และอยู่ภายใต้สัญญาอนุญาต MIT
    • DeepSeek เปิดกว้างกว่าบริษัทที่อ้างมูลค่า $157B+ เสียอีก
    • แทบไม่มีใครพูดถึง Llama ของ Meta และทุกคนก็ควรคาดว่า Llama 4 จะเปิดตัวด้วยเหตุผลบางอย่าง
    • เป้าหมายคืออย่าไปติดอยู่ตรงกลางในการแข่งลดราคาจนเหลือศูนย์