DeepSeek ने MoE ट्रेनिंग और inference के लिए ओपन सोर्स DeepEP लाइब्रेरी जारी की

DeepEP

DeepEP, Mixture-of-Experts (MoE) और expert parallelism (EP) के लिए एक communication लाइब्रेरी है। यह हाई-स्पीड, लो-लेटेंसी all-to-all GPU kernels प्रदान करती है, जिन्हें MoE dispatch और combine के रूप में जाना जाता है। यह FP8 सहित low-precision operations को भी सपोर्ट करती है। DeepSeek-V3 पेपर में प्रस्तावित group-limited gating algorithm के अनुरूप, यह asymmetric-domain bandwidth forwarding को optimize करने वाले kernels प्रदान करती है, जो NVLink domain से RDMA domain तक डेटा forward करते हैं। ये kernels उच्च throughput देते हैं, इसलिए training और inference prefill workloads के लिए उपयुक्त हैं। साथ ही यह SM(Streaming Multiprocessors) count control को सपोर्ट करती है। latency-sensitive inference decoding के लिए, DeepEP में pure RDMA का उपयोग करने वाले low-latency kernels शामिल हैं, जो latency को न्यूनतम करते हैं। यह लाइब्रेरी hook-based communication-computation overlap method भी पेश करती है, जो SM resources को occupy नहीं करती।

प्रदर्शन

NVLink और RDMA forwarding का उपयोग करने वाले सामान्य kernels

• H800 पर लगभग 160 GB/s अधिकतम NVLink bandwidth और CX7 InfiniBand 400 Gb/s RDMA network card (~50 GB/s अधिकतम bandwidth) के साथ सामान्य kernels का परीक्षण किया गया।
• DeepSeek-V3/R1 pretraining configuration (प्रति batch 4096 tokens, 7168 hidden, top-4 groups, top-8 experts, FP8 dispatching और BF16 combining) का पालन किया गया।

pure RDMA का उपयोग करने वाले low-latency kernels

• H800 पर CX7 InfiniBand 400 Gb/s RDMA network card (~50 GB/s अधिकतम bandwidth) के साथ low-latency kernels का परीक्षण किया गया।
• सामान्य DeepSeek-V3/R1 production configuration (प्रति batch 128 tokens, 7168 hidden, top-8 experts, FP8 dispatching और BF16 combining) का पालन किया गया।

क्विक स्टार्ट

आवश्यकताएँ

• Hopper GPU (भविष्य में अधिक architectures या devices को सपोर्ट किया जा सकता है)
• Python 3.8 या उससे ऊपर
• CUDA 12.3 या उससे ऊपर
• PyTorch 2.1 या उससे ऊपर
• node के भीतर communication के लिए NVLink
• node के बीच communication के लिए RDMA network

NVSHMEM dependencies डाउनलोड और इंस्टॉल करना

DeepEP, संशोधित NVSHMEM पर निर्भर करता है। इसे इंस्टॉल करने के लिए installation guide देखें।

नेटवर्क कॉन्फ़िगरेशन

DeepEP का InfiniBand network पर पूरी तरह परीक्षण किया गया है, और सैद्धांतिक रूप से यह RDMA over Converged Ethernet (RoCE) के साथ भी compatible है।

ट्रैफ़िक आइसोलेशन

InfiniBand, Virtual Lanes (VL) के माध्यम से traffic isolation को सपोर्ट करता है। अलग-अलग प्रकार के traffic के बीच interference रोकने के लिए, virtual lanes में workloads को निम्न प्रकार से अलग करने की सिफारिश की जाती है:

• सामान्य kernels का उपयोग करने वाले workloads
• low-latency kernels का उपयोग करने वाले workloads
• अन्य workloads

DeepEP में NVSHMEM_IB_SL environment variable सेट करके virtual lane allocation को नियंत्रित किया जा सकता है।

adaptive routing

Adaptive routing, InfiniBand switches द्वारा प्रदान की जाने वाली एक उन्नत routing सुविधा है, जो traffic को कई paths में समान रूप से वितरित कर सकती है। वर्तमान में low-latency kernels adaptive routing को सपोर्ट करते हैं, लेकिन सामान्य kernels इसे सपोर्ट नहीं करते (हालांकि जल्द सपोर्ट मिल सकता है)। सामान्य inter-node kernels पर adaptive routing सक्षम करने से deadlock या data corruption जैसी समस्याएँ हो सकती हैं। low-latency kernels के लिए adaptive routing सक्षम करने पर routing collisions के कारण होने वाली network congestion पूरी तरह हटाई जा सकती है, लेकिन अतिरिक्त latency जुड़ती है। सर्वोत्तम प्रदर्शन के लिए निम्न configuration की सिफारिश की जाती है:

• अधिक network load वाले environments में adaptive routing सक्षम करें
• कम network load वाले environments में static routing का उपयोग करें

congestion control

Production environment में कोई महत्वपूर्ण congestion नहीं देखी गई, इसलिए congestion control को disabled रखा गया है।

इंटरफ़ेस और उदाहरण

मॉडल training या inference prefill में उदाहरणों का उपयोग

