DeepSeek ने MoE training और inference के लिए open source DeepEP लाइब्रेरी जारी की

 (github.com/deepseek-ai)
1 पॉइंट द्वारा GN⁺ 2025-02-26 | 1 टिप्पणियां | WhatsApp पर शेयर करें
  • DeepEP आधुनिक ML training·inference में Expert Parallelism (EP) पर केंद्रित एक high-performance communication लाइब्रेरी है, जो MoE dispatch/combine के लिए all-to-all GPU kernels और FP8 जैसी low-precision support प्रदान करती है
  • V2 रिलीज़ में EP को पूरी तरह refactor किया गया है, जिससे V1 की तुलना में कहीं कम SM resources के साथ समान या बेहतर performance मिलती है, और backend को NVSHMEM से हल्के NCCL Gin में बदला गया है
  • V3 configuration-आधारित tests में 8K tokens per batch, hidden 7168, top 8 experts, FP8 dispatch, BF16 combine शर्तों पर मापा गया, जहाँ V2 ने V1 के मुकाबले अधिकतम 1.3x peak performance और अधिकतम 4x SM reduction दर्ज की
  • सभी kernels हल्के JIT modules के रूप में runtime पर compile होते हैं, इसलिए installation के समय CUDA compilation की ज़रूरत नहीं होती, और V2 high-throughput·low-latency API को एक ही ElasticBuffer interface में एकीकृत करता है
  • Hopper SM90 GPU, Python 3.8+, CUDA 12.3+, PyTorch 2.10+, NCCL 2.30.4+, NVLink, और node के बीच RDMA network की आवश्यकता है, जबकि Engram·PP·CP experimental features हैं

DeepEP क्या प्रदान करता है

  • DeepEP(DeepEveryParallel) आधुनिक machine learning training और inference के लिए एक high-performance communication लाइब्रेरी है
  • इसका मौजूदा मुख्य फ़ोकस Expert Parallelism (EP) है, और यह MoE dispatch और combine के लिए high-throughput·low-latency all-to-all GPU kernels प्रदान करता है
  • यह FP8 सहित low-precision communication को support करता है
  • इसमें pipeline parallelism (PP), context parallelism (CP), और remote memory access (Engram) के लिए experimental primitives भी शामिल हैं
  • सभी kernels हल्के JIT(Just-In-Time) modules के रूप में runtime पर compile होते हैं, इसलिए installation प्रक्रिया में CUDA compilation की आवश्यकता नहीं होती
  • हल्के design के बावजूद, इसका लक्ष्य कई configurations में hardware bandwidth limit के बराबर या उससे अधिक performance हासिल करना है

V2 रिलीज़ के मुख्य बदलाव

  • V2, Expert Parallelism का पूरी तरह refactor किया गया release है
    • इसे V1 की तुलना में कई गुना कम SM resources के साथ अत्यधिक performance देने के लिए design किया गया है
    • यह बड़े scale-up और scale-out domains को support करता है
    • backend को NVSHMEM से हल्के NCCL Gin backend में बदल दिया गया है
  • नए features इस प्रकार हैं
    • पूर्ण JIT compilation
    • header-only·lightweight NCCL Gin backend
    • मौजूदा NCCL communicator का reuse संभव
    • EPv2 में high-throughput API और low-latency API को एक ही ElasticBuffer interface में एकीकृत किया गया
    • नया GEMM layout प्रदान किया गया
    • अधिकतम EP2048 तक बड़े scale-up और scale-out domains का support
    • SM और QP की संख्या analytically calculate की जाती है, इसलिए auto-tuning की अब आवश्यकता नहीं
    • hybrid mode और direct mode दोनों का support जारी
    • V3-जैसी legacy training में SM usage को 24 से घटाकर 4~6 करते हुए समान या बेहतर performance बनाए रखी गई
    • RDMA-आधारित 0 SM Engram
    • RDMA-आधारित 0 SM PP
    • Copy Engine-आधारित 0 SM CP

