DeepSeek เปิดตัวไลบรารีโอเพนซอร์ส DeepEP สำหรับการฝึกและอนุมาน MoE

 (github.com/deepseek-ai)
1 คะแนน โดย GN⁺ 2025-02-26 | 1 ความคิดเห็น | แชร์ทาง WhatsApp
  • DeepEP เป็นไลบรารีสื่อสารประสิทธิภาพสูงที่มุ่งเน้น Expert Parallelism (EP) ในการฝึกและอนุมาน ML สมัยใหม่ โดยมี all-to-all GPU kernel สำหรับ MoE dispatch/combine และรองรับความแม่นยำต่ำ เช่น FP8
  • รีลีส V2 รีแฟกเตอร์ EP ครั้งใหญ่ ทำให้ได้ประสิทธิภาพเทียบเท่าหรือดีกว่า V1 ด้วย ทรัพยากร SM ที่น้อยกว่ามาก และเปลี่ยน backend จาก NVSHMEM ไปเป็น NCCL Gin ที่เบากว่า
  • ในการทดสอบตามคอนฟิก V3 วัดภายใต้เงื่อนไข 8K tokens per batch, hidden 7168, top 8 experts, FP8 dispatch และ BF16 combine โดย V2 ทำได้สูงสุด 1.3 เท่าของประสิทธิภาพพีก เทียบกับ V1 และลดการใช้ SM ได้สูงสุด 4 เท่า
  • kernel ทั้งหมดคอมไพล์ที่ runtime ด้วย โมดูล JIT ขนาดเบา จึงไม่ต้องคอมไพล์ CUDA ตอนติดตั้ง และ V2 รวม API แบบ throughput สูงกับ latency ต่ำไว้ในอินเทอร์เฟซ ElasticBuffer เดียว
  • ต้องใช้ Hopper SM90 GPU, Python 3.8+, CUDA 12.3+, PyTorch 2.10+, NCCL 2.30.4+, NVLink และเครือข่าย RDMA ระหว่าง node โดย Engram, PP และ CP เป็น ฟีเจอร์เชิงทดลอง

ขอบเขตที่ DeepEP มีให้

  • DeepEP (DeepEveryParallel) เป็นไลบรารีสื่อสารประสิทธิภาพสูงสำหรับการฝึกและการอนุมานแมชชีนเลิร์นนิงสมัยใหม่
  • ฟังก์ชันหลักในปัจจุบันคือ Expert Parallelism (EP) โดยมี all-to-all GPU kernel แบบ throughput สูงและ latency ต่ำสำหรับ MoE dispatch และ combine
  • รองรับ การสื่อสารความแม่นยำต่ำ รวมถึง FP8
  • ยังมี primitive เชิงทดลองสำหรับ Pipeline Parallelism (PP), Context Parallelism (CP) และ remote memory access (Engram)
  • kernel ทั้งหมดคอมไพล์ที่ runtime ด้วยโมดูล JIT (Just-In-Time) ขนาดเบา จึงไม่ต้องคอมไพล์ CUDA ระหว่างขั้นตอนติดตั้ง
  • แม้จะออกแบบมาให้เบา แต่ยังตั้งเป้าให้ได้ประสิทธิภาพแตะหรือเกินขีดจำกัด bandwidth ของฮาร์ดแวร์ในหลายคอนฟิก

