DeepSeek, MoE eğitimi ve çıkarımı için açık kaynak DeepEP kütüphanesini yayımladı

DeepEP

DeepEP, Mixture-of-Experts (MoE) ve expert parallelism (EP) için bir iletişim kütüphanesidir. Yüksek hızlı, düşük gecikmeli all-to-all GPU kernel'ları sunar ve bunlar MoE dispatch ve combine olarak bilinir. Ayrıca FP8 dahil düşük hassasiyetli işlemleri destekler. DeepSeek-V3 makalesinde önerilen grup sınırlı gating algoritmasına uygun olarak, asimetrik alan bant genişliği aktarımını optimize eden kernel'lar sunar ve veriyi NVLink alanından RDMA alanına aktarır. Bu kernel'lar yüksek throughput sağlayarak eğitim ve çıkarım prefill işleri için uygundur. Ayrıca SM (Streaming Multiprocessors) sayısı kontrolünü destekler. Gecikmeye duyarlı çıkarım decoding'i için DeepEP, gecikmeyi en aza indirmek üzere saf RDMA kullanan düşük gecikmeli kernel'lar içerir. Bu kütüphane, SM kaynaklarını işgal etmeyen hook tabanlı bir iletişim-hesaplama örtüşmesi yöntemi sunar.

Performans

NVLink ve RDMA aktarımı kullanan genel kernel'lar

• H800 üzerinde, yaklaşık 160 GB/s azami NVLink bant genişliği ve CX7 InfiniBand 400 Gb/s RDMA ağ kartı (~50 GB/s azami bant genişliği) ile genel kernel'lar test edildi.
• DeepSeek-V3/R1 ön eğitim ayarları izlendi (batch başına 4096 token, 7168 hidden, top-4 grup, top-8 uzman, FP8 dispatch ve BF16 combine).

Saf RDMA kullanan düşük gecikmeli kernel'lar

• H800 üzerinde, CX7 InfiniBand 400 Gb/s RDMA ağ kartı (~50 GB/s azami bant genişliği) ile düşük gecikmeli kernel'lar test edildi.
• Tipik DeepSeek-V3/R1 üretim ayarları izlendi (batch başına 128 token, 7168 hidden, top-8 uzman, FP8 dispatch ve BF16 combine).

Hızlı başlangıç

Gereksinimler

• Hopper GPU (ileride daha fazla mimari veya cihaz desteklenebilir)
• Python 3.8+
• CUDA 12.3+
• PyTorch 2.1+
• Düğüm içi iletişim için NVLink
• Düğümler arası iletişim için RDMA ağı

NVSHMEM bağımlılığını indirme ve kurma

DeepEP, değiştirilmiş bir NVSHMEM sürümüne bağlıdır. Kurulum kılavuzuna bakarak yüklemeniz gerekir.

Ağ yapılandırması

DeepEP, InfiniBand ağında tamamen test edilmiştir ve teorik olarak RDMA over Converged Ethernet (RoCE) ile de uyumludur.

Trafik izolasyonu

InfiniBand, Virtual Lanes (VL) aracılığıyla trafik izolasyonunu destekler. Farklı trafik türleri arasında girişimi önlemek için, iş yüklerini şu şekilde sanal lane'lere ayırmanız önerilir:

• Genel kernel kullanan işler
• Düşük gecikmeli kernel kullanan işler
• Diğer işler

DeepEP'te, sanal lane atamasını kontrol etmek için NVSHMEM_IB_SL ortam değişkeni ayarlanabilir.

Uyarlamalı yönlendirme

Uyarlamalı yönlendirme, InfiniBand switch'lerinde sunulan gelişmiş bir yönlendirme özelliğidir ve trafiğin birden fazla yol arasında dengeli dağıtılmasını sağlar. Şu anda düşük gecikmeli kernel'lar uyarlamalı yönlendirmeyi desteklerken genel kernel'lar desteklememektedir (yakında desteklenebilir). Normal düğümler arası kernel'larda uyarlamalı yönlendirmeyi etkinleştirmek deadlock veya veri bozulması sorunlarına yol açabilir. Düşük gecikmeli kernel'larda ise uyarlamalı yönlendirme etkinleştirildiğinde, yönlendirme çakışmalarından kaynaklanan ağ tıkanıklığı tamamen ortadan kaldırılabilir; ancak ek gecikme oluşur. En iyi performans için şu yapılandırma önerilir:

• Ağ yükünün yüksek olduğu ortamlarda uyarlamalı yönlendirmeyi etkinleştirin
• Ağ yükünün düşük olduğu ortamlarda statik yönlendirme kullanın

Tıkanıklık kontrolü

Üretim ortamında anlamlı bir tıkanıklık gözlemlenmediği için tıkanıklık kontrolü devre dışı bırakılmıştır.

Arayüz ve örnekler

Model eğitimi veya çıkarım prefill aşamasında örnek kullanım

