DeepSeek merilis library 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 adalah library komunikasi berperforma tinggi yang berfokus pada Expert Parallelism (EP) dalam pelatihan dan inferensi ML modern, menyediakan kernel GPU all-to-all untuk dispatch/combine MoE serta dukungan presisi rendah seperti FP8
  • Rilis V2 merefaktor EP secara menyeluruh sehingga menghasilkan performa setara atau lebih baik dengan resource SM jauh lebih sedikit dibanding V1, dan mengganti backend dari NVSHMEM ke NCCL Gin yang lebih ringan
  • Dalam pengujian berdasarkan konfigurasi V3 dengan 8K tokens per batch, hidden 7168, top 8 experts, FP8 dispatch, dan BF16 combine, V2 mencatat hingga 1,3x performa puncak dan penghematan SM hingga 4x dibanding V1
  • Semua kernel dikompilasi saat runtime sebagai modul JIT ringan dan tidak memerlukan kompilasi CUDA saat instalasi; V2 menyatukan API throughput tinggi dan latensi rendah dalam satu antarmuka ElasticBuffer
  • Membutuhkan GPU Hopper SM90, Python 3.8+, CUDA 12.3+, PyTorch 2.10+, NCCL 2.30.4+, NVLink, serta jaringan RDMA antar-node; Engram, PP, dan CP adalah fitur eksperimental

Cakupan yang disediakan DeepEP

  • DeepEP(DeepEveryParallel) adalah library komunikasi berperforma tinggi untuk pelatihan dan inferensi machine learning modern
  • Fitur utamanya saat ini adalah Expert Parallelism (EP), menyediakan kernel GPU all-to-all ber-throughput tinggi dan berlatensi rendah untuk dispatch dan combine MoE
  • Mendukung komunikasi presisi rendah, termasuk FP8
  • Juga mencakup primitive eksperimental untuk pipeline parallelism (PP), context parallelism (CP), dan remote memory access (Engram)
  • Semua kernel dikompilasi saat runtime sebagai modul JIT(Just-In-Time) ringan, sehingga tidak memerlukan kompilasi CUDA dalam proses instalasi
  • Meski dirancang ringan, targetnya adalah performa yang menyamai atau melampaui batas bandwidth hardware di berbagai konfigurasi

Perubahan utama pada rilis V2

  • V2 adalah rilis yang merefaktor Expert Parallelism sepenuhnya
    • Dirancang untuk mencapai performa ekstrem dengan resource SM beberapa kali lebih sedikit dibanding V1
    • Mendukung domain scale-up dan scale-out yang lebih besar
    • Backend diganti dari NVSHMEM ke NCCL Gin backend yang lebih ringan
  • Fitur baru meliputi
    • Kompilasi JIT penuh
    • NCCL Gin backend yang ringan dan header-only
    • Dapat menggunakan kembali NCCL communicator yang sudah ada
    • Di EPv2, API throughput tinggi dan API latensi rendah disatukan dalam satu antarmuka ElasticBuffer
    • Menyediakan layout GEMM baru
    • Mendukung domain scale-up dan scale-out yang lebih besar hingga EP2048
    • Tidak lagi memerlukan auto-tuning karena jumlah SM dan QP dihitung secara analitis
    • Tetap mendukung mode hybrid dan mode direct
    • Pada pelatihan legacy mirip V3, penggunaan SM dikurangi dari 24 menjadi 4~6 sambil mempertahankan performa setara atau lebih baik
    • 0 SM Engram berbasis RDMA
    • 0 SM PP berbasis RDMA
    • 0 SM CP berbasis Copy Engine

Batasan dan fitur yang sedang dikerjakan

  • V2 memiliki konsumsi ukuran buffer yang lebih besar dibanding V1
  • EP latensi rendah RDMA 0 SM tidak lagi didukung
  • Engram, PP, CP adalah fitur eksperimental
  • Fitur yang sedang dikerjakan meliputi
    • Elastic GPU & CPU buffers, ruang alamat virtual kontinu yang memetakan campuran memori fisik GPU dan CPU
      • Arahnya adalah memungkinkan Engram atau EP tidak seimbang yang sepenuhnya otomatis dan transparan
    • Pekerjaan untuk menangani ketidakseimbangan beban dan mengurangi ukuran buffer perantara dengan memanfaatkan EP replay
    • Pembaruan all-gather dan implementasi reduce-scatter untuk DP dan TP
  • Dokumentasi V1 berbasis NVSHMEM tersedia di docs/legacy.md

