DeepSeek merilis pustaka open source DeepEP untuk pelatihan dan inferensi MoE

 (github.com/deepseek-ai)
1 poin oleh GN⁺ 2025-02-26 | 1 komentar | Bagikan ke WhatsApp

DeepEP

DeepEP adalah pustaka komunikasi untuk Mixture-of-Experts (MoE) dan expert parallelism (EP). Pustaka ini menyediakan kernel all-to-all GPU berkecepatan tinggi dan latensi rendah, yang dikenal sebagai dispatch dan combine pada MoE. DeepEP juga mendukung komputasi presisi rendah termasuk FP8. Sesuai dengan algoritme group-limited gating yang diusulkan dalam makalah DeepSeek-V3, DeepEP menyediakan kernel yang dioptimalkan untuk transfer bandwidth domain asimetris, dengan meneruskan data dari domain NVLink ke domain RDMA. Kernel ini memberikan throughput tinggi sehingga cocok untuk pelatihan dan pekerjaan prefill inferensi. Selain itu, DeepEP juga mendukung kontrol jumlah SM (Streaming Multiprocessors). Untuk decoding inferensi yang sensitif terhadap latensi, DeepEP menyertakan kernel latensi rendah yang meminimalkan latensi dengan menggunakan RDMA murni. Pustaka ini juga memperkenalkan metode overlap komunikasi-komputasi berbasis hook yang tidak mengambil alih sumber daya SM.

Performa

Kernel umum yang menggunakan forwarding NVLink dan RDMA

  • Pada H800, kernel umum diuji dengan bandwidth maksimum NVLink sekitar 160 GB/s dan kartu jaringan RDMA CX7 InfiniBand 400 Gb/s (~50 GB/s bandwidth maksimum).
  • Mengikuti konfigurasi pretraining DeepSeek-V3/R1 (4096 token per batch, hidden size 7168, top-4 group, top-8 expert, dispatch FP8, dan combine BF16).

Kernel latensi rendah yang menggunakan RDMA murni

  • Pada H800, kernel latensi rendah diuji dengan kartu jaringan RDMA CX7 InfiniBand 400 Gb/s (~50 GB/s bandwidth maksimum).
  • Mengikuti konfigurasi produksi DeepSeek-V3/R1 yang umum (128 token per batch, hidden size 7168, top-8 expert, dispatch FP8, dan combine BF16).

Mulai cepat

Persyaratan

  • GPU Hopper (dukungan untuk lebih banyak arsitektur atau perangkat mungkin ditambahkan nanti)
  • Python 3.8 atau lebih baru
  • CUDA 12.3 atau lebih baru
  • PyTorch 2.1 atau lebih baru
  • NVLink untuk komunikasi intra-node
  • Jaringan RDMA untuk komunikasi antar-node

Mengunduh dan memasang dependensi NVSHMEM

DeepEP bergantung pada NVSHMEM yang telah dimodifikasi. Anda harus merujuk ke panduan instalasi untuk memasangnya.

Konfigurasi jaringan

DeepEP telah diuji sepenuhnya pada jaringan InfiniBand, dan secara teori juga kompatibel dengan RDMA over Converged Ethernet (RoCE).

Isolasi traffic

InfiniBand mendukung isolasi traffic melalui Virtual Lanes (VL). Untuk mencegah interferensi antarjenis traffic yang berbeda, disarankan untuk memisahkan pekerjaan ke virtual lane sebagai berikut:

  • Pekerjaan yang menggunakan kernel umum
  • Pekerjaan yang menggunakan kernel latensi rendah
  • Pekerjaan lainnya

Di DeepEP, Anda dapat mengatur alokasi virtual lane dengan menetapkan variabel lingkungan NVSHMEM_IB_SL.

Adaptive routing

Adaptive routing adalah fitur routing lanjutan yang disediakan oleh switch InfiniBand, yang dapat mendistribusikan traffic secara merata ke beberapa jalur. Saat ini kernel latensi rendah mendukung adaptive routing, tetapi kernel umum belum mendukungnya (mungkin akan didukung segera). Mengaktifkan adaptive routing pada kernel umum antar-node dapat menyebabkan deadlock atau kerusakan data. Untuk kernel latensi rendah, mengaktifkan adaptive routing dapat sepenuhnya menghilangkan kemacetan jaringan akibat konflik routing, tetapi menambah latensi ekstra. Untuk performa optimal, konfigurasi berikut direkomendasikan:

  • Aktifkan adaptive routing di lingkungan dengan beban jaringan tinggi
  • Gunakan static routing di lingkungan dengan beban jaringan rendah

Pengendalian kemacetan

Karena tidak ada kemacetan yang signifikan yang diamati di lingkungan produksi, pengendalian kemacetan dinonaktifkan.

Antarmuka dan contoh

