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

 (github.com/deepseek-ai)
1 puan yazan GN⁺ 2025-02-26 | 1 yorum | WhatsApp'ta paylaş
  • DeepEP, modern ML eğitimi ve çıkarımında uzman paralelleştirmeye (EP) odaklanan yüksek performanslı bir iletişim kütüphanesidir; MoE dispatch/combine için all-to-all GPU kernel’ları ve FP8 gibi düşük hassasiyet desteği sunar
  • V2 sürümü, EP’yi baştan sona yeniden düzenleyerek V1’e kıyasla çok daha az SM kaynağıyla eşdeğer veya daha iyi performans sağlar; backend’i NVSHMEM’den daha hafif NCCL Gin’e taşır
  • V3 yapılandırmasına göre yapılan testlerde batch başına 8K token, hidden 7168, top 8 experts, FP8 dispatch, BF16 combine koşulları ölçüldü; V2, V1’e kıyasla en fazla 1,3 kat tepe performans ve en fazla 4 kat SM tasarrufu kaydetti
  • Tüm kernel’lar hafif JIT modülleri olarak çalışma zamanında derlenir; kurulum sırasında CUDA derlemesi gerekmez ve V2, yüksek aktarım hızlı ve düşük gecikmeli API’leri tek bir ElasticBuffer arayüzünde birleştirir
  • Hopper SM90 GPU, Python 3.8+, CUDA 12.3+, PyTorch 2.10+, NCCL 2.30.4+, NVLink ve düğümler arası RDMA ağı gerekir; Engram, PP ve CP deneysel özelliklerdir

DeepEP’nin sunduğu kapsam

  • DeepEP(DeepEveryParallel), modern makine öğrenimi eğitimi ve çıkarımı için yüksek performanslı bir iletişim kütüphanesidir
  • Şu anki ana işlevi uzman paralelliği (Expert Parallelism, EP) olup MoE dispatch ve combine için yüksek aktarım hızlı, düşük gecikmeli all-to-all GPU kernel’ları sağlar
  • FP8 dahil düşük hassasiyetli iletişimi destekler
  • Pipeline paralelliği (PP), context paralelliği (CP) ve uzak bellek erişimi (Engram) için deneysel primitive’ler de içerir
  • Tüm kernel’lar hafif JIT(Just-In-Time) modülleri olarak çalışma zamanında derlenir; kurulum sürecinde CUDA derlemesi gerekmez
  • Hafif tasarımına rağmen çeşitli yapılandırmalarda donanım bant genişliği sınırına yaklaşan veya onu aşan performansı hedefler

V2 sürümündeki temel değişiklikler

  • V2, Expert Parallelism’i tamamen yeniden düzenleyen bir sürümdür
    • V1’e göre birkaç kat daha az SM kaynağıyla uç performans elde edecek şekilde tasarlanmıştır
    • Daha büyük scale-up ve scale-out alanlarını destekler
    • Backend, NVSHMEM’den daha hafif NCCL Gin backend’ine taşınmıştır
  • Yeni özellikler şunlardır
    • Tam JIT derleme
    • Header-only ve hafif NCCL Gin backend
    • Mevcut NCCL communicator’ların yeniden kullanılabilmesi
    • EPv2’de yüksek aktarım hızlı API ile düşük gecikmeli API’nin tek bir ElasticBuffer arayüzünde birleştirilmesi
    • Yeni GEMM yerleşimi sağlanması
    • En fazla EP2048’e kadar daha büyük scale-up ve scale-out alanlarının desteklenmesi
    • SM ve QP sayılarının analitik olarak hesaplanması sayesinde artık otomatik ayara ihtiyaç duyulmaması
    • Hem hibrit modun hem de doğrudan modun desteklenmeye devam etmesi
    • V3 benzeri eski eğitimde SM kullanımını 24’ten 4~6’ya düşürürken eşdeğer veya daha iyi performansın korunması
    • RDMA tabanlı 0 SM Engram
    • RDMA tabanlı 0 SM PP
    • Copy Engine tabanlı 0 SM CP

Kısıtlar ve geliştirilmekte olan özellikler

  • V2, V1’e göre daha fazla buffer boyutu tüketir
  • 0 SM RDMA düşük gecikmeli EP artık desteklenmez
  • Engram, PP, CP deneysel özelliklerdir
  • Geliştirilmekte olan özellikler şunlardır
    • GPU ve CPU fiziksel belleğini karma eşleyen, kesintisiz bir sanal adres alanı olan Elastic GPU & CPU buffers
      • Tam otomatik ve şeffaf Engram ya da dengesiz EP’yi mümkün kılmaya yöneliktir
    • Yük dengesizliklerini yönetmek ve ara buffer boyutunu azaltmak için EP replay’den yararlanma çalışmaları
    • DP ve TP için all-gather güncellemeleri ve reduce-scatter uygulaması
  • NVSHMEM tabanlı V1 belgeleri docs/legacy.md adresinde sunulur

