vLLM kết hợp Novita AI ra mắt Chord: Toán tử INT4 MoE giúp tăng tốc Kimi K2.x lên tới 2.15 lần

TL;DR

Novita AI đã mã nguồn mở Chord, một toán tử CUDA MoE W4A16 hiệu năng cao dành cho các kích hoạt (activations) BF16, trọng số INT4 và thang đo (scales) group-32. Được xây dựng cho các hình dạng phục vụ (serving shapes) của Kimi K2.x, đường dẫn được lập chỉ mục (indexed path) của nó hiển thị gốc nhập khẩu humming tương thích với Humming, được chọn bằng tham số --quantization humming trên các phiên bản vLLM tương thích. Việc tích hợp các toán tử nhóm (grouped operators) với backend Humming của vLLM vẫn đang trong quá trình thực hiện.

Đo lường theo từng lớp so với đường dẫn tương ứng của Humming công khai:

Hình 1. Độ trễ mỗi lệnh gọi so với Humming công khai trên sáu kịch bản đo lường, giá trị thấp hơn là tốt hơn. Đọc từng bảng riêng biệt: bảng giải mã (decode) B300 so sánh với mặc định của Humming chưa được tinh chỉnh vì Humming công khai không cung cấp bảng tinh chỉnh cho SM100/SM103, trong khi mọi bảng H200 đều là so sánh giữa các phiên bản đã được tinh chỉnh. Biểu đồ từ kho lưu trữ Chord; bảng đầy đủ có trong docs/performance.md.

Ý tưởng đằng sau những con số này là một nhân (kernel) W4A16 MoE không thể phù hợp cho mọi yêu cầu. Số lượng token được định tuyến cho mỗi chuyên gia (expert) thay đổi theo nhiều bậc độ lớn giữa giai đoạn tiền điền (prefill) và giải mã (decode), và chính số lượng đó, chứ không phải tổng số token, mới quyết định lịch trình nào hiệu quả hơn. Chord chọn lịch trình dựa trên hình dạng thực tế mà nó nhận được.

Đây là các phép đo ở cấp độ nhân, không phải là lời hứa về mức tăng hiệu năng end-to-end tương tự cho mọi khối lượng công việc. Các bảng đầy đủ, định nghĩa hình dạng và phương pháp đo thời gian nằm trong docs/performance.md và docs/benchmarking.md. Mã nguồn và các bảng nhân trong bài viết này tham chiếu đến commit 7ca91d8 (ngày 14 tháng 9 năm 2026).

Hai họ nhân (kernel families)

Nhánh chính hiện tại cung cấp hai họ độc lập:

Tích hợp vLLM

Cài đặt gói và chọn backend Humming hiện có:

Bản phân phối cung cấp cả gốc mô-đun chord và humming. Facade lười (lazy facade) của vLLM giải quyết humming.{dtypes,config,layer,schema,utils.weight}; đường dẫn được lập chỉ mục mặc định có thể sử dụng tích hợp hiện có này mà không cần bản vá framework cụ thể cho Chord trên các nhánh có hỗ trợ group-scale WNA16 được ghi chú bên dưới. Lược đồ được cung cấp hỗ trợ uint4, group-32, thang đo BF16 và định dạng checkpoint INT4 group-32 được nén (pack-quantized) của compressed-tensors được sử dụng bởi Kimi K2.x; các lược đồ lượng tử hóa không được hỗ trợ sẽ gây lỗi khi tải thay vì âm thầm chọn sai nhân.

Tích hợp nhóm với backend Humming của vLLM đang được thực hiện (WIP). API toán tử nhóm độc lập được hiển thị bên dưới. TP8 vẫn giữ cấu hình h200_tp8 được lập chỉ mục vì một trọng số TP8 phải phục vụ cả hai giai đoạn.

Các chi tiết triển khai khác:

Tối ưu hóa nhân

Nhân được lập chỉ mục (Indexed kernels)

Họ được lập chỉ mục bắt nguồn từ commit 4351af3 của inclusionAI/humming công khai. Các chế độ khối lượng công việc dưới đây thúc đẩy các cấu hình nhân khác nhau, được chọn trước khi trọng số được đóng gói:

Hình 2. Khối lượng công việc tiền điền và giải mã điển hình. Nhãn 9–15 hàng/chuyên gia minh họa một trường hợp kiểm thử giải mã; 80 token/chuyên gia là ngưỡng heuristic block-M cho tiền điền. Không cái nào định nghĩa một công tắc runtime giữa tiền điền và giải mã: các cấu hình và bố cục trọng số được cố định khi tải mô hình, trong khi số lượng token tinh chỉnh lịch trình trong mỗi cấu hình.

Tiền điền H200 và TP8

Quy tắc tok_e được cố tình đơn giản hóa để giải thích nhưng cụ thể cho hình dạng MoE. Dưới khoảng 80 token được định tuyến cho mỗi chuyên gia, bộ giải quyết giữ lại tìm kiếm số lượng khối cơ sở; trên điểm đó, nó định cỡ block_m xung quanh các hàng được đệm (padded rows) của mỗi chuyên gia và giới hạn thanh ghi. TP8 sử dụng các cửa sổ phẳng hơn vì kích thước trung gian hẹp của nó để lại ít ô N hơn để lấp đầy một SM:

Vòng lặp chính WGMMA cũng thực hiện quản lý phụ thuộc bất đồng bộ theo lô. Thay vì đợi mọi nhóm lệnh, nó thực hiện commit sau một lần lặp warp-K và giữ một nhóm đang chạy trong khi quá trình tải bộ nhớ chia sẻ tiếp theo và giải lượng tử hóa INT4 bắt đầu:

Các thanh ghi trọng số đệm kép (double-buffered) cho phép quá trình tải và giải lượng tử hóa tiếp theo chồng lấp với nhóm WGMMA đang thực thi. Bộ tích lũy không được tiêu thụ cho đến phần kết (epilogue), và quá trình xả cuối cùng vẫn đợi mọi thao tác WGMMA đang thực thi.

Giải mã được lập chỉ mục H200 và Blackwell

Tại một vài hàng được định tuyến cho mỗi chuyên gia, đường dẫn WGMMA bị giới hạn bởi rào cản (barrier-bound). Cấu hình giải mã hoán đổi các toán hạng MMA để các trọng số đã giải lượng tử hóa chiếm toán hạng MMA-M, sử dụng m16n8k16 và hỗ trợ 4 CTA/SM với block-M 8. Một lịch trình ô-token bán tĩnh đo được 186 µs so với 216 µs cho lịch trình hoàn toàn động ở mức 9–15 token/chuyên gia. Việc hợp nhất giải lượng tử hóa trừ-rồi-chia-tỉ-lệ vào quá trình trích xuất nibble bảo toàn thứ tự làm tròn BF16 chưa hợp nhất. Cùng một họ lệnh MMA được biên dịch cho SM100/SM103; các hình dạng giải mã Blackwell lớn hơn sử dụng các ô MMA rộng hơn không bị hoán đổi. Không cần nhân tcgen05 cho các số lượng token này.

Nhân SM90 nhóm (Grouped SM90 kernels)

Backend nhóm là một họ nhân khác, không phải là tên gọi khác của nhân được lập chỉ mục. Nó chuyên biệt hóa cơ sở hạ tầng Hopper GEMM của DeepGEMM cho W4A16 và điều chỉnh nó cho JIT và trình khởi chạy của Chord. Cả hai chế độ đều sử dụng TMA, WGMMA chuyên biệt warp và giải lượng tử hóa group-32, nhưng cách định tuyến và bố cục trọng số vật lý của chúng khác nhau:

Hình 3. Nơi lưu trữ phần đệm. Indexed để các kích hoạt không được đệm; các chỉ số định tuyến của nó mang theo các sentinel đệm. Contiguous đệm mỗi chuyên gia đến ranh giới 128 hàng; masked dự trữ một ngân sách hàng cố định cho mỗi chuyên gia.

Các điểm truy cập toán tử nhóm (chỉ dành cho Chord API):

Ở đây expected_m là một số nguyên Python dương được sử dụng để chọn khởi chạy; masked_m giữ các số lượng hợp lệ cho mỗi chuyên gia. Đầu ra được che (masked output) là phẳng ngay cả khi a3 là ba chiều, và người tiêu dùng phải bỏ qua các hàng vượt quá số lượng hợp lệ của mỗi chuyên gia.