Hasil pengukuran performa

  • Pengujian dilakukan sesuai konfigurasi V3 dengan kondisi berikut
    • 8K tokens per batch
    • hidden dimension 7168
    • top 8 experts
    • FP8 dispatch
    • BF16 combine
  • Hasil utama adalah sebagai berikut
    • 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
  • Nilai pengukuran adalah bandwidth logis; misalnya 90 GB/s pada EP 8 x 2 mencakup traffic local rank
  • V2 mencapai hingga 1,3x performa puncak dibanding V1 dan menghemat jumlah SM hingga 4x
  • Hasil untuk konfigurasi EP yang lebih besar saat ini dihilangkan, dan pengguna disarankan melakukan benchmark sendiri
  • Berdasarkan pengalaman internal, kernel diperkirakan akan terus memenuhi bandwidth hardware meskipun skalanya membesar
  • Data performa V1 ada di docs/legacy.md

Instalasi dan persyaratan

  • Persyaratannya adalah sebagai berikut
    • GPU SM90 Hopper atau arsitektur yang mendukung SM90 PTX ISA
    • Python 3.8 atau lebih baru
    • CUDA 12.3 atau lebih baru untuk GPU SM90
    • PyTorch 2.10 atau lebih baru
    • NCCL 2.30.4 atau lebih baru
    • NVLink untuk komunikasi dalam node
    • Jaringan RDMA untuk komunikasi antar-node
  • NCCL disarankan diinstal melalui pip agar DeepEP dapat mendeteksinya secara otomatis di environment Python
pip install "nvidia-nccl-cu13>=2.30.4" --no-deps
  • Untuk mendukung metode legacy, DeepEP juga bergantung pada NVSHMEM; lihat NVSHMEM Installation Guide untuk instruksi instalasi
  • Contoh menjalankan build dan test saat pengembangan adalah sebagai berikut
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
  • Instalasi dilakukan dengan perintah berikut
python setup.py install
  • Setelah instalasi, deep_ep dapat di-import dan digunakan dari proyek Python

Antarmuka berpusat pada ElasticBuffer

  • Di V2, semua operasi EP disatukan di bawah satu antarmuka ElasticBuffer
    • API throughput tinggi dan latensi rendah ditangani melalui antarmuka yang sama
    • Buffer dapat diinisialisasi dengan menentukan konfigurasi MoE secara langsung
    • Jumlah SM dan QP optimal dihitung secara analitis
  • Contoh inisialisasi buffer menggunakan ElasticBuffer.get_buffer_size_hint() untuk menghitung ukuran yang dibutuhkan dan memeriksa apakah buffer yang ada dapat digunakan kembali
  • Saat membuat buffer baru, parameter seperti num_max_tokens_per_rank, hidden, num_topk, use_fp8_dispatch ditentukan
  • _buffer.get_theoretical_num_sms(num_experts, num_topk) digunakan untuk mendapatkan jumlah SM teoretis yang akan dipakai oleh kernel komunikasi
  • Jika num_sms ditentukan langsung dalam pemanggilan dispatch dan combine, nilai tersebut dapat menimpa nilai yang dihitung

Pola penggunaan untuk pelatihan, prefill, dan decoding

  • Dalam pelatihan atau prefill inferensi, MoE dispatch merutekan token ke expert terkait di semua rank
    • Mendukung input BF16 dan FP8
    • handle berisi metadata routing yang diperlukan untuk pemanggilan combine berikutnya
    • handle.num_recv_tokens_per_expert_list menyediakan jumlah token per expert yang dibutuhkan untuk GEMM
  • Backward pass MoE dispatch sebenarnya ditangani dengan combine
  • MoE combine melakukan reduce atas output expert kembali ke rank asal
  • Backward pass MoE combine sebenarnya ditangani dengan dispatch
  • Overlap komunikasi dan komputasi dikelola melalui antarmuka EventOverlap
    • Komputasi independen dapat dilakukan saat komunikasi berlangsung
    • Sebelum menggunakan hasil, compute stream disinkronkan dengan event.current_stream_wait()
  • Pada decoding inferensi, ElasticBuffer yang sama juga digunakan
    • Saat gating decision tidak berubah, metadata routing digunakan kembali dengan cached_handle
    • Pola ini menghindari perhitungan ulang layout dan sinkronisasi CPU