Performans ölçüm sonuçları

  • V3 yapılandırmasına uygun olarak şu koşullarda test edilmiştir
    • batch başına 8K token
    • hidden dimension 7168
    • top 8 experts
    • FP8 dispatch
    • BF16 combine
  • Başlıca sonuçlar şunlardır
    • 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
  • Ölçülen değerler mantıksal bant genişliğidir; örneğin EP 8 x 2 için 90 GB/s, local rank trafiğini içerir
  • V2, V1’e kıyasla en fazla 1,3 kat tepe performansa ulaşır ve en fazla 4 kat SM sayısı tasarrufu sağlar
  • Daha büyük EP yapılandırmalarının sonuçları şu anda atlanmıştır; kullanıcıların kendi benchmark’larını yapması önerilir
  • Dahili deneyime göre kernel’ların ölçek büyüdüğünde de donanım bant genişliğini doyurmaya devam etmesi beklenir
  • V1 performans verileri docs/legacy.md içindedir

Kurulum ve gereksinimler

  • Gereksinimler şunlardır
    • Hopper SM90 GPU veya SM90 PTX ISA destekli mimari
    • Python 3.8 veya üzeri
    • SM90 GPU için CUDA 12.3 veya üzeri
    • PyTorch 2.10 veya üzeri
    • NCCL 2.30.4 veya üzeri
    • Düğüm içi iletişim için NVLink
    • Düğümler arası iletişim için RDMA ağı
  • DeepEP’nin Python ortamında NCCL’yi otomatik algılayabilmesi için NCCL’nin pip ile kurulması önerilir
pip install "nvidia-nccl-cu13>=2.30.4" --no-deps
  • Eski yöntem desteği için NVSHMEM’e de bağımlıdır; kurulum talimatları için NVSHMEM Installation Guide belgesine bakın
  • Geliştirme sırasında build ve test çalıştırma örnekleri şöyledir
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
  • Kurulum aşağıdaki komutla yapılır
python setup.py install
  • Kurulumdan sonra Python projesinde deep_ep import edilerek kullanılabilir

ElasticBuffer merkezli arayüz

  • V2’de tüm EP işlemleri tek bir ElasticBuffer arayüzü altında birleştirilmiştir
    • Yüksek aktarım hızı ve düşük gecikme API’leri aynı arayüzle ele alınır
    • Buffer, MoE ayarları doğrudan belirtilerek başlatılabilir
    • En uygun SM sayısı ve QP sayısı analitik olarak hesaplanır
  • Buffer başlatma örneğinde ElasticBuffer.get_buffer_size_hint() ile gerekli boyut hesaplanır ve mevcut buffer’ın yeniden kullanılıp kullanılamayacağı kontrol edilir
  • Yeni buffer oluşturulurken num_max_tokens_per_rank, hidden, num_topk, use_fp8_dispatch vb. belirtilir
  • _buffer.get_theoretical_num_sms(num_experts, num_topk) ile iletişim kernel’ında kullanılacak teorik SM sayısı elde edilir
  • dispatch ve combine çağrılarında num_sms doğrudan belirtilirse hesaplanan değer geçersiz kılınabilir

Eğitim, prefill ve decoding kullanım kalıpları

  • Eğitim veya çıkarım prefill aşamasında MoE dispatch, token’ları tüm rank’lerdeki ilgili expert’e yönlendirir
    • BF16 ve FP8 girdileri destekler
    • handle, daha sonraki combine çağrısı için gereken yönlendirme metaverisini içerir
    • handle.num_recv_tokens_per_expert_list, GEMM için gereken expert başına token sayısını sağlar
  • MoE dispatch’in backward pass’i pratikte combine ile işlenir
  • MoE combine, expert çıktılarını özgün rank’e reduce eder
  • MoE combine’ın backward pass’i pratikte dispatch ile işlenir
  • İletişim ve hesaplamanın örtüşmesi EventOverlap arayüzüyle yönetilir
    • İletişim sürerken bağımsız hesaplama yapılabilir
    • Sonuçlar kullanılmadan önce event.current_stream_wait() ile compute stream senkronize edilir
  • Çıkarım decoding’de de aynı ElasticBuffer kullanılır
    • gating decision değişmediğinde yönlendirme metaverisi cached_handle ile yeniden kullanılır
    • Bu kalıp, yerleşimin yeniden hesaplanmasını ve CPU senkronizasyonunu önler