การเปลี่ยนแปลงหลักในรีลีส V2

  • V2 เป็นรีลีสที่รีแฟกเตอร์ Expert Parallelism ใหม่ทั้งหมด
    • ออกแบบมาเพื่อให้ได้ประสิทธิภาพสุดขีดด้วยทรัพยากร SM น้อยกว่า V1 หลายเท่า
    • รองรับโดเมน scale-up และ scale-out ที่ใหญ่ขึ้น
    • backend เปลี่ยนจาก NVSHMEM ไปเป็น NCCL Gin backend ที่เบากว่า
  • ฟีเจอร์ใหม่มีดังนี้
    • JIT compilation ทั้งหมด
    • NCCL Gin backend แบบ header-only และน้ำหนักเบา
    • สามารถนำ NCCL communicator เดิมกลับมาใช้ซ้ำได้
    • ใน EPv2 รวม API แบบ throughput สูงและ latency ต่ำไว้ในอินเทอร์เฟซ ElasticBuffer เดียว
    • มี GEMM layout ใหม่
    • รองรับโดเมน scale-up และ scale-out ที่ใหญ่ขึ้นได้ถึง EP2048
    • คำนวณจำนวน SM และ QP เชิงวิเคราะห์ ทำให้ไม่จำเป็นต้อง auto-tuning อีกต่อไป
    • ยังคงรองรับทั้ง hybrid mode และ direct mode
    • ในการฝึก legacy ลักษณะคล้าย V3 ลดการใช้ SM จาก 24 ตัวเหลือ 4–6 ตัว โดยยังรักษาประสิทธิภาพเทียบเท่าหรือดีกว่าเดิม
    • 0 SM Engram ที่ใช้ RDMA
    • 0 SM PP ที่ใช้ RDMA
    • 0 SM CP ที่ใช้ Copy Engine

ข้อจำกัดและฟีเจอร์ที่กำลังพัฒนา

  • V2 มี การใช้ขนาด buffer มากกว่า V1
  • ไม่รองรับ 0 SM RDMA low-latency EP อีกต่อไป
  • Engram, PP, CP เป็นฟีเจอร์เชิงทดลอง
  • ฟีเจอร์ที่กำลังพัฒนามีดังนี้
    • Elastic GPU & CPU buffers ซึ่งเป็นพื้นที่ virtual address ต่อเนื่องที่ผสมการ map หน่วยความจำกายภาพของ GPU และ CPU
      • มุ่งไปในทิศทางที่ทำให้ Engram หรือ EP ที่ไม่สมดุลทำงานได้แบบอัตโนมัติเต็มรูปแบบและโปร่งใส
    • งานที่ใช้ EP replay เพื่อจัดการ load imbalance และลดขนาด intermediate buffer
    • การอัปเดต all-gather และการทำ reduce-scatter สำหรับ DP และ TP
  • เอกสาร V1 ที่ใช้ NVSHMEM มีให้ที่ docs/legacy.md

ผลการวัดประสิทธิภาพ

  • ทดสอบตามคอนฟิก V3 ภายใต้เงื่อนไขต่อไปนี้
    • 8K tokens ต่อ batch
    • 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
  • ค่าที่วัดเป็น logical bandwidth เช่น 90 GB/s ของ EP 8 x 2 รวม local rank traffic ไว้ด้วย
  • V2 ทำได้สูงสุด 1.3 เท่าของประสิทธิภาพพีก เมื่อเทียบกับ V1 และลดจำนวน SM ได้สูงสุด 4 เท่า
  • ขณะนี้ละผลลัพธ์ของคอนฟิก EP ที่ใหญ่กว่าไว้ และแนะนำให้ผู้ใช้ทำ benchmark เอง
  • จากประสบการณ์ภายใน คาดว่า kernel จะยัง saturate bandwidth ของฮาร์ดแวร์ต่อไปแม้สเกลจะใหญ่ขึ้น
  • ข้อมูลประสิทธิภาพของ V1 อยู่ที่ docs/legacy.md

การติดตั้งและข้อกำหนด

  • ข้อกำหนดมีดังนี้
    • Hopper SM90 GPU หรือสถาปัตยกรรมที่รองรับ SM90 PTX ISA
    • Python 3.8 ขึ้นไป
    • CUDA 12.3 ขึ้นไปสำหรับ SM90 GPU
    • PyTorch 2.10 ขึ้นไป
    • NCCL 2.30.4 ขึ้นไป
    • NVLink สำหรับการสื่อสารภายใน node
    • เครือข่าย RDMA สำหรับการสื่อสารระหว่าง node
  • แนะนำให้ติดตั้ง NCCL ผ่าน pip เพื่อให้ DeepEP ตรวจพบได้อัตโนมัติในสภาพแวดล้อม Python
pip install "nvidia-nccl-cu13>=2.30.4" --no-deps
  • เพื่อรองรับเมธอด legacy ยังพึ่งพา NVSHMEM ด้วย และดูคำแนะนำการติดตั้งได้ที่ NVSHMEM Installation Guide
  • ตัวอย่างการ 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
  • ติดตั้งด้วยคำสั่งต่อไปนี้