सीमाएँ और जारी features

  • V2 में V1 की तुलना में buffer size consumption अधिक है
  • 0 SM RDMA low-latency EP अब support नहीं किया जाता
  • Engram, PP, CP experimental features हैं
  • जिन features पर काम जारी है, वे इस प्रकार हैं
    • GPU और CPU physical memory की mixed mapping वाला सतत virtual address space, यानी Elastic GPU & CPU buffers
      • इसका उद्देश्य पूरी तरह automatic·transparent Engram या imbalanced EP को संभव बनाना है
    • EP replay का उपयोग करके load imbalance को संभालने और intermediate buffer size घटाने का काम
    • DP और TP के लिए all-gather update और reduce-scatter implementation
  • NVSHMEM-आधारित V1 documentation docs/legacy.md पर उपलब्ध है

performance measurement के परिणाम

  • V3 configuration के अनुरूप निम्न शर्तों पर test किया गया
    • प्रति batch 8K tokens
    • hidden dimension 7168
    • top 8 experts
    • FP8 dispatch
    • BF16 combine
  • मुख्य परिणाम इस प्रकार हैं
    • SM90, CX7, EP 8 x 2: dispatch 90 GB/s RDMA, combine 81 GB/s RDMA, 12 SM
    • SM90, CX7, EP 8 x 4: dispatch 61 GB/s RDMA, combine 61 GB/s RDMA, 6 SM
    • SM100, CX7, EP 8 x 2: dispatch 90 GB/s RDMA, combine 91 GB/s RDMA, 12 SM
    • SM100, EP 8: dispatch 726 GB/s NVLink, combine 740 GB/s NVLink, 64 SM
    • SM100, EP 8: dispatch 643 GB/s NVLink, combine 675 GB/s NVLink, 24 SM
  • मापे गए values logical bandwidth हैं; उदाहरण के लिए EP 8 x 2 के 90 GB/s में local rank traffic शामिल है
  • V2 ने V1 की तुलना में अधिकतम 1.3x peak performance हासिल की और अधिकतम 4x SM count कम किया
  • बड़े EP configurations के परिणाम फिलहाल छोड़े गए हैं, और users को स्वयं benchmark करने की सलाह दी गई है
  • आंतरिक अनुभव के आधार पर, kernel के scale बढ़ने पर भी hardware bandwidth saturation जारी रहने की उम्मीद है
  • V1 performance data docs/legacy.md में है

installation और requirements

  • requirements इस प्रकार हैं
    • Hopper SM90 GPU या SM90 PTX ISA support करने वाली architecture
    • Python 3.8 या उससे ऊपर
    • SM90 GPU के लिए CUDA 12.3 या उससे ऊपर
    • PyTorch 2.10 या उससे ऊपर
    • NCCL 2.30.4 या उससे ऊपर
    • node के भीतर communication के लिए NVLink
    • node के बीच communication के लिए RDMA network
  • NCCL के लिए pip installation की सिफारिश की जाती है ताकि Python environment में DeepEP इसे auto-detect कर सके
pip install "nvidia-nccl-cu13>=2.30.4" --no-deps
  • legacy method support के लिए यह NVSHMEM पर भी निर्भर करता है, और installation निर्देश NVSHMEM Installation Guide में दिए गए हैं
  • development के दौरान build और test चलाने के उदाहरण इस प्रकार हैं
python setup.py build
ln -s build/lib.linux-x86_64-cpython-38/deep_ep_cpp.cpython-38-x86_64-linux-gnu.so

python tests/elastic/test_ep.py
python tests/elastic/test_agrs.py
python tests/elastic/test_engram.py
python tests/elastic/test_pp.py
  • installation निम्न command से किया जाता है
python setup.py install
  • installation के बाद Python project में deep_ep को import करके उपयोग किया जा सकता है

