vLLM 및 Novita AI, Kimi K2.x용 더 빠른 INT4 MoE 커널 출시: Chord
TL;DR
Novita AI는 Chord를 공개했습니다. 이는 NVIDIA B300 GPU에서 공개된 Humming 백엔드 대비 최대 2.15×, H200 GPU에서 최대 **1.33×**의 성능 향상을 제공하는 고성능 W4A16(INT4 가중치, BF16 활성화) MoE CUDA 연산자입니다. 이 연산자는 기존 humming 임포트 루트를 통해 vLLM과 통합되지만, 완전한 그룹 커널 통합은 아직 진행 중입니다.
핵심 기여
Chord는 Kimi K2.x 모델의 특정 서빙 형태에 최적화된 두 가지 독립적인 커널 패밀리—인덱스 및 그룹—를 제공합니다. 인덱스 커널은 H200 EP8 프리필, H200 TP8 단일 인스턴스 서빙, B300 EP8 디코딩 워크로드를 대상으로 하며, 그룹 커널(연속 프리필 및 마스크된 디코딩)은 SM90 GPU를 대상으로 하며 DeepGEMM 기반 구현을 사용합니다.
"이 숫자들의 핵심 아이디어는 하나의 W4A16 MoE 커널이 모든 요청에 적합할 수 없다는 점입니다. 프리필과 디코딩 간 전달된 토큰 수는 수량 주기로 달라지며, 이 수치가 총 토큰 수가 아니라 어떤 스케줄이 승리할지를 결정합니다."
성능 요약
인덱스 커널은 H200 및 B300에서 Humming을 능가
| 시나리오 | GPU | 성능 향상 (Chord vs. Humming) |
|---|---|---|
| H200 EP8 프리필 | H200 | 1.11–1.20× |
| H200 TP8 단일 인스턴스 서빙 | H200 | 1.17–1.33× |
| H200 EP8 디코딩 (게이트/업) | H200 | 1.16–1.24× (다운은 최대 1.31×) |
| B300 EP8 디코딩 (튜닝되지 않은 Humming) | B300 | 1.81–2.15× |
이 수치는 레이어별 커널 타이밍이며, 엔드투엔드 지연 보장이 아닙니다. 전체 표, 형태 정의 및 벤치마킹 방법론은 리포지토리의 docs/performance.md 및 docs/benchmarking.md에 문서화되어 있습니다.
그룹 커널은 SM90에서 유사한 성능 향상 제공
| 시나리오 | GPU | 성능 향상 |
|---|---|---|
| H200 EP8 그룹 프리필 (128행/전문가) | H200 | 1.31× |
| H200 EP8 그룹 디코딩 (전문가당 32토큰) | H200 | 1.35× |
그룹 커널은 다른 가중치 레이아웃(연속 또는 마스크)과 지속적인 워프 전용 메인 루프를 사용하며, TMA 로드, 디퀀타이제이션 및 WGMMA 실행을 파이프라인화합니다.
vLLM과의 통합
설치 및 활성화
pip install git+https://github.com/novitalabs/chord.git
vllm serve <kimi-k2.x-int4-model> --quantization humming
# 또는 vLLM 설정 파일에서 moe_backend="humming"으로 설정
- 패키지는
chord와humming모듈 루트를 모두 설치합니다. vLLM의 지연 패서드는 W4A16 그룹 스케일 형식을 지원하는 브랜치에서는 Chord 전용 패치 없이도humming.{dtypes,config,layer,…}를 해결합니다. - 지원되지 않는 양자화 방식은 오류를 발생시키며, 잘못된 커널로 자동 전환되지 않습니다.
- 그룹 통합은 현재 진행 중이며, 현재는 인덱스 경로만 즉시 사용 가능합니다.
런타임 고려사항
- 프로파일 선택(EP8 대 TP8)은 vLLM이 전달하는 프로젝션 형태에서 추론되며, 인덱스 커널에는 추가 shard 인수가 필요하지 않습니다.
- 인덱스 커널은 vLLM의 과도하게 할당된
moe_align_block_size버퍼에서 작동할 수 있으며, 신뢰할 수 있는 라우팅 검증이 비활성화된 경우 CUDA 그래프 캡처가 가능합니다. - 환경 변수
VLLM_HUMMING_USE_F16_ACCUM과VLLM_BATCH_INVARIANT는 비활성화 상태를 유지해야 합니다. 두 백엔드 모두 해당 계산 옵션을 구현하지 않습니다. - 오래된 vLLM 브랜치는 Chord 형식을 인식하기 위해
_supports_quant_scheme에 그룹-32 양자화 키를 추가해야 할 수 있습니다.
커널 수준 최적화
인덱스 커널 기법
- 배치된
wait<1>WGMMA 파이프라인화는 가중치 로드, 디퀀타이제이션 및 이전 WGMMA 그룹을 겹치게 하여 업프로젝션에서 3–6 %, 다운프로젝션에서 1–5 %의 성능 향상을 제공합니다. - 전문가당 토큰 수(
tok_e) 히ュ리스틱은 전달된 전문가당 토큰 수에 따라 블록-M 크기를 선택하며, 80 토큰/전문가를 기준으로 기본 블록 수 탐색과 패딩된 행 인식 레이아웃 간 전환을 수행합니다. - 제한된 2-CTA/SM 윈도우는 중간 크기 타일의 주거 워프를 증가시켜
cp.async수집 지연을 숨깁니다. - 형태 인식 스트림-K 게이팅은 M×N 그리드가 포화 상태일 때 스트림-K를 비활성화하여 불필요한 분할/감소 오버헤드를 피합니다.
그룹 SM90 커널 기법
- 연속 프리필은 각 전문가를 128행 경계로 패딩하고,
BLOCK_K=64와 전치된 BF16 스케일을 사용합니다. - 마스크된 디코딩은 전문가당 고정된 행 예산을 예약하고,
BLOCK_K=128로 가중치를 패킹하며, 예상 토큰 수에서BLOCK_M을 선택합니다. - 웨이브 인식 BN 선택은 충분한 N-타일이 존재할 때만 BN256을 선택하고, 그렇지 않으면 고용도를 유지하기 위해 BN128을 사용합니다.
- 지연 조정된 버퍼링된-K 깊이는 일반적인 디코딩에 대해 약 512개의 버퍼링된 K 요소를 목표로 하며, 큰 마스크된 타일의 경우 공유 메모리를 소진하지 않도록 약 768까지 확장합니다.
엔드투엔드 서빙 영향
8×H200 구성(Kimi-K2.6, FP8 KV 캐시, ShareGPT 워크로드)에 대한 별도의 서빙 보고서는 미미하지만 일관된 개선을 보였습니다:
| 지표 | Humming | Chord | Δ |
|---|---|---|---|
| 평균 TTFT | 2022 ms | 1849 ms | –8.6 % |
| 프리필 처리량 | 20,716 토큰/초 | 22,712 토큰/초 | +9.6 % |
| 디코딩 처리량 (배치 8) | 483 토큰/초 | 503 토큰/초 | +4.1 % |
| 디코딩 처리량 (배치 64) | 1650 토큰/초 | 1740 토큰/초 | +5.5 % |
| 디코딩 처리량 (배치 128) | 2514 토큰/초 | 2715 토큰/초 | +8.0 % |
| 보고서는 OCRBench 및 GSM8K에서 정확도 저하 없음을 관찰했습니다. |
로드맵
- vLLM의 Humming 백엔드와의 완전한 그룹 통합을 완료하여 기존 프레임워크를 통해 연속 및 마스크된 연산자를 노출합니다.
- B200/B300 GPU용 EP8 프리필 커널 공개; 내부 테스트에서 희망적인 성능을 보였으며, 후속 릴리스에서 공유될 예정입니다.
시작하기
Chord는 https://github.com/novitalabs/chord에서 사용 가능합니다. 리포지토리는 다음을 포함합니다:
- 시작 가이드 – 설치 및 기본 사용법.
- 최적화 문서 – 커널 히ュ리스틱에 대한 상세 설명.
- 성능 표 – 시나리오별 완전한 지연 수치.
- 튜닝 및 벤치마킹 방법론 – 재현 가능한 테스트 스크립트. 커뮤니티의 기여, 이슈 보고 및 벤치마크 결과를 환영합니다.
감사의 말
Chord의 인덱스 경로는 inclusionAI/Humming (커밋 4351af3)에 기반합니다. 그룹 SM90 백엔드는 DeepGEMM의 Hopper GEMM 인프라를 적응했습니다. 프로젝트는 Apache-2.0 라이선스 하에 배포됩니다. Chord 개발에 기여한 Novita AI 팀과 통합을 지원한 vLLM 유지보수자 및 커뮤니티에 감사드립니다.