Ortam değişkenleri ve build sırasında sabitlenen değerler

  • Genel ayarlar
    • EP_BUFFER_DEBUG: buffer başlatma, SM yaklaştırması ve backend debug bilgilerini yazdırır
    • EP_SUPPRESS_NCCL_CHECK: NCCL sürüm uyuşmazlığı kontrolünü bastırır
    • EP_AVOID_RECORD_STREAM: çıktı tensor’larının record_stream kullanımından kaçınır
    • EP_NUM_TOPK_IDX_BITS: top-k index kodlama bit sayısını geçersiz kılar
  • Ağ ayarları
    • EP_NIC_NAME: NIC özellik sorgusu için kullanılacak varsayılan NIC adı; varsayılan mlx5_0
    • EP_OVERRIDE_RDMA_SL: RDMA service level index’ini geçersiz kılar
    • EP_DISABLE_GIN: NCCL Gin backend’ini devre dışı bırakır
  • JIT ayarları
    • EP_JIT_CACHE_DIR: derlenmiş kernel cache dizini; varsayılan $HOME/.deep_ep
    • EP_JIT_NVCC_COMPILER: NVCC derleyici yolu
    • EP_JIT_CPP_STANDARD: C++ standart sürümü; varsayılan 20
    • EP_JIT_DUMP_PTX, EP_JIT_DUMP_SASS, EP_JIT_DUMP_ASM: PTX/SASS/ASM çıktısı dökümüyle ilgili ayarlar
  • Bazı ortam değişkenleri persistent olarak çalışır
    • Build sırasında yakalanır ve kurulum paketinin varsayılan değeri olarak dahil edilir
    • Import sırasında mevcut ortam değişkeniyle geçersiz kılınmazsa bu varsayılan otomatik uygulanır
    • Hedef değişkenler EP_JIT_CACHE_DIR, EP_JIT_PRINT_COMPILER_COMMAND, EP_NUM_TOPK_IDX_BITS, EP_NCCL_ROOT_DIR’dir
  • Ek ayrıntılar için test_ep.py veya Python belgelerine bakın

Ağ yapılandırma önerileri

  • DeepEP, InfiniBand ağlarında eksiksiz test edilmiştir
  • Teorik olarak RDMA over Converged Ethernet, yani RoCE ile de uyumludur
  • Trafik yalıtımı
    • InfiniBand’in Virtual Lanes özelliğiyle desteklenir
    • expert-parallel workload ile diğer workload’ların farklı virtual lane’lere ayrılması önerilir
    • V2’de virtual lane ataması sl_idx argümanı veya EP_OVERRIDE_RDMA_SL ortam değişkeniyle kontrol edilebilir
  • Adaptive routing
    • InfiniBand switch’lerinin trafiği birden çok yola dengeli biçimde dağıtmasını sağlayan gelişmiş bir yönlendirme özelliğidir
    • Ek gecikme oluşsa bile tüm ağ yükü koşullarında etkinleştirilmesi önerilir
  • Congestion control
    • Maksimum bant genişliğine zarar verdiği için devre dışı bırakılır
    • Tıkanıklık kaçınılmazsa ilgili workload’un düşük öncelikli bir virtual lane’e atanması önerilir
  • PCI atomic mode
    • Donanım destekliyorsa RDMA atomic operation performansını iyileştirmek için NIC’in PCI_ATOMIC_MODE ayarının yapılması önerilir
sudo mlxconfig -y -d mlx5_$i set PCI_ATOMIC_MODE=4

Deneysel branch’ler ve topluluk fork’ları

  • Deneysel branch’ler
    • Zero-copy: PyTorch tensor’ları ile iletişim buffer’ları arasındaki kopyalamayı kaldırarak genel kernel’ların SM kullanımını büyük ölçüde azaltır
    • Eager: Düşük gecikmeli protokol kullanarak RDMA atomic OP’nin eklediği extra RTT latency’yi ortadan kaldırır
    • Hybrid-EP: TMA instructions kullanan yeni backend uygulaması, minimum SM kullanımı, daha büyük NVLink alanı desteği, single-batch fine-grained iletişim/hesaplama örtüşmesi, PCIe kernel’ı ve NVFP4 desteği içerir
    • AntGroup-Opt: AntGroup Network Platform Department tarafından yazılmış optimizasyon serisidir
    • Mori-EP: MORI backend’i tabanlı ROCm/AMD GPU düşük gecikmeli mod desteği
    • nvDev: Compute Fabric Transport gibi en yeni CUDA özelliklerini içeren V2 tabanlı branch
  • Topluluk fork’ları
    • uccl/uccl-ep: DeepEP’nin Nvidia ve AMD gibi heterojen GPU’larda, EFA, Broadcom ve CX7 gibi NIC’lerde çalışmasını destekler
    • Infrawaves/DeepEP_ibrc_dual-ports_multiQP: IBRC transport’a multi-QP çözümü ve dual-port NIC desteği ekler
    • antgroup/DeepXTrace: slow rank’i verimli ve hassas biçimde bulan tanısal analiz aracı
    • ROCm/mori: Wide EP, KVCache transfer, Collectives gibi performans hassasiyetli yapay zeka workload’ları için AMD’nin yeni nesil iletişim kütüphanesi