ElasticBuffer केंद्रित interface

  • V2 में सभी EP operations को एक ही ElasticBuffer interface के तहत एकीकृत किया गया है
    • high-throughput और low-latency API को एक ही interface में संभाला जाता है
    • buffer को MoE settings सीधे specify करके initialize किया जा सकता है
    • optimal SM count और QP count analytically calculate किए जाते हैं
  • buffer initialization के उदाहरण में ElasticBuffer.get_buffer_size_hint() से आवश्यक size calculate किया जाता है और यह जाँचा जाता है कि मौजूदा buffer का reuse संभव है या नहीं
  • नया buffer बनाते समय num_max_tokens_per_rank, hidden, num_topk, use_fp8_dispatch आदि specify किए जाते हैं
  • _buffer.get_theoretical_num_sms(num_experts, num_topk) से communication kernel के लिए theoretical SM count प्राप्त किया जा सकता है
  • dispatch और combine calls में num_sms को सीधे specify करके calculated value को override किया जा सकता है

training·prefill·decoding usage patterns

  • training या inference prefill में MoE dispatch tokens को सभी ranks के संबंधित expert तक route करता है
    • यह BF16 और FP8 inputs को support करता है
    • handle में बाद की combine call के लिए आवश्यक routing metadata होता है
    • handle.num_recv_tokens_per_expert_list GEMM के लिए expert-वार token count देता है
  • MoE dispatch का backward pass वास्तव में combine द्वारा संभाला जाता है
  • MoE combine expert outputs को मूल rank पर reduce करता है
  • MoE combine का backward pass वास्तव में dispatch द्वारा संभाला जाता है
  • communication और computation की overlap को EventOverlap interface द्वारा manage किया जाता है
    • communication चलते समय independent computation किया जा सकता है
    • परिणाम उपयोग करने से पहले event.current_stream_wait() से compute stream को synchronize करना होता है
  • inference decoding में भी वही ElasticBuffer उपयोग होता है
    • जब gating decision नहीं बदलता, तो cached_handle से routing metadata का reuse किया जा सकता है
    • यह pattern layout recalculation और CPU synchronization से बचाता है

environment variables और build-time fixed values

  • सामान्य settings
    • EP_BUFFER_DEBUG: buffer initialization, SM approximation, backend debugging जानकारी output करता है
    • EP_SUPPRESS_NCCL_CHECK: NCCL version mismatch check को suppress करता है
    • EP_AVOID_RECORD_STREAM: output tensor के record_stream से बचता है
    • EP_NUM_TOPK_IDX_BITS: top-k index encoding bit count को override करता है
  • networking settings
    • EP_NIC_NAME: NIC properties query करने के लिए default NIC name, default mlx5_0
    • EP_OVERRIDE_RDMA_SL: RDMA service level index को override करता है
    • EP_DISABLE_GIN: NCCL Gin backend को disable करता है
  • JIT settings
    • EP_JIT_CACHE_DIR: compiled kernel cache directory, default $HOME/.deep_ep
    • EP_JIT_NVCC_COMPILER: NVCC compiler path
    • EP_JIT_CPP_STANDARD: C++ standard version, default 20
    • EP_JIT_DUMP_PTX, EP_JIT_DUMP_SASS, EP_JIT_DUMP_ASM: PTX·SASS output dump संबंधी settings
  • कुछ environment variables persistent रूप में काम करते हैं
    • build के समय capture होकर installed package की default values में शामिल हो जाते हैं
    • import के समय अगर current environment variable से override न किया जाए, तो ये default values अपने-आप लागू हो जाती हैं
    • संबंधित variables हैं EP_JIT_CACHE_DIR, EP_JIT_PRINT_COMPILER_COMMAND, EP_NUM_TOPK_IDX_BITS, EP_NCCL_ROOT_DIR
  • अतिरिक्त विवरण test_ep.py या Python documentation में देखें