python setup.py install
  • หลังติดตั้งแล้ว สามารถ import deep_ep ในโปรเจกต์ Python เพื่อใช้งานได้

อินเทอร์เฟซที่มี ElasticBuffer เป็นศูนย์กลาง

  • ใน V2 การทำงาน EP ทั้งหมดถูกรวมไว้ภายใต้อินเทอร์เฟซ ElasticBuffer เดียว
    • จัดการ API ทั้งแบบ throughput สูงและ latency ต่ำด้วยอินเทอร์เฟซเดียวกัน
    • สามารถ initialize buffer โดยระบุการตั้งค่า MoE ได้โดยตรง
    • คำนวณจำนวน SM และ QP ที่เหมาะสมเชิงวิเคราะห์
  • ตัวอย่างการ initialize buffer คือใช้ ElasticBuffer.get_buffer_size_hint() เพื่อคำนวณขนาดที่ต้องใช้ และตรวจสอบว่าสามารถ reuse buffer เดิมได้หรือไม่
  • เมื่อสร้าง buffer ใหม่ ให้ระบุ num_max_tokens_per_rank, hidden, num_topk, use_fp8_dispatch เป็นต้น
  • ใช้ _buffer.get_theoretical_num_sms(num_experts, num_topk) เพื่อรับจำนวน SM เชิงทฤษฎีที่จะใช้กับ communication kernel
  • หากระบุ num_sms โดยตรงในการเรียก dispatch และ combine จะ override ค่าที่คำนวณไว้

รูปแบบการใช้งานสำหรับการฝึก, prefill และ decoding

  • ในการฝึกหรือ inference prefill นั้น MoE dispatch จะ route token ไปยัง expert ที่เกี่ยวข้องในทุก rank
    • รองรับอินพุต BF16 และ FP8
    • handle เก็บ routing metadata ที่จำเป็นสำหรับการเรียก combine ภายหลัง
    • handle.num_recv_tokens_per_expert_list ให้จำนวน token ต่อ expert ที่จำเป็นสำหรับ GEMM
  • backward pass ของ MoE dispatch ถูกประมวลผลด้วย combine ในทางปฏิบัติ
  • MoE combine จะ reduce output ของ expert กลับไปยัง rank เดิม
  • backward pass ของ MoE combine ถูกประมวลผลด้วย dispatch ในทางปฏิบัติ
  • การซ้อนทับกันของ communication และ computation จัดการผ่านอินเทอร์เฟซ EventOverlap
    • สามารถทำ computation อิสระระหว่างที่ communication กำลังดำเนินอยู่ได้
    • ก่อนใช้ผลลัพธ์ ให้ synchronize compute stream ด้วย event.current_stream_wait()
  • ใน inference decoding ก็ใช้ ElasticBuffer เดียวกัน
    • เมื่อ gating decision ไม่เปลี่ยน สามารถ reuse routing metadata ด้วย cached_handle
    • รูปแบบนี้หลีกเลี่ยงการคำนวณ layout ใหม่และการ synchronize กับ CPU