Lisans ve atıf

  • DeepEP V2, NCCL Gin backend’i üzerine inşa edilmiştir
  • Depo kodu MIT License ile yayımlanmıştır
  • Atıf öğesi DeepEP: an efficient expert-parallel communication library olup yıl 2025 olarak belirtilir

İlgili okumalar

1 yorum

 
GN⁺ 2025-02-26
Hacker News yorumları

  • Belgelenmemiş PTX komutu ld.global.nc.L1::no_allocate.L2::256B keşfedilip kullanılarak uç düzey performans elde edilmiş
    Bu komut, tutarlı olmayan salt okunur PTX değiştiricisi .nc ile geçici GPU belleğine eriştiği için tanımsız davranışa yol açabilir
    Ancak Hopper mimarisinde .L1::no_allocate ile birlikte kullanıldığında doğruluğun korunduğu test edilmiş ve performansın çok daha iyi olduğu söyleniyor

    • Gerçekçi olarak NVIDIA daha sonra yeni bir mimaride bu belgelenmemiş komutun davranışını, kasıtlı olsun ya da olmasın, ince biçimde değiştirip her şeyi altüst edebilir mi?

  • Kendimi şekerci dükkânına girmiş çocuk gibi hissediyorum
    Sadece makalelere bakarak düzgün şekilde tersine mühendislik yapmak çok uzun sürecek pek çok numara var ve umarım bu haftaki yayınlar, MoE'nin temel akademik model olarak kullanıldığı bir rönesans başlatır

    • Bu açıdan bakınca, gerçek en ileri model pratiği ile akademik modeller arasında neler döndüğünü anlamıyorum
      İlki GPT-4'ten beri zaten tamamen MoE iken, açık modellerde DeepSeek V3 ve Mixtral dışında genelde hep yoğun modeller görüyoruz

  • Bu ekibi gerçekten sevmemek elde değil
    Herkes için açık kaynağın sınırlarını zorluyorlar

    • Open AI™ gibi ayrı yazılmış olan taraftayım
    • Aslında açık kaynak sayılmaz
      Gerçek açık kaynak bir model görmek istiyorsanız AI2'nin OLMo 2 modeline bakın: https://allenai.org/blog/olmo2
      Modeli yeniden üretmek için gereken her şeyi, verinin kendisi dahil, gerçekten paylaşıyorlar
      Bağlantıda da “Tam anlamıyla açık bilim, açık ağırlıklardan daha fazlasını gerektirir; bu yüzden ağırlıklar, veri, kod, reçete, ara checkpoint'ler ve instruction tuning modellerini içeren yeni OLMo güncellemelerini daha geniş dil modelleme topluluğuyla paylaşmaktan mutluyuz” deniyor

  • Zuckerberg, Meta'nın yapay zekayı açık kaynak olarak yayımladığını iddia etmeyi bırakmalı
    TV reklamı bile yapıyorlar ama gerçekte sadece ağırlıkları yayımlıyorlar ve kod yok
    Gerçek açık kaynak yapay zeka yalnızca DeepSeek'te var

    • Teknik olarak DeepSeek de OLMo ya da Open Euro kadar açık kaynak değil
      Çünkü veriyi yayımlamadılar
    • DeepSeek kesinlikle gerçek açık kaynak değil
      Açık kaynak olması için OSI'nin listelediği gibi gerçek bir açık kaynak lisansı kullanması ve ön/son eğitim kodu, tuning ile ilgili kodlar, değerlendirme kodu, güvenlik/sansürle ilgili her şey ve muhtemelen tüm eğitim verisini paylaşması gerekir
      Aksi halde ağırlıkları yeniden üretmek mümkün olmaz ve ağırlık paylaşmak derlenmiş bir programı paylaşmaya benzer
      Bildiğim kadarıyla rekabetçi tek gerçek açık kaynak model AI2'nin OLMo 2'si: https://allenai.org/blog/olmo2
      Yakın zamanda cihaz içi çıkarım yapan bir uygulama da yayımladılar; o da açık kaynak: https://allenai.org/blog/olmoe-app
      Tülu 3 adlı başka bir modelleri daha var ve DeepSeek V3'ten daha iyi performans verdiği söyleniyor: https://allenai.org/blog/tulu-3-405B
    • Meta, PyTorch'u 10 yılı aşkın süredir geliştirip cilalıyor
      LLM eğitmek için gereken neredeyse her şey, en son teknikler dahil, onun içinde var
      Daha ne gerekiyor? Meta altyapısına özel kod parçaları mı?
    • PyTorch da buna dahil mi?
    • Açık ağırlıklar = ikili veri yığını
      FREEWARE / SHAREWARE modeline geri dönmüş oluyoruz
      “Açık” ağırlıklar için böyle bir ifade kullanılmalı

  • Verimli ve optimize edilmiş all-to-all iletişim, NVLink ve RDMA üzerinden node içi ve node'lar arası destek, eğitim ve çıkarım prefilling için yüksek throughput çekirdekleri, çıkarım decoding için düşük gecikmeli çekirdekler, yerel FP8 dispatch desteği ve hesaplama-iletişim çakışması için esnek GPU kaynak kontrolü sunuyor
    X: https://x.com/deepseek_ai/status/1894211757604049133

  • DeepSeek'in çalışmalarının motivasyonu yanlış olabilir
    Örneğin ABD'nin yapay zekadaki ilk hareket avantajını sıfırlamaya çalışan devlet destekli bir girişim gibi olabilir; ancak dünya geneline net etkisi düpedüz harika
    En kötü ihtimalle, yani yanlış nedenlerle yapıyor olsalar bile, DeepSeek'e teşekkürler; OpenAI'nin yıllardır dünyaya yaptığını söyleyip yapmadığı şeyi gerçekten yapıyorlar

    • Uluslararası ilişkiler alanında doğru ve yanlış o kadar da geçerli değil
      Bunu açık kaynak olarak yayımlamak, gelişmiş Nvidia GPU ihracat yasaklarından daha mı “yanlış”?
      DeepSeek'in açık kaynak yayını muhtemelen Çin Komünist Partisi'nin onayıyla hem Çin Komünist Partisi'nin hem de daha geniş açık kaynak yapay zeka topluluğunun işine yarayan bir hamle sadece; bunu ilkesel bir duruş gibi almamak gerek
      Başka ülkelerin rekabet avantajını ortadan kaldırmanın yollarını aramak, büyük küçük tüm hükümetlerin başlıca faaliyetlerinden biridir

  • MIT lisansı altında, gerçek Open AI™ şirketinden gelen açık kaynak yayının 2. turu bu
    DeepSeek bir kez daha, kendisine “Open” diyen 157 milyar dolarlık şirketten daha açık olduğunu gösteriyor
    Meta'nın Llama'sı hakkında neredeyse kimse konuşmuyor ama herkes Llama 4'ün çıkarım yetenekleriyle geleceğini beklemeli
    Amaç, sıfıra doğru yarışın ortasında ezilip yok olmamak

  • ABD, DeepSeek'in gerçekten sadece H800 kullanıp kullanmadığını doğrulamak için Singapur'daki GPU faturalarını didiklerken, dünyanın geri kalanı bu optimizasyonları tam H100'lerde çalıştırabilecek mi yani?
    ABD yaptırımları yüzünden H100 edinmek ya da erişmek zordu ve ABD kibri yüzünden kendi emirlerinin tüm dünyayı kapsadığına inanıyormuş gibi yapmaya da devam ederken?
    Bunu böyle anlamak doğru mu?

  • Herkesin beklediği PTX bu sefer dahil mi?

    • Evet, csrc/kernels dizininde bir kısmı var
      asm diye aratırsanız kullanıldığı yerleri bulabilirsiniz
    • Herkesin beklediği PTX'in neden bu kadar önemli olduğunu geri kalanımıza açıklamak gerekiyor

  • İnsanların CPU performansını son damlasına kadar sıkmak için assembly hack'lediği ya da belgelenmemiş komutlar aradığı 80'ler ve 90'ları hatırlatıyor
    Bir gün derleyiciler yeterince optimize olacak ya da GPU'lar öyle güçlü hale gelecek ki, bugün CPU'larda olduğu gibi bu tür numaralar artık büyük fark yaratmayacak gibi geliyor