network configuration recommendations

  • DeepEP का InfiniBand network पर पूरी तरह test किया गया है
  • सैद्धांतिक रूप से यह RDMA over Converged Ethernet यानी RoCE के साथ भी compatible है
  • traffic isolation
    • InfiniBand में Virtual Lanes के माध्यम से support किया जाता है
    • expert-parallel workload और अन्य workload को अलग virtual lanes में विभाजित करने की सिफारिश की जाती है
    • V2 में sl_idx argument या EP_OVERRIDE_RDMA_SL environment variable से virtual lane allocation नियंत्रित किया जा सकता है
  • Adaptive routing
    • यह InfiniBand switches की advanced routing capability है जो traffic को कई paths में समान रूप से बाँटती है
    • अतिरिक्त latency होने पर भी सभी network load conditions में इसे enable करने की सिफारिश की जाती है
  • Congestion control
    • अधिकतम bandwidth के लिए हानिकारक होने के कारण इसे disable किया जाना चाहिए
    • यदि congestion अपरिहार्य हो, तो संबंधित workload को low-priority virtual lane में assign करने की सिफारिश की जाती है
  • PCI atomic mode
    • यदि hardware support करता हो, तो RDMA atomic operation performance सुधारने के लिए NIC के PCI_ATOMIC_MODE को सेट करने की सिफारिश की जाती है
sudo mlxconfig -y -d mlx5_$i set PCI_ATOMIC_MODE=4

experimental branches और community forks

  • experimental branches
    • Zero-copy: PyTorch tensors और communication buffer के बीच copy हटाकर सामान्य kernels की SM usage को काफी कम करता है
    • Eager: low-latency protocol का उपयोग करके RDMA atomic OP द्वारा जुड़ने वाली extra RTT latency को हटाता है
    • Hybrid-EP: TMA instructions का उपयोग करने वाला नया backend implementation, न्यूनतम SM usage, बड़े NVLink domains का support, single-batch fine-grained communication·computation overlap, PCIe kernels, और NVFP4 support शामिल करता है
    • AntGroup-Opt: AntGroup Network Platform Department द्वारा लिखी गई optimization series
    • Mori-EP: MORI backend पर आधारित ROCm/AMD GPU low-latency mode support
    • nvDev: Compute Fabric Transport जैसे नवीनतम CUDA features शामिल करने वाली V2-आधारित branch
  • community forks
    • uccl/uccl-ep: Nvidia·AMD जैसे heterogeneous GPU और EFA·Broadcom·CX7 जैसे NIC पर DeepEP चलाने का support
    • Infrawaves/DeepEP_ibrc_dual-ports_multiQP: IBRC transport में multi-QP solution और dual-port NIC support जोड़ता है
    • antgroup/DeepXTrace: slow rank को प्रभावी और सटीक ढंग से खोजने वाला diagnostic analyzer
    • ROCm/mori: Wide EP, KVCache transfer, Collectives आदि performance-sensitive AI workloads के लिए AMD की अगली पीढ़ी की communication लाइब्रेरी

license और citation

  • DeepEP V2, NCCL Gin backend पर निर्मित है
  • repository code MIT License के तहत जारी किया गया है
  • citation entry DeepEP: an efficient expert-parallel communication library है, और वर्ष 2025 दिया गया है

संबंधित पढ़ाई