Menggunakan contoh untuk pelatihan model atau prefill inferensi

Kernel umum dapat digunakan pada tahap pelatihan model atau prefill inferensi (tanpa bagian backward).

Menggunakan contoh untuk decoding inferensi

Kernel latensi rendah dapat digunakan pada tahap decoding inferensi.

Hal yang perlu diperhatikan

  • Demi performa ekstrem, ditemukan dan digunakan instruksi PTX yang tidak terdokumentasi, ld.global.nc.L1::no_allocate.L2::256B. Instruksi ini memicu perilaku yang tidak terdefinisi saat mengakses memori GPU volatile menggunakan modifier PTX read-only yang tidak koheren. Namun, berdasarkan pengujian dengan .L1::no_allocate pada arsitektur Hopper, performanya meningkat jauh lebih baik. Jika kernel tidak berfungsi di platform lain, Anda dapat menonaktifkannya dengan menambahkan DISABLE_AGGRESSIVE_PTX_INSTRS=1 ke setup.py atau melaporkan issue.
  • Untuk performa yang lebih baik di klaster, disarankan menjalankan semua pengujian dan menggunakan konfigurasi auto-tuning yang optimal. Konfigurasi default dioptimalkan pada klaster internal DeepSeek.

Lisensi

Repositori kode ini dirilis di bawah lisensi MIT, dan kode yang merujuk ke NVSHMEM (termasuk csrc/kernels/ibgda_device.cuh dan third-party/nvshmem.patch) tunduk pada NVSHMEM SLA.

Bacaan terkait

1 komentar

 
GN⁺ 2025-02-26
Komentar Hacker News
  • Menemukan dan menggunakan instruksi PTX yang tidak terdokumentasi demi performa ekstrem. Instruksi ini dapat menyebabkan perilaku yang tidak terdefinisi karena menggunakan modifier PTX hanya-baca yang tidak konsisten untuk mengakses memori GPU volatil. Namun, akurasi yang diuji dengan .L1::no_allocate pada arsitektur Hopper akan terjamin dan performanya akan jauh lebih baik
  • Zuckerberg harus berhenti mengklaim bahwa Meta melakukan open source AI. Mereka hanya merilis bobot, bukan kodenya. Satu-satunya AI open source yang sesungguhnya adalah DeepSeek
  • Rasanya seperti anak kecil di toko permen. Beberapa trik ini akan butuh waktu terlalu lama untuk direkayasa balik dengan benar hanya berdasarkan makalah. Semoga rilis minggu ini memulai renaisans penggunaan MoE sebagai model akademik dasar
  • Mustahil untuk tidak menyukai orang-orang ini. Mereka benar-benar mendorong batas open source untuk kita semua. Terima kasih sudah berbagi
    • Komunikasi all-to-all yang efisien dan teroptimasi
    • Dukungan intra-node dan antar-node melalui NVLink dan RDMA
    • Kernel throughput tinggi untuk prefilling pelatihan dan inferensi
    • Kernel latensi rendah untuk decoding inferensi
    • Dukungan dispatch FP8 native
    • Kontrol sumber daya GPU yang fleksibel untuk overlap komputasi-komunikasi
  • Motif di balik pekerjaan DeepSeek mungkin salah (misalnya, upaya yang didukung negara untuk menghapus keunggulan kepemimpinan AS dalam AI). Namun, secara global hasilnya sungguh fantastis
    • Bahkan dalam skenario terburuk (jika mereka melakukan ini karena alasan yang salah), tetap terima kasih kepada DeepSeek. Mereka benar-benar melakukan apa yang telah dibohongi OpenAI kepada seluruh dunia selama bertahun-tahun
    • Benar-benar luar biasa
  • Penasaran apakah PTX yang semua orang nantikan akhirnya disertakan kali ini
  • Instruksi PTX yang disebut dalam laporan teknis seharusnya ditautkan ke kodenya di sini
  • Sementara AS melacak kuitansi GPU di Singapura untuk memastikan DeepSeek hanya memakai H800, seluruh dunia bisa menjalankan optimasi ini pada H100 penuh
    • Bertanya-tanya apakah mereka hanya berpura-pura bahwa mendapatkan atau mengakses H100 itu sulit karena kesombongan mereka yang percaya bahwa sanksi AS dan perintah mereka berlaku untuk seluruh dunia
  • Ini adalah rilis open source kedua dari perusahaan "Open AI™" yang sebenarnya, dan berada di bawah lisensi MIT
    • DeepSeek lebih terbuka dibanding perusahaan yang mengklaim $157B+
    • Hampir tak ada yang membicarakan Llama milik Meta, dan semua orang seharusnya berharap Llama 4 dirilis disertai alasan yang jelas
    • Tujuannya adalah agar tidak terjebak di tengah dalam perlombaan menuju nol