DeepSeek ra mắt Tile Kernels: thư viện GPU kernel TileLang chạm giới hạn phần cứng

TL;DR

DeepSeek vừa mở mã nguồn TileKernels — một thư viện GPU kernel viết bằng TileLang, tập trung vào các thao tác nặng xung quanh Mixture-of-Experts: top-k gating, token routing, cast FP8/FP4/E5M6, RMSNorm gradient reduction và Sinkhorn normalization. DeepSeek khẳng định nhiều kernel đã chạy trong pipeline training và inference nội bộ, và "approach the limit of hardware performance regarding compute intensity and memory bandwidth". Giấy phép MIT, chạy trên GPU NVIDIA Hopper (SM90) và Blackwell (SM100).

Có gì mới

TileKernels không phải là một model — đây là tập hợp kernel "đường ngang" phục vụ mô hình MoE quy mô lớn. Repo chia thành các nhóm:

• Gating — top-k expert selection cho MoE
• MoE Routing — token-to-expert mapping + weight normalization
• Quantization — cast FP8 / FP4 / E5M6 fused
• Engram — gating kernel fused RMSNorm và gradient reduction
• Manifold HyperConnection — Sinkhorn normalization
• Transpose — batched transpose
• Modeling — wrapper PyTorch autograd function

Đáng chú ý: TileKernels không chứa kernel GEMM hay attention — hai thứ đó đã nằm ở các repo chị em DeepGEMM (FP8 GEMM fine-grained scaling) và FlashMLA (Multi-head Latent Attention). TileKernels phụ trách phần routing, chuẩn hoá, quantization — những chỗ dễ trở thành nút cổ chai khi scale MoE.

Vì sao quan trọng

Khi một model như DeepSeek-V3 có 160 expert và 236B parameter (21B activated), routing + normalization không còn là "phần phụ" — chúng chiếm phần lớn ngân sách memory bandwidth. Kernel hand-written CUDA cho các pha này cực đắt về thời gian kỹ sư. TileLang đưa bài toán về ~80 dòng Python mỗi kernel mà vẫn đạt gần hiệu năng assembly viết tay. TileKernels biến pattern này thành một thư viện production-ready do chính team đã train được model 236B chia sẻ.

Với hệ sinh thái inference engine (vLLM, SGLang) và lab đang muốn tự train MoE, đây là một gift box: code tham chiếu từ team hiểu rõ nhất cách ép kiệt một cụm H800.

Số liệu kỹ thuật

Con số hiệu năng định lượng nằm ở các repo chị em. FlashMLA (attention) trên H800 SXM5 đạt 3000 GB/s ở cấu hình memory-bound và 660 TFLOPS compute-bound với dense MLA decoding; sparse MLA đạt 410 TFLOPS decode và 640 TFLOPS prefill. Bản thân TileLang cho thấy kernel Python ~80 dòng vượt Triton khoảng 1.98× và PyTorch 3.76× trong các benchmark chính thức của dự án.

So sánh với các kernel stack khác

Ai được lợi

• Lab train MoE LLM trên cụm H100/H800/B200 — có thể thay thẳng vào path gating và tiết kiệm nhiều tuần profile.
• Inference engine (vLLM, SGLang) hỗ trợ model DeepSeek-V3/V3.2 — dùng TileKernels cho pha routing sẽ thu ngắn inter-token latency.
• Nhà nghiên cứu operator — TileLang cho phép prototype operator mới trong vài giờ thay vì vài tuần CUDA.
• Team FP8/FP4 low-precision training — kernel cast fused tránh round-trip memory khi chuyển precision.

Giới hạn & giá

TileKernels miễn phí hoàn toàn (MIT). Nhưng rào kỹ thuật không thấp:

• Chỉ chạy trên Hopper hoặc Blackwell — A100, consumer RTX, AMD đều không được.
• Yêu cầu toolchain rất mới: CUDA 13.1+, PyTorch 2.10+, TileLang 0.1.9+.
• Không có GEMM / attention trong repo — phải kết hợp DeepGEMM + FlashMLA để có full stack.
• README không công bố con số FLOPS chi tiết cho từng kernel; claim "approach hardware limit" ở dạng định tính.

Sắp tới

TileKernels là mảnh ghép thứ ba trong bộ đồ nghề kernel của DeepSeek (sau FlashMLA tháng 2/2025 và DeepGEMM tháng 3/2025). Mô hình DeepSeek-V3.2-Exp đã ship kernel ở cả hai dạng TileLang và CUDA từ 29/9/2025. Với việc TileLang liên tục mở rộng hỗ trợ phần cứng (Apple Metal tháng 10/2025, AMD MI300X, Huawei Ascend), kỳ vọng hợp lý là thế hệ model DeepSeek kế tiếp sẽ được bọc hoàn toàn bằng stack TileLang, và TileKernels sẽ đón thêm kernel chuyên biệt cho các operator mới như sparse attention hay linear attention.

Nguồn: deepseek-ai/TileKernels, tile-ai/tilelang, DeepSeek V3.2-Exp news, FlashMLA.