Chế độ được ghi lại trong trọng số đã chuẩn bị và được kiểm tra tại thời điểm điều phối, vì vậy việc vô tình đưa trọng số đã đóng gói tiền điền vào nhân giải mã sẽ gây lỗi rõ ràng. Điều phối nhóm sở hữu việc tìm kiếm bố cục SM90 và không chấp nhận các ghi đè block_m hoặc tuning_config của indexed. Việc phân giải nhân và tải cubin được ghi nhớ theo mô tả (và các ghi đè tinh chỉnh CHORD_W4A16_*), loại bỏ việc tìm kiếm lặp lại ở phía host, vốn đo được khoảng 30 µs trong các lần khởi chạy giải mã nhỏ.

Vòng lặp chính nhóm là liên tục và chuyên biệt warp: một nhóm warp sản xuất sử dụng TMA để dàn dựng kích hoạt, trọng số đã đóng gói và các ô thang đo, trong khi các nhóm warp tiêu thụ thực thi WGMMA và ghi kết quả BF16. Đường dẫn chuyển tiếp nhìn thấy các byte INT4 đã được hoán vị và các thang đo MN-major, và mô tả được lưu trong bộ nhớ đệm ánh xạ mỗi hình dạng (mode, M, N, K, expert_count) tới cubin của nó mà không cần lặp lại tìm kiếm bố cục trên mỗi lệnh gọi giải mã.

Heuristic nhóm có một vài lựa chọn cụ thể cho khối lượng công việc W4A16:

Những quy tắc này là lý do tại sao nhóm không sử dụng lại bảng tinh chỉnh của indexed: backend nhóm chọn (BM, BN, BK, cluster, stages) từ chế độ và hình dạng thực tế tại thời điểm điều phối. Trong phạm vi H200 EP8 được trích dẫn ở trên, lợi thế tiền điền thu hẹp ở mức 512 hàng/chuyên gia vì cả hai triển khai đều tiến gần đến cùng một giới hạn thông lượng; các lựa chọn về ô và đường ống quan trọng nhất ở các khối có kích thước nhỏ và trung bình.

Các phép đo

Các bảng nhân sử dụng triton.testing.do_bench và so sánh từng đường dẫn Chord với backend Humming công khai tương ứng trên cùng một GPU. Các so sánh indexed sử dụng cùng hình dạng và bản vẽ định tuyến; các so sánh nhóm khớp với số lượng hàng mỗi chuyên gia. Chạy hai bộ kiểm thử để kiểm tra đầu ra của Chord so với tham chiếu PyTorch thuần túy và in các bảng thời gian của nó trên các GPU được hỗ trợ:

Tóm tắt dưới đây cộng thời gian gọi gate/up và down từ các bảng đầy đủ. Tốc độ tăng tốc là Humming (gate_up + down) / Chord (gate_up + down); nó loại trừ việc định tuyến, kích hoạt và giao tiếp.

So sánh B300 được định tính một cách có chủ đích: Humming công khai không có bảng tinh chỉnh SM100/SM103, vì vậy thời gian mặc định của nó là một tham chiếu chưa được tinh chỉnh. Các tỷ lệ indexed của H200 là so sánh giữa các phiên bản đã được tinh chỉnh.

Cả hai họ đều được đo lường so với cùng một phiên bản Humming công khai, 4351af3. Các hàng nhóm so sánh với các đường dẫn grouped_contiguous/grouped_masked của chính Humming thay vì đường dẫn indexed, vì đó là hợp đồng mà backend này thay thế. Humming hiển thị cả hai dưới dạng các giá trị GemmType được điều phối thông qua nhân chung của nó thay vì các tệp CUDA riêng biệt, và benchmarks/bench_humming.py chọn chúng bằng --gemm_type grouped_contiguous hoặc --gemm_type grouped_masked. Số lượng hàng mỗi chuyên gia được khớp ở cả hai phía tại các bội số của ranh giới ô 128 hàng — --balanced ở phía Humming và các trường hợp căn chỉnh trong tests/test_w4a16_grouped.py — vì vậy mỗi hàng là cùng một hình dạng GEMM cho cả hai triển khai và không có ô nào bị lãng phí cho phần đệm.

Phục vụ end-to-end

Một báo cáo phục vụ trước đó đã đo lường đường dẫn TP8 indexed trên Kimi-K2.6 với 8×H200, TP8 + DCP8, bộ nhớ đệm FP8 KV và các yêu cầu ShareGPT. Cả hai nhà cung cấp đều sử dụng cùng lệnh --quantization humming.