Environment variables และค่าคงที่ตอน build

  • การตั้งค่าทั่วไป
    • EP_BUFFER_DEBUG: แสดงข้อมูลดีบักของ buffer initialization, การประมาณ SM และ backend
    • EP_SUPPRESS_NCCL_CHECK: ระงับการตรวจสอบ NCCL version mismatch
    • EP_AVOID_RECORD_STREAM: หลีกเลี่ยง record_stream ของ output tensor
    • EP_NUM_TOPK_IDX_BITS: override จำนวนบิตสำหรับ encoding top-k index
  • การตั้งค่า networking
    • EP_NIC_NAME: ชื่อ NIC เริ่มต้นที่จะใช้ query คุณสมบัติ NIC ค่าเริ่มต้นคือ mlx5_0
    • EP_OVERRIDE_RDMA_SL: override RDMA service level index
    • EP_DISABLE_GIN: ปิดใช้งาน NCCL Gin backend
  • การตั้งค่า JIT
    • EP_JIT_CACHE_DIR: ไดเรกทอรี cache ของ kernel ที่คอมไพล์แล้ว ค่าเริ่มต้นคือ $HOME/.deep_ep
    • EP_JIT_NVCC_COMPILER: path ของ NVCC compiler
    • EP_JIT_CPP_STANDARD: เวอร์ชัน C++ standard ค่าเริ่มต้นคือ 20
    • EP_JIT_DUMP_PTX, EP_JIT_DUMP_SASS, EP_JIT_DUMP_ASM: การตั้งค่าที่เกี่ยวกับการ dump output PTX/SASS
  • environment variable บางตัวทำงานแบบ persistent
    • ถูก capture ตอน build และรวมเป็นค่าเริ่มต้นของแพ็กเกจที่ติดตั้ง
    • หากไม่ override ด้วย environment variable ปัจจุบันตอน import ค่าเริ่มต้นนี้จะถูกใช้โดยอัตโนมัติ
    • ตัวแปรเป้าหมายคือ EP_JIT_CACHE_DIR, EP_JIT_PRINT_COMPILER_COMMAND, EP_NUM_TOPK_IDX_BITS, EP_NCCL_ROOT_DIR
  • รายละเอียดเพิ่มเติมดูได้ที่ test_ep.py หรือเอกสาร Python

คำแนะนำการตั้งค่าเครือข่าย

  • DeepEP ผ่านการทดสอบครบถ้วนบนเครือข่าย InfiniBand
  • ตามทฤษฎีแล้วเข้ากันได้กับ RDMA over Converged Ethernet หรือ RoCE ด้วย
  • การแยก traffic
    • รองรับผ่าน Virtual Lanes ของ InfiniBand
    • แนะนำให้แยก expert-parallel workload และ workload อื่นไปอยู่คนละ virtual lane
    • ใน V2 สามารถควบคุมการจัดสรร virtual lane ได้ด้วยอาร์กิวเมนต์ sl_idx หรือ environment variable EP_OVERRIDE_RDMA_SL
  • Adaptive routing
    • เป็นฟีเจอร์ routing ขั้นสูงที่ทำให้สวิตช์ InfiniBand กระจาย traffic ไปหลายเส้นทางอย่างสม่ำเสมอ
    • แนะนำให้เปิดใช้งานในทุกสภาวะโหลดของเครือข่าย แม้อาจมี latency เพิ่มขึ้น
  • Congestion control
    • ปิดใช้งานเพราะส่งผลเสียต่อ bandwidth สูงสุด
    • หากหลีกเลี่ยง congestion ไม่ได้ แนะนำให้จัด workload นั้นไว้ใน virtual lane ที่มี priority ต่ำ
  • PCI atomic mode
    • หากฮาร์ดแวร์รองรับ แนะนำให้ตั้งค่า PCI_ATOMIC_MODE ของ NIC เพื่อปรับปรุงประสิทธิภาพ RDMA atomic operation
sudo mlxconfig -y -d mlx5_$i set PCI_ATOMIC_MODE=4

Experimental branch และ community fork

  • Experimental branch
    • Zero-copy: ลบการ copy ระหว่าง PyTorch tensor กับ communication buffer เพื่อลดการใช้ SM ของ kernel ทั่วไปอย่างมาก
    • Eager: ใช้ low-latency protocol เพื่อลบ extra RTT latency ที่ RDMA atomic OP เพิ่มเข้ามา
    • Hybrid-EP: implementation backend ใหม่ที่ใช้ TMA instructions, การใช้ SM ต่ำสุด, รองรับโดเมน NVLink ที่ใหญ่ขึ้น, การซ้อนทับ communication/computation แบบละเอียดระดับ single-batch, PCIe kernel และรองรับ NVFP4
    • AntGroup-Opt: ชุด optimization ที่เขียนโดย AntGroup Network Platform Department
    • Mori-EP: รองรับโหมด latency ต่ำบน ROCm/AMD GPU โดยอิง backend MORI
    • nvDev: branch ที่อิง V2 ซึ่งรวมฟีเจอร์ CUDA ล่าสุด เช่น Compute Fabric Transport
  • Community fork
    • uccl/uccl-ep: รองรับการรัน DeepEP บน GPU ต่างค่าย เช่น Nvidia และ AMD รวมถึง NIC เช่น EFA, Broadcom และ CX7
    • Infrawaves/DeepEP_ibrc_dual-ports_multiQP: เพิ่มโซลูชัน multi-QP และการรองรับ dual-port NIC ให้กับ IBRC transport
    • antgroup/DeepXTrace: diagnostic analyzer สำหรับค้นหา slow rank ได้อย่างมีประสิทธิภาพและแม่นยำ
    • ROCm/mori: ไลบรารีสื่อสารรุ่นถัดไปของ AMD สำหรับ AI workload ที่ไวต่อประสิทธิภาพ เช่น Wide EP, KVCache transfer และ Collectives