सामान्य kernels का उपयोग मॉडल training या inference prefill चरण (backward part के बिना) में किया जा सकता है।

inference decoding में उदाहरणों का उपयोग

low-latency kernels का उपयोग inference decoding चरण में किया जा सकता है।

ध्यान देने योग्य बातें

• अत्यधिक प्रदर्शन के लिए, undocumented PTX instruction ld.global.nc.L1::no_allocate.L2::256B खोजकर उपयोग किया गया है। यह instruction non-coherent read-only PTX modifier का उपयोग करके volatile GPU memory तक पहुँचने पर undefined behavior पैदा करता है। हालांकि, Hopper architecture पर .L1::no_allocate के साथ परीक्षण में प्रदर्शन काफी बेहतर पाया गया। यदि अन्य platforms पर kernels काम नहीं करते, तो setup.py में DISABLE_AGGRESSIVE_PTX_INSTRS=1 जोड़कर इसे disable किया जा सकता है या एक issue दर्ज किया जा सकता है।
• cluster में बेहतर प्रदर्शन के लिए सभी tests चलाने और optimal auto-tuning configuration का उपयोग करने की सिफारिश की जाती है। default configuration को DeepSeek के internal cluster में optimize किया गया है।

लाइसेंस

यह code repository MIT license के तहत जारी की गई है, और NVSHMEM को संदर्भित करने वाला code (csrc/kernels/ibgda_device.cuh और third-party/nvshmem.patch सहित) पर NVSHMEM SLA लागू होता है।

[GN⁺]

Hacker News टिप्पणियाँ

• अत्यधिक प्रदर्शन के लिए दस्तावेज़ों में दर्ज न होने वाले PTX निर्देश खोजे और इस्तेमाल किए गए। यह निर्देश असंगत read-only PTX modifier का उपयोग करके volatile GPU memory तक पहुँचता है, जिससे undefined behavior हो सकता है। हालांकि Hopper architecture पर .L1::no_allocate के साथ परखी गई शुद्धता सुनिश्चित है और प्रदर्शन काफी बेहतर होगा
• Zuckerberg को यह दावा करना बंद कर देना चाहिए कि Meta AI को open source कर रहा है। वे code नहीं, सिर्फ weights जारी कर रहे हैं। सच में open source AI तो सिर्फ DeepSeek है
• ऐसा महसूस हो रहा है जैसे कोई बच्चा कैंडी की दुकान में हो। इनमें से कुछ ट्रिक्स को papers के आधार पर सही तरीके से reverse engineer करने में बहुत समय लगेगा। उम्मीद है कि इस हफ्ते की घोषणाएँ MoE को डिफ़ॉल्ट academic model के रूप में इस्तेमाल करने वाले renaissance की शुरुआत करेंगी
• इन लोगों से प्यार न करना मुश्किल है। ये सच में हम सबके लिए open source की सीमाएँ आगे बढ़ा रहे हैं। साझा करने के लिए धन्यवाद
  • efficient और optimized all-to-all communication
  • NVLink और RDMA के ज़रिए intra-node और inter-node support
  • training और inference prefill के लिए high-throughput kernels
  • inference decoding के लिए low-latency kernels
  • native FP8 dispatch support
  • compute-communication overlap के लिए flexible GPU resource control
• DeepSeek के काम के पीछे की प्रेरणा गलत हो सकती है (जैसे AI में अमेरिका की बढ़त खत्म करने की राज्य-प्रायोजित कोशिश)। लेकिन वैश्विक स्तर पर इसका परिणाम बस शानदार है
  • सबसे खराब स्थिति में भी (अगर वे यह काम गलत कारणों से कर रहे हों) DeepSeek का धन्यवाद। वे वास्तव में वही कर रहे हैं जिसके बारे में OpenAI ने सालों तक दुनिया से झूठ बोला
  • वाकई कमाल है
• सोच रहा हूँ कि जिस PTX का सब इंतज़ार कर रहे थे, क्या वह इस बार शामिल है
• technical report में जिस PTX instruction का ज़िक्र है, उसे यहाँ code से जोड़ा जाना चाहिए
• जब अमेरिका सिंगापुर में GPU रसीदों को ट्रैक कर रहा है ताकि यह सत्यापित कर सके कि DeepSeek ने सिर्फ H800 का उपयोग किया या नहीं, तब बाकी दुनिया full H100 पर इन optimizations को चला सकती है
  • समझ नहीं आता कि अमेरिकी प्रतिबंधों और उनके आदेशों को पूरी दुनिया पर लागू मानने के अहंकार की वजह से क्या वे यह दिखावा कर रहे हैं कि H100 पाना या उस तक पहुँचना मुश्किल था
• यह असली "Open AI™" कंपनी की दूसरी open source release है, और यह MIT license के तहत है
  • DeepSeek उस कंपनी से ज़्यादा खुली है जो $157B+ का दावा करती है
  • लगभग कोई भी Meta की Llama के बारे में बात नहीं कर रहा, और सबको उम्मीद करनी चाहिए कि Llama 4 किसी वजह से जारी की जाएगी
  • लक्ष्य यह है कि zero तक की race में बीच में फँसकर न रह जाएँ