1 टिप्पणियां

 
GN⁺ 2025-02-26
Hacker News टिप्पणियाँ

  • डॉक्यूमेंट न किया गया PTX इंस्ट्रक्शन ld.global.nc.L1::no_allocate.L2::256B ढूंढकर इस्तेमाल किया गया, जिससे बेहद उच्च प्रदर्शन मिला
    यह इंस्ट्रक्शन non-coherent read-only PTX modifier .nc के साथ volatile GPU memory को एक्सेस करता है, इसलिए undefined behavior पैदा कर सकता है
    लेकिन Hopper architecture पर .L1::no_allocate के साथ इस्तेमाल करने पर correctness बनी रहती है, ऐसा टेस्ट किया गया है, और प्रदर्शन काफी बेहतर बताया गया है

    • व्यवहारिक रूप से क्या NVIDIA बाद में किसी नए architecture में इस डॉक्यूमेंटेशन से बाहर वाले इंस्ट्रक्शन का behavior थोड़ा बदलकर, चाहे जानबूझकर या अनजाने में, पूरा खेल पलट सकती है?

  • यह किसी बच्चे के candy store में घुसने जैसा लगता है
    सिर्फ पेपर देखकर ठीक से reverse engineer करने में बहुत समय लगने वाले कई tricks हैं, और उम्मीद है कि इस हफ्ते की रिलीज़ें MoE को डिफ़ॉल्ट अकादमिक मॉडल की तरह इस्तेमाल करने वाले पुनर्जागरण की शुरुआत करेंगी

    • इस नज़रिए से देखें तो यह समझना मुश्किल है कि असल state-of-the-art मॉडल प्रैक्टिस और अकादमिक मॉडलों के बीच क्या चल रहा है
      पहले वाले तो GPT-4 से ही लगभग सब MoE हैं, लेकिन खुले मॉडल DeepSeek V3 और Mixtral को छोड़कर अक्सर हमेशा dense models ही होते हैं

  • इस टीम को पसंद किए बिना रहना मुश्किल है
    ये सभी के लिए open source की सीमाओं को आगे बढ़ा रहे हैं

    • Open AI™ की तरह space डालकर लिखा गया वाला
    • असल में यह open source नहीं है
      सचमुच का open source मॉडल देखना हो तो AI2 का OLMo 2 देखें: https://allenai.org/blog/olmo2
      मॉडल को reproduce करने के लिए ज़रूरी हर चीज़, data तक, वास्तव में साझा की गई है
      ऊपर के लिंक में भी कहा गया है, “पूर्ण open science के लिए सिर्फ open weights से अधिक की ज़रूरत होती है, इसलिए हमें wider language modeling community के साथ weights, data, code, recipes, intermediate checkpoints, और instruction-tuned models सहित नए OLMo updates साझा करते हुए खुशी हो रही है”

  • Zuckerberg को यह दावा करना बंद करना चाहिए कि Meta AI को open source के रूप में जारी कर रही है
    TV ads तक चल रहे हैं, लेकिन वास्तव में सिर्फ weights जारी किए जाते हैं और code नहीं है
    असली open source AI सिर्फ DeepSeek है

    • सख्ती से कहें तो DeepSeek भी OLMo या Open Euro जितना open source नहीं है
      क्योंकि उसने data जारी नहीं किया है
    • DeepSeek स्पष्ट रूप से असली open source नहीं है
      open source होने के लिए OSI द्वारा सूचीबद्ध जैसे वास्तविक open source license का इस्तेमाल होना चाहिए, और pre/post-training code, tuning से जुड़ा code, evaluation code, safety/censorship से जुड़ी हर चीज़, और शायद पूरा training data भी साझा करना होगा
      नहीं तो weights को reproduce नहीं किया जा सकता, और weights साझा करना compiled program साझा करने जैसा है
      मेरी जानकारी में प्रतिस्पर्धी और सचमुच open source मॉडल सिर्फ AI2 का OLMo 2 है: https://allenai.org/blog/olmo2
      हाल ही में on-device inference करने वाला एक app भी जारी किया गया है, और वह भी open source है: https://allenai.org/blog/olmoe-app
      Tülu 3 नाम का एक और मॉडल भी है, जिसके बारे में कहा गया है कि वह DeepSeek V3 से बेहतर प्रदर्शन करता है: https://allenai.org/blog/tulu-3-405B
    • Meta PyTorch को 10 साल से अधिक समय से निखार रही है
      LLM को train करने के लिए ज़रूरी लगभग सब कुछ, नई तकनीकों सहित, उसमें मौजूद है
      फिर और क्या चाहिए? क्या Meta infrastructure के लिए विशेष code वाले हिस्से?
    • क्या PyTorch भी इसमें शामिल है?
    • खुले weights = binary blob
      यानी फिर से FREEWARE / SHAREWARE मॉडल पर लौटना
      “open” weights के लिए ऐसे ही शब्द इस्तेमाल किए जाने चाहिए

  • यह कुशल और अनुकूलित all-to-all communication, NVLink और RDMA के माध्यम से intra-node और inter-node support, training और inference prefilling के लिए high-throughput kernels, inference decoding के लिए low-latency kernels, native FP8 dispatch support, और compute-communication overlap के लिए flexible GPU resource control प्रदान करता है
    X: https://x.com/deepseek_ai/status/1894211757604049133

  • DeepSeek के काम की प्रेरणा गलत भी हो सकती है
    उदाहरण के लिए, अमेरिका की AI first-mover बढ़त को शून्य करने की राज्य-प्रायोजित कोशिश जैसी, लेकिन दुनिया भर के सभी लोगों के लिए इसका net effect बस शानदार है
    सबसे खराब स्थिति में भी, यानी अगर यह गलत कारणों से किया जा रहा हो, तब भी DeepSeek को धन्यवाद, क्योंकि वह वह काम वास्तव में कर रही है जिसके बारे में OpenAI ने सालों तक दुनिया से झूठ बोला कि वह कर रही है

    • अंतरराष्ट्रीय संबंधों के क्षेत्र में सही और गलत इतना ज़ोर से लागू नहीं होता
      क्या इसे open source के रूप में जारी करना advanced Nvidia GPU export ban से भी ज़्यादा “गलत” है?
      DeepSeek का open source release शायद बस चीनी कम्युनिस्ट पार्टी की सहमति के तहत, चीनी कम्युनिस्ट पार्टी और व्यापक open source AI community दोनों के लिए एक साथ फ़ायदेमंद चीज़ रहा हो, और इसे किसी सिद्धांतवादी रुख की तरह नहीं लेना चाहिए
      दूसरे देशों की competitive advantage को खत्म करने के तरीके ढूँढना, छोटे-बड़े सभी governments की मुख्य गतिविधियों में से एक है

  • MIT license के तहत यह वास्तविक Open AI™ कंपनी की ओर से आया open source release का दूसरा राउंड है
    एक बार फिर DeepSeek, खुद को “Open” कहने वाली 157 बिलियन डॉलर की कंपनी से अधिक खुली है
    Meta की Llama के बारे में लगभग कोई बात नहीं कर रहा, लेकिन सभी को उम्मीद करनी चाहिए कि Llama 4 reasoning capabilities के साथ आएगा
    लक्ष्य यह है कि शून्य की ओर भागती इस race के बीच पिसकर खत्म न हो जाएँ

  • जब अमेरिका यह जाँचने के लिए Singapore की GPU रसीदें खंगाल रहा है कि DeepSeek ने सिर्फ H800 इस्तेमाल किए या नहीं, तब बाकी दुनिया इन optimizations को पूरे H100 पर चला सकती है?
    जबकि अमेरिकी प्रतिबंधों के कारण H100 हासिल करना या उन तक पहुँचना मुश्किल था, और अमेरिका का अहंकार यह मानता रहा कि उसके आदेश पूरी दुनिया पर लागू होते हैं?
    क्या मैं इसे सही समझ रहा हूँ?

  • क्या वह PTX जिसके लिए सब इंतज़ार कर रहे थे, इस बार शामिल है?

    • हाँ, csrc/kernels directory में कुछ हिस्सा है
      asm खोजेंगे तो जहाँ इस्तेमाल हुआ है वह मिल जाएगा
    • बाकी लोगों के लिए यह समझाना ज़रूरी है कि वह PTX, जिसका सब इंतज़ार कर रहे थे, इतना महत्वपूर्ण क्यों है

  • यह मुझे 80~90 के दशक की याद दिलाता है, जब लोग CPU performance निचोड़ने के लिए assembly hack करते थे या undocumented instructions ढूँढते थे
    कभी न कभी compiler इतने optimize हो जाएँगे, या GPU इतने शक्तिशाली हो जाएँगे, कि आज के CPU की तरह ये tricks अब इतना बड़ा फ़र्क नहीं पैदा करेंगी