Variabel lingkungan dan nilai tetap saat build

  • Pengaturan umum
    • EP_BUFFER_DEBUG: menampilkan informasi debug untuk inisialisasi buffer, aproksimasi SM, dan backend
    • EP_SUPPRESS_NCCL_CHECK: menekan pemeriksaan ketidakcocokan versi NCCL
    • EP_AVOID_RECORD_STREAM: menghindari record_stream pada tensor output
    • EP_NUM_TOPK_IDX_BITS: menimpa jumlah bit encoding indeks top-k
  • Pengaturan jaringan
    • EP_NIC_NAME: nama NIC default yang digunakan untuk mengambil atribut NIC, nilai default mlx5_0
    • EP_OVERRIDE_RDMA_SL: menimpa index service level RDMA
    • EP_DISABLE_GIN: menonaktifkan NCCL Gin backend
  • Pengaturan JIT
    • EP_JIT_CACHE_DIR: direktori cache kernel terkompilasi, default $HOME/.deep_ep
    • EP_JIT_NVCC_COMPILER: path compiler NVCC
    • EP_JIT_CPP_STANDARD: versi standar C++, default 20
    • EP_JIT_DUMP_PTX, EP_JIT_DUMP_SASS, EP_JIT_DUMP_ASM: pengaturan terkait dump output PTX/SASS
  • Sebagian variabel lingkungan bekerja secara persistent
    • Ditangkap saat build dan dimasukkan sebagai nilai default paket instalasi
    • Jika tidak ditimpa oleh variabel lingkungan saat import, nilai default ini diterapkan otomatis
    • Variabel targetnya adalah EP_JIT_CACHE_DIR, EP_JIT_PRINT_COMPILER_COMMAND, EP_NUM_TOPK_IDX_BITS, EP_NCCL_ROOT_DIR
  • Detail tambahan dapat dilihat di test_ep.py atau dokumentasi Python

Rekomendasi konfigurasi jaringan

  • DeepEP telah diuji sepenuhnya pada jaringan InfiniBand
  • Secara teoretis, DeepEP juga kompatibel dengan RDMA over Converged Ethernet, yaitu RoCE
  • Isolasi traffic
    • Didukung melalui Virtual Lanes pada InfiniBand
    • Disarankan memisahkan workload expert-parallel dan workload lain ke virtual lane yang berbeda
    • Di V2, alokasi virtual lane dapat dikendalikan dengan argumen sl_idx atau variabel lingkungan EP_OVERRIDE_RDMA_SL
  • Adaptive routing
    • Ini adalah fitur routing lanjutan yang membuat switch InfiniBand mendistribusikan traffic secara merata ke beberapa jalur
    • Disarankan untuk diaktifkan pada semua kondisi beban jaringan, meskipun menimbulkan latensi tambahan
  • Congestion control
    • Dinonaktifkan karena merugikan bandwidth maksimum
    • Jika kemacetan tidak terhindarkan, disarankan mengalokasikan workload tersebut ke virtual lane berprioritas rendah
  • PCI atomic mode
    • Jika hardware mendukung, disarankan mengatur PCI_ATOMIC_MODE pada NIC untuk meningkatkan performa RDMA atomic operation
sudo mlxconfig -y -d mlx5_$i set PCI_ATOMIC_MODE=4

Branch eksperimental dan fork komunitas

  • Branch eksperimental
    • Zero-copy: menghapus penyalinan antara tensor PyTorch dan buffer komunikasi, sehingga sangat mengurangi penggunaan SM pada kernel umum
    • Eager: menggunakan protokol latensi rendah untuk menghilangkan extra RTT latency yang ditambahkan oleh RDMA atomic OP
    • Hybrid-EP: implementasi backend baru menggunakan TMA instructions, penggunaan SM minimum, dukungan domain NVLink yang lebih besar, overlap komunikasi dan komputasi fine-grained single-batch, kernel PCIe, serta dukungan NVFP4
    • AntGroup-Opt: seri optimasi yang ditulis oleh AntGroup Network Platform Department
    • Mori-EP: dukungan mode latensi rendah ROCm/AMD GPU berbasis backend MORI
    • nvDev: branch berbasis V2 yang mencakup fitur CUDA terbaru seperti Compute Fabric Transport
  • Fork komunitas
    • uccl/uccl-ep: mendukung menjalankan DeepEP pada GPU heterogen seperti Nvidia dan AMD serta NIC seperti EFA, Broadcom, dan CX7
    • Infrawaves/DeepEP_ibrc_dual-ports_multiQP: menambahkan solusi multi-QP dan dukungan NIC dual-port pada IBRC transport
    • antgroup/DeepXTrace: analyzer diagnostik yang menemukan slow rank secara efisien dan presisi
    • ROCm/mori: library komunikasi generasi berikutnya AMD untuk workload AI yang sensitif terhadap performa, seperti Wide EP, KVCache transfer, dan Collectives

