DeepSeek เปิดซอร์ส kernel สำหรับ Huawei Ascend 950 ทั้งชุด — FlashMLA ทำได้ 95% ของ peak ตอน prefill, TileLang 0.1.15 มี backend Ascend ในตัว แต่ FlashMLA รุ่นใหม่ตัดการรองรับ Hopper

DeepSeek ปล่อยโค้ดชุดใหญ่สำหรับชิป Huawei Ascend 950 เมื่อ 30 กันยายน 2026 ทั้งหมดลงในรีโปบน GitHub ภายในเวลาไล่เลี่ยกัน คือ attention kernel ใน FlashMLA, kernel ด้าน MoE และ quantization ใน TileKernels, TopK kernel ใน DeepSelect และ backend Ascend ใน TileLang 0.1.15 ซึ่ง PR ที่เพิ่ม backend นี้ใช้ชื่อว่า “[Public Release 9/30]” ตรงกับวันเดียวกัน Reuters รายงานโดยอ้างโพสต์ WeChat ของ DeepSeek ว่า Huawei “สนับสนุนเต็มที่” และทั้งสองฝ่ายร่วมกันปรับ supernode ที่มีชิป Ascend 950 จำนวน 128 ตัว

ประเด็นที่สำคัญไม่ได้อยู่ที่โค้ดชิ้นใดชิ้นหนึ่ง แต่อยู่ที่ว่า kernel ที่ใช้รัน DeepSeek-V4.1 จริงในตอนนี้มีเส้นทาง Ascend ที่เป็นโอเพนซอร์สแล้วแทบครบ และเขียนบน TileLang ภาษาเดียวกับฝั่ง NVIDIA

FlashMLA: 410 TFlops บน Ascend แต่เลิกรองรับ Hopper

README ของ FlashMLA ระบุว่า sparse attention (DeepSeek Sparse Attention) บน Ascend 950 ทำได้ 410 TFlops ตอน prefill (95% ของ peak ทางทฤษฎี) และ 360 TFlops ตอน decoding (83%) พร้อมรายงานเชิงลึกที่อธิบาย pipeline และเทคนิค online softmax ใน kernel ถ้าดูตัวเลขดิบ ฝั่ง B200 ยังเร็วกว่ามาก (sparse prefill 1,350 TFlops) แต่การใช้ฮาร์ดแวร์ได้ถึง 95% ของ peak แสดงว่าคอขวดของ Ascend ไม่ได้อยู่ที่ซอฟต์แวร์ฝั่ง attention แล้ว

จุดที่ทีมที่ใช้ FlashMLA อยู่ต้องระวัง: รุ่นนี้ ตัดการรองรับสถาปัตยกรรม Hopper (H100/H200) ตัดรุ่นโมเดลเก่า (V3 / V3.2 / V4.0) และ เปลี่ยนฟอร์แมต FP8/FP4 KV cache จนใช้ร่วมกับเวอร์ชันก่อนไม่ได้ ถ้ายังต้องรันบน Hopper ต้อง pin ไว้ที่ commit ba89a34 ตามที่ README บอก ส่วน fused kernel แบบ norm-RoPE-attn-RoPE-cast และ dense attention ยังรองรับเฉพาะ CUDA ฝั่ง Ascend ต้องใช้ CANN 9.2.0 ขึ้นไป และ torch_npu

TileKernels, DeepSelect และ TileLang

• TileKernels — kernel สำหรับ MoE routing, FP8/FP4 quantization และ Engram ได้ backend ตัวที่สองที่ เลือกอัตโนมัติตอนรัน โค้ด Python ชุดเดียวจึงรันได้ทั้งบน NVIDIA และ Huawei NPU README ขอบคุณ Huawei ที่ช่วยด้านวิศวกรรม
• DeepSelect — TopK kernel สำหรับ lightning indexer ของ DSA และ sampler อ้างว่าเร็วกว่า torch.topk 2–20 เท่า ฝั่ง Ascend ตอนนี้รองรับเฉพาะ bfloat16 (sampler ที่ใช้ float32 ยังมีแต่บน CUDA)
• TileLang 0.1.15 — เพิ่ม tilelang.ascend.language และ target="ascend" ที่ generate โค้ดให้ Ascend 950 โดยตรง จัดตาราง Cube/Vector unit ให้อัตโนมัติ และเขียนแบบ SIMD/SIMT ผสมกันได้ ส่วน Ascend รุ่น A2/A3 ยังอยู่ในโปรเจกต์ TileLang-Ascend ที่ชุมชนดูแล

รีโปทั้งสามของ DeepSeek (FlashMLA, TileKernels, DeepSelect) ใช้ไลเซนส์ MIT ส่วน TileLang เป็นโปรเจกต์ของ tile-ai ที่เริ่มจากนักวิจัยมหาวิทยาลัยปักกิ่ง

มุมมอง AI Curator

ข้อได้เปรียบใหญ่ที่สุดของ NVIDIA ไม่ใช่ตัวชิป แต่คือการที่ kernel สำคัญทุกตัวเขียนด้วย CUDA มาก่อน งานรอบนี้ของ DeepSeek โจมตีตรงจุดนั้น เพราะเขียน kernel ด้วย TileLang ซึ่งเป็นภาษากลางที่ generate ได้ทั้งสองฝั่ง ทำให้ค่าใช้จ่ายในการย้ายฮาร์ดแวร์ต่ำลงมาก ถ้าเทคนิคนี้ใช้ได้กับ DeepSeek ก็น่าจะใช้ได้กับชิปค่ายอื่นในอนาคตด้วย

สำหรับทีมในไทยและอาเซียนที่มีทั้งคลาวด์จีน (Huawei Cloud มี region ในไทย) และคลาวด์ฝั่งสหรัฐให้เลือก นี่เป็นครั้งแรกที่ inference stack ของโมเดลเปิดระดับแนวหน้ามีเส้นทางบน Ascend ที่ตรวจสอบโค้ดได้เองทั้งหมด แต่ควรรอผลจากคนนอกที่รัน end-to-end บน Ascend 950 จริงก่อน เพราะตัวเลข 95% เป็นของ kernel เดียว ไม่ใช่ throughput ของทั้งระบบ เรื่องที่ควรทำทันทีคือเช็กว่าทีมไหน build FlashMLA จาก main บนเครื่อง H100 อยู่บ้าง เพราะการ upgrade ครั้งถัดไปจะใช้ไม่ได้ทันที

แหล่งข่าว

• FlashMLA — README และ News 2026.09.30 (GitHub)
• A Deep Dive Into the Ascend Sparse Attention Forward Kernel — FlashMLA docs
• TileLang v0.1.15 release notes — Ascend 950 backend
• China’s AI industry closes ranks as Deepseek ships open-source software for Huawei’s Ascend chips — The Decoder