License และการอ้างอิง

  • DeepEP V2 สร้างขึ้นบน NCCL Gin backend
  • โค้ดใน repository เผยแพร่ภายใต้ MIT License
  • รายการอ้างอิงคือ DeepEP: an efficient expert-parallel communication library และระบุปีเป็น 2025

บทความที่เกี่ยวข้อง

1 ความคิดเห็น

 
GN⁺ 2025-02-26
ความคิดเห็นจาก Hacker News

  • ทำประสิทธิภาพได้สุดขีดด้วยการค้นหาและใช้ คำสั่ง PTX ที่ไม่ได้ระบุไว้ในเอกสาร ld.global.nc.L1::no_allocate.L2::256B
    คำสั่งนี้เข้าถึงหน่วยความจำ GPU แบบ volatile ด้วยตัวปรับแต่ง PTX แบบอ่านอย่างเดียวที่ไม่สอดคล้องกัน .nc จึงอาจทำให้เกิดพฤติกรรมที่ไม่ได้นิยามไว้
    อย่างไรก็ตาม มีการทดสอบแล้วว่าบนสถาปัตยกรรม Hopper เมื่อใช้ร่วมกับ .L1::no_allocate จะรับประกันความถูกต้อง และว่ากันว่าประสิทธิภาพดีขึ้นมาก

    • ในทางปฏิบัติ NVIDIA จะสามารถเปลี่ยนพฤติกรรมของ คำสั่งนอกเอกสาร นี้อย่างละเอียดอ่อนในสถาปัตยกรรมใหม่ภายหลัง จนพลิกกระดานได้ ไม่ว่าจะตั้งใจหรือไม่ก็ตามหรือเปล่า?

  • รู้สึกเหมือนเด็กได้เข้าไปในร้านขายลูกกวาด
    มีทริกมากมายที่ถ้าจะ reverse engineer ให้ถูกต้องจากการดูแค่เปเปอร์คงใช้เวลานานเกินไป และหวังว่าการเปิดเผยในสัปดาห์นี้จะเปิดยุคเรอเนซองส์ที่ใช้ MoE เป็นโมเดลวิชาการพื้นฐาน

    • ถ้ามองจากมุมนี้ ก็ไม่เข้าใจว่าเกิดอะไรขึ้นระหว่าง งานปฏิบัติจริงของโมเดลล้ำสมัย กับโมเดลเชิงวิชาการ
      ฝ่ายแรกตั้งแต่ GPT-4 เป็นต้นมาก็เป็น MoE กันหมดแล้ว แต่โมเดลเปิดเผยต่อสาธารณะมักเป็นโมเดล dense เสมอ ยกเว้น DeepSeek V3 กับ Mixtral

  • ทีมนี้ทำให้ไม่รักไม่ได้จริง ๆ
    กำลังผลัก ขอบเขตของโอเพนซอร์ส เพื่อทุกคน

    • เป็นฝั่งที่เขียนเว้นวรรคเหมือน Open AI™ น่ะ
    • จริง ๆ แล้วไม่ใช่โอเพนซอร์ส
      ถ้าอยากดูโมเดลโอเพนซอร์สจริง ๆ ให้ดู OLMo 2 ของ AI2: https://allenai.org/blog/olmo2
      พวกเขาแชร์ทุกอย่างที่จำเป็นต่อการสร้างโมเดลซ้ำ รวมถึงตัวข้อมูลเองด้วยจริง ๆ
      ในลิงก์ข้างต้นก็กล่าวว่า “เพราะวิทยาศาสตร์แบบเปิดอย่างสมบูรณ์ต้องการมากกว่าแค่น้ำหนักที่เผยแพร่ต่อสาธารณะ เราจึงยินดีที่จะแชร์การอัปเดต OLMo ใหม่กับชุมชนโมเดลภาษาที่กว้างขึ้น ซึ่งรวมถึงน้ำหนัก ข้อมูล โค้ด สูตรการฝึก checkpoint ระหว่างทาง และโมเดลที่ปรับแต่งด้วยคำสั่ง”

  • Zuckerberg ควรเลิกอ้างว่า Meta เปิดเผย AI เป็นโอเพนซอร์สได้แล้ว
    ถึงขั้นทำโฆษณาทีวี แต่จริง ๆ แล้ว เปิดแค่น้ำหนัก และไม่มีโค้ด
    AI โอเพนซอร์สจริง ๆ มีแค่ DeepSeek เท่านั้น

    • ถ้าพูดให้เคร่งครัด DeepSeek ก็ไม่ได้โอเพนซอร์สเท่า OLMo หรือ Open Euro
      เพราะไม่ได้เปิดเผยข้อมูล
    • DeepSeek ไม่ใช่โอเพนซอร์สจริง ๆ แน่นอน
      การจะเป็นโอเพนซอร์สต้องใช้ ไลเซนส์โอเพนซอร์สจริง ๆ อย่างที่ OSI ระบุไว้ และต้องแชร์โค้ดก่อนและหลังการฝึก โค้ดที่เกี่ยวกับการปรับจูน โค้ดประเมินผล ทุกอย่างที่เกี่ยวกับความปลอดภัยและการเซ็นเซอร์ และอาจรวมถึงข้อมูลฝึกทั้งหมดด้วย
      ไม่เช่นนั้นก็ไม่สามารถสร้างน้ำหนักซ้ำได้ และการแชร์น้ำหนักก็คล้ายกับการแชร์โปรแกรมที่คอมไพล์แล้ว
      เท่าที่รู้ โมเดลที่แข่งขันได้และเป็นโอเพนซอร์สจริง ๆ มีเพียง OLMo 2 ของ AI2: https://allenai.org/blog/olmo2
      ล่าสุดยังเปิดเผยแอปที่ทำ inference บนอุปกรณ์ด้วย ซึ่งก็เป็นโอเพนซอร์สเช่นกัน: https://allenai.org/blog/olmoe-app
      ยังมีอีกโมเดลชื่อ Tülu 3 และว่ากันว่าประสิทธิภาพดีกว่า DeepSeek V3: https://allenai.org/blog/tulu-3-405B
    • Meta ขัดเกลา PyTorch มานานกว่า 10 ปีแล้ว
      แทบทุกอย่างที่จำเป็นในการฝึก LLM รวมถึงเทคโนโลยีล่าสุด อยู่ในนั้นแล้ว
      ยังต้องการอะไรอีก? ส่วนโค้ดที่เฉพาะเจาะจงกับโครงสร้างพื้นฐานของ Meta หรือ?
    • นับ PyTorch ด้วยไหม?
    • น้ำหนักที่เปิดเผย = ก้อน binary
      เท่ากับย้อนกลับไปเป็นโมเดล FREEWARE / SHAREWARE
      ควรใช้ถ้อยคำแบบนี้กับน้ำหนักที่ “เปิดเผยต่อสาธารณะ”

  • มี การสื่อสาร all-to-all ที่มีประสิทธิภาพและปรับแต่งมาอย่างดี รองรับทั้งภายในโหนดและระหว่างโหนดผ่าน NVLink และ RDMA, kernel throughput สูงสำหรับการฝึกและ inference prefill, kernel latency ต่ำสำหรับ inference decoding, รองรับการ dispatch แบบ FP8 ดั้งเดิม และการควบคุมทรัพยากร GPU ที่ยืดหยุ่นเพื่อทับซ้อนการคำนวณกับการสื่อสาร
    X: https://x.com/deepseek_ai/status/1894211757604049133

  • แรงจูงใจของงาน DeepSeek อาจผิดก็ได้
    เช่น อาจเป็น ความพยายามที่รัฐสนับสนุน เพื่อลดข้อได้เปรียบการครองนำด้าน AI ของสหรัฐฯ ให้เหลือศูนย์ แต่ผลสุทธิต่อทุกคนทั่วโลกนั้นยอดเยี่ยมล้วน ๆ
    ในกรณีเลวร้ายที่สุด ต่อให้ทำด้วยเหตุผลที่ผิด ก็ต้องขอบคุณ DeepSeek เพราะพวกเขากำลังทำสิ่งที่ OpenAI โกหกคนทั้งโลกมาหลายปีว่าจะทำ

    • ในขอบเขตความสัมพันธ์ระหว่างประเทศ ความถูกผิดไม่ได้ใช้ได้มากขนาดนั้น
      การเปิดสิ่งนี้เป็นโอเพนซอร์ส “ผิด” มากกว่าการ ห้ามส่งออก Nvidia GPU ระดับสูง หรือ?
      การเปิดโอเพนซอร์สของ DeepSeek น่าจะเป็นเพียงสิ่งที่เกิดขึ้นภายใต้ความยินยอมของพรรคคอมมิวนิสต์จีน และเป็นผลดีพร้อมกันทั้งต่อพรรคคอมมิวนิสต์จีนกับชุมชน AI โอเพนซอร์สที่กว้างขึ้น ไม่ควรถูกมองว่าเป็นจุดยืนเชิงหลักการใด ๆ
      การหาวิธีทำลายข้อได้เปรียบในการแข่งขันของประเทศอื่นเป็นกิจกรรมหลักของรัฐบาลทุกแห่ง ไม่ว่าใหญ่หรือเล็ก

  • เป็นรอบที่ 2 ของการเปิดโอเพนซอร์สโดย บริษัท Open AI™ ตัวจริง ภายใต้ไลเซนส์ MIT
    อีกครั้งที่ DeepSeek เปิดกว้างกว่าบริษัทมูลค่า 157,000 ล้านดอลลาร์ที่อ้างว่าตัวเอง “Open”
    แทบไม่มีใครพูดถึง Llama ของ Meta แต่ทุกคนควรคาดว่า Llama 4 จะมาพร้อมความสามารถด้าน reasoning
    เป้าหมายคือไม่ถูกบีบตายกลางการแข่งขันที่มุ่งสู่ศูนย์

  • ระหว่างที่สหรัฐฯ คุ้ย ใบเสร็จ GPU ในสิงคโปร์เพื่อดูว่า DeepSeek ใช้แค่ H800 จริงไหม นั่นหมายความว่าส่วนที่เหลือของโลกสามารถนำ optimization เหล่านี้ไปรันบน H100 แบบเต็ม ๆ ได้ใช่ไหม?
    ทั้งที่เพราะมาตรการคว่ำบาตรของสหรัฐฯ ทำให้หา H100 หรือเข้าถึงได้ยาก และเพราะความโอหังของสหรัฐฯ ก็ยังแสร้งทำต่อไปว่าคำสั่งของตนครอบคลุมทั้งโลก?
    เข้าใจแบบนี้ถูกไหม?

  • PTX ที่ทุกคนรอคอยรวมอยู่ในครั้งนี้ไหม?

    • ใช่ มีบางส่วนอยู่ในไดเรกทอรี csrc/kernels
      ค้นหา asm ก็จะเจอจุดที่ใช้
    • ต้องอธิบายให้คนอื่น ๆ ฟังว่าทำไม PTX ที่ทุกคนรอคอยถึงสำคัญขนาดนั้น

  • ทำให้นึกถึง ยุค 80–90 ที่ผู้คนแฮ็ก assembly หรือค้นหาคำสั่งที่ไม่ได้ระบุไว้ในเอกสารเพื่อรีดประสิทธิภาพจาก CPU
    สักวัน compiler อาจปรับแต่งได้ดีพอ หรือ GPU อาจทรงพลังเกินไป จนทริกแบบนี้ไม่สร้างความแตกต่างมากอีกต่อไปเหมือน CPU ทุกวันนี้