Lisensi dan sitasi

  • DeepEP V2 dibangun di atas Gin backend NCCL
  • Kode repositori dirilis dengan MIT License
  • Item sitasinya adalah DeepEP: an efficient expert-parallel communication library, dengan tahun ditulis 2025

Bacaan terkait

1 komentar

 
GN⁺ 2025-02-26
Komentar Hacker News

  • Mereka menemukan dan memakai instruksi PTX yang tidak terdokumentasi ld.global.nc.L1::no_allocate.L2::256B untuk mendapatkan performa ekstrem
    Instruksi ini mengakses memori GPU yang volatile dengan modifier PTX read-only non-koheren .nc, sehingga bisa menyebabkan perilaku tak terdefinisi
    Namun pada arsitektur Hopper, saat dipakai bersama .L1::no_allocate, kabarnya sudah diuji menjamin kebenaran dan performanya jauh lebih baik

    • Secara realistis, mungkinkah NVIDIA nantinya, pada arsitektur baru, mengubah secara halus perilaku instruksi di luar dokumentasi ini dan membalik keadaan, entah disengaja atau tidak?

  • Rasanya seperti anak kecil masuk ke toko permen
    Ada banyak trik yang akan memakan waktu terlalu lama untuk di-reverse engineering dengan benar hanya dari paper, dan semoga rilis minggu ini membuka renaisans yang menjadikan MoE sebagai model akademik default

    • Dari sudut pandang ini, saya tidak paham apa yang terjadi antara praktik model frontier yang sebenarnya dan model akademik
      Yang pertama sejak GPT-4 sudah semuanya MoE, tetapi model terbuka, selain DeepSeek V3 dan Mixtral, sering kali selalu model dense

  • Tim ini benar-benar sulit untuk tidak disukai
    Mereka mendorong batas open source demi semua orang

    • Maksudnya yang ditulis terpisah seperti Open AI™
    • Sebenarnya ini bukan open source
      Kalau ingin melihat model yang benar-benar open source, lihat OLMo 2 dari AI2: https://allenai.org/blog/olmo2
      Mereka benar-benar membagikan semua yang dibutuhkan untuk mereproduksi model, sampai datanya sendiri
      Di tautan di atas juga disebutkan, “Karena open science yang lengkap membutuhkan lebih dari sekadar bobot terbuka, kami senang dapat membagikan pembaruan OLMo baru kepada komunitas language modeling yang lebih luas, termasuk bobot, data, kode, resep, checkpoint perantara, dan model instruction-tuned”

  • Zuckerberg harus berhenti mengklaim bahwa Meta merilis AI sebagai open source
    Mereka bahkan membuat iklan TV, tetapi sebenarnya hanya merilis bobot tanpa kode
    AI open source yang sebenarnya hanyalah DeepSeek

    • Secara ketat, DeepSeek juga tidak se-open source OLMo atau Open Euro
      Karena mereka tidak membuka datanya
    • DeepSeek jelas bukan open source sungguhan
      Agar menjadi open source, mereka harus memakai lisensi open source yang sebenarnya seperti yang dicantumkan OSI, serta membagikan kode pra- dan pasca-pelatihan, kode terkait tuning, kode evaluasi, semua hal terkait keamanan dan sensor, dan mungkin seluruh data pelatihan
      Kalau tidak, bobotnya tidak bisa direproduksi, dan membagikan bobot mirip dengan membagikan program yang sudah dikompilasi
      Sejauh yang saya tahu, satu-satunya model open source sungguhan yang kompetitif adalah OLMo 2 dari AI2: https://allenai.org/blog/olmo2
      Baru-baru ini mereka juga merilis aplikasi untuk inferensi di perangkat, dan ini juga open source: https://allenai.org/blog/olmoe-app
      Ada juga model lain bernama Tülu 3, yang katanya performanya lebih baik daripada DeepSeek V3: https://allenai.org/blog/tulu-3-405B
    • Meta telah memoles PyTorch selama lebih dari 10 tahun
      Hampir semua yang dibutuhkan untuk melatih LLM, termasuk teknologi terbaru, ada di sana
      Apa lagi yang dibutuhkan? Bagian kode yang khusus untuk infrastruktur Meta?
    • Apakah PyTorch juga termasuk?
    • Bobot terbuka = gumpalan biner
      Kita kembali ke model FREEWARE / SHAREWARE
      Istilah seperti ini seharusnya dipakai untuk bobot “terbuka”

  • Menyediakan komunikasi all-to-all yang efisien dan dioptimalkan, dukungan intra-node dan antar-node melalui NVLink dan RDMA, kernel throughput tinggi untuk training dan prefill inferensi, kernel latensi rendah untuk decoding inferensi, dukungan dispatch FP8 native, serta kontrol sumber daya GPU yang fleksibel untuk overlap komputasi-komunikasi
    X: https://x.com/deepseek_ai/status/1894211757604049133

  • Motivasi di balik kerja DeepSeek bisa saja keliru
    Misalnya, bisa berupa upaya yang disponsori negara untuk meniadakan keunggulan awal AS dalam AI, tetapi efek bersihnya bagi semua orang di seluruh dunia tetap saja luar biasa
    Dalam skenario terburuk sekalipun, yaitu jika mereka melakukannya karena alasan yang salah, saya tetap berterima kasih kepada DeepSeek; mereka benar-benar melakukan hal yang selama bertahun-tahun OpenAI klaim secara palsu akan dilakukan untuk dunia

    • Dalam ranah hubungan internasional, benar dan salah tidak terlalu berlaku sebesar itu
      Apakah merilis ini sebagai open source lebih “salah” daripada larangan ekspor GPU Nvidia kelas atas?
      Rilis open source DeepSeek mungkin hanya sesuatu yang, dengan persetujuan Partai Komunis Tiongkok, baik bagi Partai Komunis Tiongkok sekaligus komunitas AI open source yang lebih luas, dan tidak boleh dianggap sebagai posisi yang berprinsip
      Mencari cara menghilangkan keunggulan kompetitif negara lain adalah aktivitas utama semua pemerintahan, besar maupun kecil

  • Ini adalah putaran kedua rilis open source dari perusahaan Open AI™ yang sebenarnya di bawah lisensi MIT
    Sekali lagi, DeepSeek lebih terbuka daripada perusahaan senilai 157 miliar dolar yang mengklaim dirinya “Open”
    Hampir tidak ada yang membicarakan Llama dari Meta, tetapi semua orang seharusnya memperkirakan Llama 4 akan hadir dengan kemampuan reasoning
    Tujuannya adalah tidak tergencet di tengah perlombaan menuju nol

  • Saat AS mengubek-ubek kuitansi GPU di Singapura untuk memastikan apakah DeepSeek hanya memakai H800, apakah ini berarti seluruh dunia bisa menjalankan optimasi ini di H100 utuh?
    Sementara karena sanksi AS H100 sulit didapat atau diakses, dan karena arogansi AS mereka terus berpura-pura percaya bahwa perintah mereka berlaku di seluruh dunia?
    Apakah pemahaman seperti ini benar?

  • Apakah PTX yang semua orang tunggu-tunggu kali ini sudah disertakan?

    • Benar, ada sebagian di direktori csrc/kernels
      Cari asm untuk menemukan tempat penggunaannya
    • Orang-orang lain perlu penjelasan mengapa PTX yang semua orang tunggu-tunggu itu begitu penting

  • Ini mengingatkan pada era 80–90-an ketika orang-orang meretas assembly atau mencari instruksi yang tidak terdokumentasi untuk memeras performa CPU
    Suatu hari nanti, compiler mungkin akan cukup optimal atau GPU akan menjadi terlalu kuat, sehingga trik seperti ini tidak lagi membuat perbedaan besar seperti CPU masa kini