Genel kernel'lar, model eğitimi veya çıkarımın prefill aşamasında (geri yayılım kısmı olmadan) kullanılabilir.

Çıkarım decoding aşamasında örnek kullanım

Düşük gecikmeli kernel'lar çıkarımın decoding aşamasında kullanılabilir.

Dikkat edilmesi gerekenler

• En yüksek performans için, belgelenmemiş PTX komutu ld.global.nc.L1::no_allocate.L2::256B keşfedilip kullanıldı. Bu komut, tutarsız salt okunur PTX değiştiricisi kullanarak volatile GPU belleğine erişen tanımsız davranışa neden olur. Ancak Hopper mimarisinde .L1::no_allocate ile yapılan testlerde performansın çok daha iyi olduğu görüldü. Eğer kernel'lar başka platformlarda çalışmazsa, bunu devre dışı bırakmak için setup.py dosyasına DISABLE_AGGRESSIVE_PTX_INSTRS=1 ekleyebilir veya bir issue açabilirsiniz.
• Kümede daha iyi performans için tüm testlerin çalıştırılması ve en uygun otomatik ayar yapılandırmasının kullanılması önerilir. Varsayılan yapılandırma DeepSeek'in dahili kümesinde optimize edilmiştir.

Lisans

Bu kod deposu MIT lisansı altında yayımlanmıştır ve NVSHMEM'e referans veren kod (csrc/kernels/ibgda_device.cuh ve third-party/nvshmem.patch dahil) NVSHMEM SLA'ya tabidir.

[GN⁺]

Hacker News yorumları

• Uç düzey performans için belgelenmemiş PTX komutlarını keşfedip kullanmışlar. Bu komutlar, tutarsız salt okunur PTX değiştiricileri kullanarak uçucu GPU belleğine eriştiğinden tanımsız davranışa yol açabilir. Ancak Hopper mimarisinde .L1::no_allocate ile doğruluğun test edildiği ve performansın çok daha iyi olduğu belirtiliyor
• Zuckerberg, Meta'nın yapay zekayı açık kaynak yaptığını iddia etmeyi bırakmalı. Kod değil, yalnızca ağırlıkları yayımlıyorlar. Gerçek açık kaynak yapay zeka yalnızca DeepSeek'te var
• Kendimi bir şeker dükkânındaki çocuk gibi hissediyorum. Bu numaraların bazılarını yalnızca makalelere dayanarak doğru şekilde tersine mühendislik etmek çok uzun sürecek. Umarım bu haftaki duyurular, varsayılan akademik model olarak MoE kullanan bir rönesans başlatır
• Bu insanları sevmemek elde değil. Hepimiz için açık kaynağın sınırlarını gerçekten zorluyorlar. Paylaştığınız için teşekkürler
  • Verimli ve optimize edilmiş all-to-all iletişim
  • NVLink ve RDMA üzerinden düğüm içi ve düğümler arası destek
  • Eğitim ve çıkarım prefill için yüksek verimli çekirdekler
  • Çıkarım decoding için düşük gecikmeli çekirdekler
  • Yerel FP8 dispatch desteği
  • Hesaplama-iletişim örtüşmesi için esnek GPU kaynak kontrolü
• DeepSeek'in çalışmalarının arkasındaki motivasyon yanlış olabilir (ör. yapay zekada ABD'nin liderlik avantajını ortadan kaldırmaya yönelik devlet destekli bir girişim). Ancak dünya genelinde sonuçlar tek kelimeyle muhteşem
  • En kötü durumda bile (bu işi yanlış nedenlerle yapıyor olsalar bile) DeepSeek'e teşekkürler. OpenAI'nin yıllardır tüm dünyaya yalan söyleyerek yaptığını, onlar gerçekten yapıyor
  • Gerçekten inanılmaz
• Herkesin beklediği PTX'in bu kez dahil edilip edilmediğini merak ediyorum
• Teknik raporda bahsedilen PTX komutlarının buradaki koda bağlanması gerekiyor
• ABD, DeepSeek'in yalnızca H800 kullanıp kullanmadığını doğrulamak için Singapur'daki GPU faturalarının peşine düşerken, dünyanın geri kalanı bu optimizasyonları tam H100'lerde çalıştırabilir
  • ABD yaptırımları ve kendi emirlerinin tüm dünyayı kapsadığına inanma kibri yüzünden, H100 edinmenin ya da erişmenin zor olduğunu sanıyor gibi yapıp yapmadıklarını merak ediyorum
• Gerçek "Open AI™" şirketinden gelen ikinci açık kaynak sürüm bu ve MIT lisansı altında
  • DeepSeek, $157B+ iddiasındaki şirketten daha açık
  • Neredeyse hiç kimse Meta'nın Llama'sından bahsetmiyor ve herkesin Llama 4'ün nedenleriyle birlikte çıkmasını beklemesi gerekir
  • Amaç, sıfıra yarışta arada sıkışıp kalmamak