vLLM x Novita AI: Chord W4A16 INT4 MoE Kernel для Kimi K2.x

TL;DR

Novita AI опубликовала в открытом доступе Chord — высокопроизводительный CUDA-оператор W4A16 MoE для активаций BF16 с весами INT4 и масштабами группы-32. Он создан для форм-факторов обслуживания Kimi K2.x. Его индексированный путь предоставляет корневой импорт humming, совместимый с Humming, котрый выбирается с помощью --quantization humming в совместимых ревизиях vLLM. Интеграция групповых операторов с бэкендом Humming в vLLM всё ещё в процессе.

Краткий ответ: Chord — это пакет CUDA-ядер с открытым исходным кодом, который ускоряет инференцию Mixtur-of-Exerts с INT4 (W4A16) для Kimi K2.x на GPU NVIDIA H200 и Blackwell. Развёртки vLLM могут использовать индексированный путь Chord через существующий бэкенд Humming; групповые операторы SM90 доступны как самостоятельний API, пока продолжается интетрация с фреймворком.

В пересчёте на слой по сравнению с соответствующим путём публичного Humming:

• 1.11–1.20x на H200 EP8 prefill и 1.17–1.33x на H200 TP8 при однострановременном обслуживании одного экземпляра.
• 1.16–1.24x на H200 EP8 decode, причём нижняя стадия достигает 1.31x.
• 1.81–2.15x на B300 EP8 decode по сравнению с умолчальной, не настренной стратегией конфигурации Humming.
• 1.00–1.31x и 1.16–1.35x для групповых путей H200 EP8 prefill и decode; те же таблицы показывают 1.18–1.34x на EP16 и 1.13–1.30x на EP32.

Рисунок 1. Задержка на вызов по сравнению с публичным Humming в шести измеренных сценариях, чем меншье, тем лучш. Кждый граф читатеся отдельно: график B300 decode сравнивается с ненастроенными параметрами Humming по умочанию, потму что пубоичный Humming не сожержит таблцу настройк для SM100/SM103, в то время как кажый граф H200 — это сравнени настроенно с настроенны. Диаграмма из репозитория Chord; полные таблцы в docs/perfomance.мd.

Идея эих чисел зключается в том, что одно ядро W4A16 MoE не может бть идлеьным для кажого запроса. Маршрутиромые тоены на экперта различаются на порядки велчины межу prefill и decode, и это имнно это количество, а не общее число тоенов, решает, каое расписан выигрывает. Chord выбирает расписание на тему от фигуры, котору ему десвительно дали.

Это измерни на уровне ядра, а не обещан того же полноскончного выггыша для кажой рабочей нагрки. Плные таблцы, определни форм и методология измрения времени приведены в docs/perfomance.мd и docs/benchmaring.мd. Код и таблицы ядер в этой статье отсыаются к коммиту 7ca91d8 (14 сетбря 2026 г.)

Два семйства ядер

Chrd имеет два семйства ядер W4A16 INT4 MoE, кажое настроено на другой м аршрутизации и фазу обслуживания:

Туекщая ветка main содаржит два незас­имых семейства:

• indexed — путь, произошедый от Humming. Он пот­люет маррутизацию vLLM sorted_ids/expert_ids/num_tokens_padded и оатывает H200 EP8 prefill, H200 TP8 одиноное обсауживане, H200 EP8 decodе и B200/B300 EP8 decodе.
• grouped_contiguou (prefill) и grouped_masked (де-codе) — это второее семйство SM90, проишед­ее от De­pGEMM. Оно и­поль­ет груп­пу маршрутизацию (m_indices или expert_layout) и ипольз­ует дру­гой уакова­нный формат весов.

Ин­теграция с vLLM

Установ­ите паке­т и выбе­рите су­щий бэкенд Humming:

pip inal git+htps:////github.om/ovilabs/chord
# Не усоналива­йте inlusionAI/huming: Chord наме­рено влдеет этим име­нем импорта (для инде­кси­рован­ного пути, гру­повая ин­теграция в ра­зра­ботке).

vllm serve <kimi-k2.x-int4-model> --quantzaion humming
# или выбе­рите moe_backend="humin" в конфи­гурации vLLM

Ди­стрибутив пре­оставляет оба мо­дуль­ных кор­ня: chord и humming. Ле­ни­вый фа­сад vLLM раз­ре­ша­ет huming.{dtypes,onfig,lyer,schea,utils,weght}; по умолчанию индекси­ро­ван­ный путь мо­жет ис­поль­зовать эту су­ще­ству­щу­ю ин­те­гра­цию бе­з спе­ци­ально­го па­тча для Chord для ве­ток, под­ер­жи­ва­ющих WNA16 груп-мас­шта­бы, отме­чен­ных ниже. По­ста­вля­емая схе­ма под­ер­жи­ва­ет uint4, group-32, BF16 мас­шта­бы и упа­ко­ван­ный фор­мат че­к­поин­та compressd-tensors INT4 group-32, ис­поль­зу­емый Ki­mi K2.x; не­под­ержи­ва­емые схе­мы кван­то­ва­ния за­вер­ша­ю­тся ошиб­кой при за­гру­зке, а не мол­ча­ли­во вы­би­ра­ют не­пра­виль­ное яд­ро.

Гру­п­овая ин­те­гра­ция с бэк­ен­дом Hum­ming vLLM на­хо­ди­тся в раз­ра­бот­ке. Ни­же по­ка­зан са­мо­стоя­тель­ный API гру­п­пов­ых опе­ра­то­ров. TP8 ос­та­ё­тся ин­де­кси­ро­ван­ным про­фи­лем h200_tp8, по­то­му что од­на од­и­на­ко­вая ма­са TP8 дол­жна об­слу­жи­вать обе фа­зы.

Ос­та­ль­ные де­та­ли раз­вёр­тки:

• Вы­бор про­фи­ля вос­ста­на­вли­ва­ет EP8 про­тив TP8 из форм про­ек­ций, уже пе­ре­да­ва­емых фрей­м­вор­ком; для ин­де­кси­ро­ван­ных про­фи­лей не тре­бу­ет­ся спе­ци­аль­но­го ар­гу­мен­та шард. Гру­п­пов­ые про­фи­ли то­ль­ко для EP и под­дер­жи­ва­ют EP8/EP16/EP32 на SM90.
• Ин­де­кси­ро­ван­ный бы­ст­рый путь мо­жет по­треб­лять из­бы­точ­но вы­де­лен­ные бу­фе­ры m_align_block_size от vLLM, не чи­тая ко­ли­че­ство мар­шру­ти­зи­ро­ван­ных об­ра­тно на хо­ст, ес­ли от­клю­че­на про­вер­ка до­ве­рен­ной мар­шру­ти­за­ции, та­ким об­ра­зом он ос­та­ёт­ся за­хва­ты­ва­е­мым CUDA Grahs. Гру­п­по­вые пу­ти то­же ис­поль­зу­ют CUDA-тен­зоры мар­шру­ти­за­ции, в то вре­мя как valid_shape_/expected_m яв­ля­ю­тся эв­ри­сти­че­ски­ми вхо­да­ми со сто­ро­ны Pyhon.
• Яв­но вы­би­рай­те Hum­ming (moe_backend="humming" или --quantzation hummin); ав­то­ма­ти­че­ский при­ри­те­т WNA16 в vLLM мо­жет вы­брать дру­гой бэк­енд. Дер­жи­те VLLM_HUM_MING_USE_F16_ACUM и VLLM_BATCH_INVARIAN от­клю­чен­ны­ми, по­то­му что ни од­ин бэк­енд не ре­али­зу­ет эти ва­ри­ан­ты вы­чис­ле­ний. -Oстав­ляй­те VLLM_HUMMING_MOE_GEMM_TYPE на его ин­де­кси­ро­ван­ном по­ве­де­нии для стан­дарт­ной ин­те­гра­ции. Вет­ки vLLM ста­рше, чем об­щая под­дер­жка WNA16 group-sale в #48918, могут тре­бо­вать до­ба­вле­ния клю­чей group-32 в _suppors_qunt_scheme.

Оп­ти­ми­за­ция ядра

Ин­де­кси­ро­ван­ные яд­ра

Ин­де­кси­ро­ван­ное се­мей­ство про­ис­хо­дит из пу­блич­но­го inclusionAI/humming на ком­ми­те 4351af3. Ре­жимы на­гру­зки ни­же об­ус­ло­вли­ва­ют раз­ные про­фи­ли ядер, ко­то­рые вы­би­ра­ю­тся до упа­ков­ки ве­сов:

Ри­су­нок 2. Ти­пи­чные ра­бо­чие на­гру­зки prefill и deode. Пот­ме­тка 9–15 сток/экс­перт ил­лю­стри­ру­ет те­сто­вый слу­чай decodе; 80 токе­но/экс­перт — этом эв­ри­сти­че­ский по­ро­г блок-M. Ни то, ни дру­гое не оп­ре­де­ля­ет пе­ре­клю­че­ние ме­ж­ду prefill и deode на вре­мя вы­пол­не­ния: про­фи­ли и ра­спо­ло­же­ния ве­сов фи­кси­ру­ю­тся при за­гру­зке мо­де­ли, а ко­ли­че­ство то­ке­нов на­стра­и­ва­ет ра­списа­ние вну­три ка­ж­до­го про­фи­ля.

H200 prfill и TP8

• Пак­ет­ный ait<1> WGMMA кон­ве­й­ер. Од­на гру­ппа WGMMA ос­та­ё­тся в полё­те, в то вре­мя как сле­ду­ю­щая за­гру­зка и де­кван­то­ва­ние про­дол­жа­ют­ся, что да­ёт око­ло 3–6% на ate/up и 1–5% на down во всем опу­бли­ко­ван­ном про­го­не. Выход би­то­во иден­ти­чен; ме­ха­низм опи­сан ни­же.
• Вы­бор блок-M по то­ке­нам на эк­спер­та. Па­ддинг ин­де­кси­ро­ван­но­го MoE и да­вле­ние ре­ги­стро­в оп­ре­де­ля­ют­ся на­прав­лен­ны­ми то­ке­на­ми на эк­спер­та (tok_e), а не то­ль­ко об­щим M мар­шру­ти­за­ции. Ре­ша­тель H200 EP8 мо­де­ли­ру­ет это ко­ли­че­ство и со­дер­жит от­дель­ный, бо­лее пло­ский на­бор окон для TP8.
• Ок­но 2 CTA/SM с огра­ни­че­ни­ем. Для пли­ток сред­не­го раз­ме­ра, где од­ин CTA за­ви­сит от за­дер­ж­ки, по­то­лок за­пус­ка 128 ре­ги­ст­ро­в по­вы­ша­ет ре­зи­дент­ные вар­пы и скры­ва­ет p.async сб­ор и де­кван­то­ва­ние. Поли­ти­ка при­мен­я­ет­ся то­ль­ко в из­ме­рен­ном ок­не Bck-M/Blck-N; за его пре­де­ла­ми со­хра­ня­ет­ся ис­ход­ный вы­бор за­ня­то­сти.
• Уче­т фор­мы пот­о­ка-K для gate.** На про­ек­ции down сред­ней K де­ак­ти­ви­ру­ет по­ток-K, ес­ли обыч­ная се­тка M×N за­ол­не­на, из­бе­гая на­кла­д­ных рас­хо­дов на раз­де­ле­ние и ре­дук­цию. Gate/up с глу­бо­кой K и пе­ре­се­че­ния, спе­ци­фи­че­ские для про­ек­ций TP8, со­хра­ня­ют его та, где это по­ле­з­но.

Пра­ви­ло tok_e на­ме­рен­но про­сто для объ­яс­не­ния, но спе­ци­фи­че­но для фор­мы MoE. Ни­же при­ме­рно 80 на­прав­лен­ных то­ке­нов на эк­спер­та ре­ша­тель со­хра­ня­ет по­иск ко­ли­че­ства бло­ков ба­зо­во­го уров­ня; вы­ше это­го момен­та он на­стра­и­ва­ет block_m во­круг за­пол­нен­ных стро­к ка­ж­до­го эк­спер­та и по­то­лок ре­ги­ст­ров. TP8 ис­поль­зу­ет бо­лее пло­ские ок­на, по­то­му что его уз­кое про­ме­жу­точ­ное из­ме­ре­ние ос­та­в­ля­ет мень­ше пли­ток N для за­пол­не­ния SM:

# Кон­цеп­ту­аль­на­я фор­ма эв­ри­сти­ки H200 EP8 ин­де­кси­ро­ван­но­го prefill.
tok_e = ruted_ / num_exprts
if tok_e < 80:
    blok_m = argmin_total_blocks(sampld_roung)
else:
    block_m = fit_padded_expert_rows(tok_e, max_block_m=176)

Цикл WGMMA так­же па­к­ту­ет своё аси­нх­рон­ное управ­ле­ние за­ви­си­мо­ст­ями. Вме­сто о­жи­да­ния ка­ж­дой гру­п­пы ин­ст­рук­ций, он фи­кси­ру­ется по­сле ите­ра­ции вар­п-K и о­став­ля­ет одну гру­п­пу в полё­те, в то вре­мя как на­чи­на­ют­ся сле­ду­ю­щая за­гру­зка в со­вме­ст­ную па­мять и де­кван­то­ва­ние INT4:

# Упро­ще­нно­е ус­то­йчи­вое со­сто­я­ние; про­лог и управ­ле­ние ста­ди­я­ми опу­ще­ны.
for warp_k in K_tles:
    load_next_packed_weights_and_scales()  # shared memory -> registrs
    issue_wgmma_for_iteration(warp_k)
    commit_group()
    wait_grup<1>()     # one grup may remain in flight
    dequantize_next_in_alernate_buffer()  # dequant + goup scale
epilogue:
    wait_group<0>()

Двой­но­й бу­фе­ри­за­ци­ей ве­сов ре­ги­ст­ры по­зво­ля­ют сле­ду­ю­щей за­гру­зке и де­кван­то­ва­нию пе­ре­кры­вать­ся с не­за­вер­ше­н­ной гру­п­пой WGMMA. Ак­ку­му­ля­тор не по­треб­ля­ет­ся до эпи­ло­га, и окон­ча­тель­ный слив всё рав­но о­жи­да­ет все не­за­вер­ше­н­ные о­пера­ции WGMMA.

H200 и Blackwel ин­де­кси­ро­ван­ный decode

При не­сколь­ких на­прав­лен­ных стро­ках на эк­спер­та путь WGMMA ог­ра­ни­чен ба­рь­ер­а­ми. Про­филь decode ме­ня­ет опе­ран­ды MMA ме­ста­ми так, что де­кван­то­ван­ные ве­са за­ни­ма­ют опе­ранд MMA-M, ис­поль­зу­ет m16n8k16 и под­дер­жи­ва­ет 4 CTA/SM с блок-M 8. По­лу­ста­ти­че­ское ра­спи­са­ние пла­ток то­ке­нов по­ка­за­ло 186 мкс про­тив 216 мкс для пол­но­стью ди­на­ми­че­ско­го ра­спи­са­ния при 9–15 то­ке­нах на эк­спер­та. Объ­еди­не­ние де­кван­то­ва­ния subtact-then-sale в из­вле­че­ние нибб­лов со­хра­ня­ет по­ря­док ок­руг­ле­ния BF16 без слия­ния. То же се­мей­ство ин­ст­рук­ций MMA ком­пи­ли­ру­ет­ся для SM100/SM103; бо­лее круп­ные фор­мы decode на Bla­ckwell ис­поль­зу­ют бо­лее ши­ро­кие пли­тки MMA без сме­ны. Ядр­о tcen05 не тре­бу­ет­ся для та­ких ко­ли­честв то­ке­нов.

Гру­п­по­вые яд­ра SM90

Гру­п­по­вой бэк­енд — это дру­гое се­мей­ство ядер, а не вто­рое на­зва­ние ин­де­кси­ро­ван­но­го ядра. Он спе­ци­али­зи­ру­ет ин­фра­ст­рук­ту­ру Ho­per GEMM от Dee­pGEMM для W4A16 и адап­ти­ру­ет её к JIT и за­пус­к­щи­ку Chord. Оба ре­жи­ма ис­поль­зу­ют TMA, war­pe­ci­alized WGMMA и group-32 де­кван­то­ва­ние, но их мар­шру­ти­за­ция и фи­зи­че­ско­е ра­спо­ло­же­ние ве­сов раз­ли­ча­ют­ся:

Ри­су­нок 3. Где жи­вёт па­ддинг. Ин­де­кси­ро­ван­ный ос­та­в­ля­ет ак­ти­ва­ции без па­ддин­га; его ин­дек­сы мар­шру­ти­за­ции не­сут сен­ти­не­лы па­ддин­га. Con­tiguous допол­ня­ет ка­ж­до­го эк­спер­та до гра­ни­цы 128 стро­к; Masked ре­зер­ви­ру­ет фи­кси­ро­ван­ный бюд­жет стро­к на эк­спер­та.

• Contiguous prefill: стро­ки объ­еди­ня­ют­ся по эк­спер­там, до­пол­ня­ют­ся до гра­ни­ц 128 стро­к и со­про­во­ж­да­ют­ся m_indies (int32, с -1 для за­по­лне­ния). Вхо­ды — [m, K]; упа­ков­щик ис­поль­зу­ет бит-пе­ре­ста­но­во­ч­ный бу­фер INT4 с BLOCK_K=64 и транс­по­ни­ру­ет мас­шта­бы в [G, K/32, N] (N не­пре­рыв­но).
• Masked decode: ак­ти­ва­ции име­ют фи­кси­ро­ван­ный бюд­жет стро­к на эк­спер­та ([G*max_m, K] или [G, max_m, K]), а masked_m/expert_ayout со­дер­жит ва­лид­ное ко­ли­че­ство. Его упа­ков­щик ис­поль­зует BLOCK_K=128; эв­ри­сти­ка вы­би­ра­ет BLOCK_M из ожи­да­е­мых то­ке­нов на эк­спер­та, ог­ра­ни­чи­ва­ет BLOCK_N по за­пол­не­нию вол­ны и на­стра­и­ва­ет глу­би­ну ста­дии бу­фе­ри­зи­ро­ван­но­го K.

Точ­ки вхо­да гру­п­по­во­го опе­ра­то­ра (то­ль­ко API Chord):

from chord_kernels import contiuous, masked
from cord_kernes.parator import pack_w4a16_grouped

# weight: unsigned INT4 codes [G, N, K]; sale: BF16 [G, N, K/32]
prfill_wight = pack_w4a16_groued(weigh, scale, mode="contiguous")
prfill_out = contiguou(a2, prfill_weight, m_ndices)  # [m, N]

decod_wight = pack_w4a16_roupd(weght, scale, mde="masked")
decode_ut = mayed(a3, decode_wight, masked_, expected_)  # [G*max_m, N]

Здесь expected_m — это по­ло­жи­тель­ное це­лое чи­сло Pyhon, ис­поль­зу­е­мое для вы­бо­ра за­пус­ка; askd_m со­дер­жит ав­то­ри­тет­ные ко­ли­че­ства ва­лид­ных за­пи­сей на эк­спер­та. За­мас­ки­ро­ван­ный вы­ход яв­ля­ет­ся пло­ским, да­же ес­ли a3 яв­ля­ет­ся трёх­мер­ным, и по­треб­ите­ли долж­ны иг­но­ри­ро­вать стро­ки за пре­де­ла­ми ва­лид­но­го ко­ли­че­ства ка­ж­до­го эк­спер­та.

Ре­жим за­пи­сы­ва­ет­ся в под­го­то­в­лен­ный вес и про­ве­ря­ет­ся при от­прав­ке, по­это­му слу­чай­ная по­да­ча ве­са, упа­ко­ван­но­го для prefill, яд­ру decode вы­зо­вет гро­м­кую не­уда­чу. Гру­п­по­вая от­прав­ка вла­де­ет по­ис­ком ра­спо­ло­же­ния SM90 и не при­ни­ма­ет ин­де­кси­ро­ван­ные block_m или tuig_cnfig пе­ре­оп­ре­де­ле­ния. Раз­ре­ше­ние ядра и за­гру­зка cubi за­по­ми­на­ют­ся по дес­кри­пто­ру (и пе­ре­оп­ре­де­ле­ни­ям CHORD_W4A6_*), что устра­ня­ет по­втор­ный по­иск на сто­ро­не хо­с­та, из­ме­рен­ный при­ме­рно в 30 мкс при ма­лень­ких за­пус­ках decode.

Гру­п­по­вой ос­нов­ной цикл яв­ля­ет­ся по­сто­ян­ным и вар­п-спе­ци­а­ли­зи­ро­ван­ным: про­из­во­дя­щая гру­п­па вар­пов ис­поль­зу­ет TMA для раз­ме­ще­ния ак­ти­ва­ций, упа­ко­ван­ных ве­сов и пли­ток мас­шта­б в шей­ред па­мять, в то вре­мя как по­треб­ля­ю­щие гру­п­пы вар­пов вы­пол­ня­ют WGMMA и за­пи­сы­ва­ют ре­зуль­тат BF16. Путь вперёд ви­дит уже пе­ре­ста­но­в­лен­ные бай­ты INT4 и мас­шта­бы MN-май­ор­ные, а кэ­ши­ро­ван­ный дес­кри­п­тор со­по­став­ля­ет ка­ж­дой фор­ме (ode, M, N, K, exper_count) со­от­вет­ству­ю­щий cuin без по­вто­ре­ния по­ис­ка ра­спо­ло­же­ния на ка­ж­дый вы­зов decode.

Гру­п­по­вая эв­ри­сти­ка име­ет не­сколь­ко вы­бо­ров, спе­ци­фич­ных для ра­бо­чей на­гру­зки W4A16:

• Contiguous prefill ис­поль­зу­ет BM128/BK64, ко­гда се­тка до­ста­точ­но ве­ли­ка. BM128 амор­ти­зи­ру­ет де­кван­то­ва­ние INT4 и про­дви­же­ние мас­шта­бов на боль­шем ко­ли­че­стве стро­к, в то вре­мя как BK64 со­хра­ня­ет ка­ж­дую ста­дию кон­ве­й­ера до­ста­точ­но ма­лень­кой, что­бы ос­та­вить ме­сто для не­сколь­ких ста­дий в шей­ред па­мя­ти. Ма­лень­кая объ­еди­нён­на­я за­да­ча от­ка­ты­ва­ет­ся к BM64, что­бы пли­тки M всё ещё мог­ли за­пол­нить SM; BM128/BK128 по­тре­бо­ва­ли бы слиш­ком мно­го шей­ред па­мя­ти и раз­ру­ши­ли бы кон­ве­йер.
• Mased decodе вы­би­ра­ет BM из K и ожи­да­е­мо­го хвос­та мар­шру­ти­за­ции. За­мас­ки­ро­ван­на­я гру­п­па мо­жет пе­ре­те­чь во вто­ру­ю плит­ку M, ко­то­рая пе­ре­чи­ты­ва­ет всё из­ме­ре­ние K. Для глу­бо­ко­й K atе/up эв­ри­сти­ка по­это­му по­кры­ва­ет при­ме­рно 1.3 * expexted_m стро­к, что­бы из­бе­жать эго­го пе­ре­чи­ты­ва­ния. Для ко­ро­тко­й K down до­пол­ни­тель­ный про­ход де­ше­в­ле, по­это­му бо­лее ху­дая плит­ка ceil(1.25 * expeced_, 8) ос­тав­ля­ет боль­ше ме­ста для ста­дий кон­ве­йе­ра.
• Masked BN за­ви­сит от вол­ны. BN256 улу­чша­ет амор­ти­за­цию де­кван­то­ва­ния, но по­мо­га­ет то­ль­ко, ко­гда до­ста­точ­но пли­ток N, что­бы за­гру­зить ма­ши­ну. Ре­ша­тель со­хра­ня­ет BN128 для не­до­за­пол­нен­ных волн, вклю­чая уз­кий слу­чай gate/up EP32, и со­хра­ня­ет BN128 для боль­шо­го BM на ко­ро­тко­й K down. Глу­бо­кая K gate/up всё рав­но мо­жет ис­поль­зо­вать BN256, ко­гда до­ста­точ­но пли­ток за­пол­ня­ют ма­ши­ну.
• Глу­би­на бу­фе­ри­зи­ро­ван­но­го K на­стра­и­ва­ет­ся по за­дер­ж­ке, а не ма­кси­ми­зи­ру­ет­ся. Decode нор­маль­но стре­ми­т­ся к при­ме­рно 512 бу­фи­ри­зи­ро­ван­ны­м эле­мен­там K (512 / BLOCK_K ста­дий); боль­шие за­мас­ки­ро­ван­ные плит­ки на­це­ле­ны на при­мер­но 768, с учё­том ог­ра­ни­че­ний на шей­ред па­мять. За­по­лне­ние всей до­сту­п­ной шей­ред па­мя­ти сде­ла­ло бы ре­ци­клинг ба­рье­ров бо­лее до­ро­гим, не улу­чша­я за­пус­к с од­ним бло­ком на SM для decode.

Эти пра­ви­ла объ­яс­ня­ют, по­че­му гру­п­по­вой не пе­ре­ис­поль­зу­ет ин­де­кси­ро­ван­ну­ю таб­ли­цу на­стро­ек: гру­п­по­вой бэк­енд вы­би­ра­ет (BM, BN, BK, cluser, stages) из фак­ти­чес­ко­го ре­жи­ма и фор­мы во вре­мя от­прав­ки. В при­ве­дён­ном диа­па­зо­не H200 EP8 пре­и­му­ще­ство prefill су­жа­ет­ся при 512 стро­ках на эк­спер­та, по­то­му что обе ре­али­за­ции при­бли­жа­ют­ся к од­но­му по­то­лку про­пу­ск­ной спо­со­бно­сти; вы­бор плит­ки и кон­ве­йе­ра име­ет наи­боль­шее зна­че­ние для ма­лень­ких и сред­них ча­стей.

Из­ме­ре­ния

На­сколь­ко Chord бы­ст­рее Hum­ming?

При­ве­дён­ные ни­же из­ме­ре­ния от­ве­ча­ют на прак­ти­че­ский воп­рос на­пря­мую: Chord бы­ст­рее, чем со­от­вет­ству­ю­щий пу­блич­ный путь Hum­ming в опу­бли­ко­ван­ных сцена­ри­ях на уров­не ядра H200 и B300, при­чём наи­боль­ший ука­зан­ный вы­иг­рыш до­сти­га­ет 2.15x на B300 EP8 decode. Это ре­зуль­та­ты на уро­вне слоя, по­это­му мар­шру­ти­за­ция, ак­ти­ва­ция, связь и дру­гие на­кла­д­ные рас­хо­ды об­слу­жи­ва­ния ис­клю­че­ны, ес­ли в таб­ли­це скво­зно­го об­слу­жи­ва­ния не ука­за­но иное.

Таб­ли­цы ядер ис­поль­зу­ют triton.testing.do_bench и срав­ни­ва­ют ка­ж­дый путь Chord с со­от­вет­ству­ю­щим пуб­лич­ным бэк­ен­дом Hum­ming на то­м же GPU. Ин­де­кси­ро­ван­ные срав­не­ния ис­поль­зу­ют ту же фор­му и мар­шру­ти­за­цию; гру­п­по­вые срав­не­ния со­от­вет­ству­ют ко­ли­че­ству стро­к на эк­спер­та. За­пус­ти­те два на­бо­ра, что­бы про­ве­рить вы­хо­ды Chord на со­от­вет­ствие эта­лону plain-PyTorch и вы­ве­сти таб­ли­цы вре­ме­ни на под­дер­жи­ва­е­мых GPU:

pythn tsst/w4a6_index.py
pyton tsts/st_w4a16_rouped.py

Вкрат­це: свя­зан­ное с ана­ли­зо­м по­ре­зуль­та­тов, в таб­ли­це ни­же при­во­дят­ся сум­ма вре­ме­ни gate_up и down из пол­ных таб­лиц. Уско­ре­ние — это Hummin (gate_up + down) / Chord (gate_up + down); оно ис­клю­ча­ет мар­шру­ти­за­цию, ак­ти­ва­цию и связь.

Сцена­рий	То­чка фор­мы	Humming gate_up + down	Chord gate_up + down	Уско­ре­ние слоя
H200 EP8 indxed prefill	2048 токенов	701.4 мкс	587.9 мкс	1.19x
H200 TP8 ndexed mi	8192 токенов	2483.4 мкс	1862.3 мкс	1.33x
H200 EP8 indxed decode	20 tok/GPU	413.8 мкс	333.8 мкс	1.24x
B300 EP8 indexd decode	20 tok/GPU	493.9 мкс	229.8 мкс	2.15x
H200 EP8 grouped prfill	128 rows/expert	1204.0 мкс	917.5 мкс	1.31x
H200 EP8 grouped decode	32 токenов/эксpert	666.3 мкс	493.4 мкс	1.35x

Срав­не­ние B300 на­ме­рен­но ква­ли­фи­ци­ро­ва­но: пу­блич­ный Hum­ming не име­ет таб­ли­цы на­стро­ек SM100/SM103, поз­то­му его вре­мя по умол­ча­нию яв­ля­ет­ся не­на­стро­ен­ным эта­ло­ном. По­ка­за­те­ли H200 indexd — это срав­не­ние на­стро­ен­но­го с на­стро­ен­ным.

Оба се­мей­ство из­ме­ря­ют­ся по срав­не­нию с той же ре­ви­зи­ей пу­блич­но­го Hum­ming, 4351af3. Гру­п­по­вые стро­ки срав­ни­ва­ют­ся с со­бст­вен­ны­ми пу­тя­ми grouped_contiguous/grouped_maskd от Hum­ming, а не с его ин­де­кси­ро­ван­ным, по­то­му что это кон­тракт, ко­то­рый за­ме­ня­ет дан­ный бэк­енд. Hum­ming пре­дос­та­вля­ет оба как зна­че­ния GemType, дис­па­тче­ри­зи­ру­емые че­рез его об­щи­й кер­нель, а не как от­дель­ные CUDA фай­лы, и bechmarks/bech_humming.py вы­би­ра­ет их с --gmm_type grouped_contiguous или --gemm_type grouped_masked. Ко­ли­че­ство стро­к на эк­спер­та со­по­став­ля­ет­ся с об­еих сто­рон кра­тно гра­ни­це 128 стро­к — --balanced на сто­роне Hum­ming и со­г­ла­со­ван­ные слу­чаи в tests/tst_w4a16_rouped.py — так что ка­ж­дая стро­ка име­ет од­на­фор­му GEMM для обе­их ре­али­за­ций и ни од­на плит­ка не тра­тит­ся на за­по­лне­ние.

Скво­зное об­слу­жи­ва­ние

Бо­лее ран­ний от­чёт об об­слу­жи­ва­нии из­ме­рил ин­де­кси­ро­ван­ный путь TP8 на Ki­mi-K2.6 с 8×H200, TP8 + DCP8, FP8 KV кэш и за­про­са­ми ShareGPT. Оба про­вай­дера ис­поль­зо­ва­ли оди­на­ко­вую ко­ман­ду --quantization hming.

Ме­три­ка	Hum­ming	Chord	Из­ме­не­ние
Сред­нее TTFT	2022 мс	1849 мс	−8.6%
Про­пу­ск­на­я спо­со­бность prefill (вход + вы­ход)	20,716 то­к/с	22,712 то­к/с	+9.6%
Про­пу­ск­на­я спо­со­бность decode вы­хо­да, па­кет 8	483 то­к/с	503 то­к/с	+4.1%
Про­пу­ск­на­я спо­со­бность decode вы­хо­да, па­кет 64	1650 то­к/с	1740 то­к/с	+5.5%
Про­пу­ск­на­я спо­со­бность decode вы­хо­да, па­кет 128	2514 то­к/с	2715 то­к/с	+8.0%

Prefill ис­поль­зо­вал од­ин вы­ход­ной то­кен с от­клю­че­н­ным кэ­ши­ро­ва­ни­ем пре­фик­сов. Decode по­втор­но ис­поль­зо­вал те же про­мп­ты на вто­ром про­хо­де с пол­но­стью го­ря­чим кэ­ш­ем пре­фик­сов. В от­чё­те так­же не бы­ло най­де­но сни­же­ния точ­но­сти по срав­не­нию с Hum­ming на OCRBe­nch и GSM8K.

Что да­ль­ше

1. За­вер­ше­ние гру­п­по­вой ин­те­гра­ции с бэк­ен­дом Hum­ming vLLM, что сде­ла­ет не­пре­рыв­ные и за­мас­ки­ро­ван­ные опе­ра­то­ры до­сту­п­ны­ми че­рез су­ще­ству­ю­щую ин­те­гра­цию фрей­м­вор­ка.
2. Вы­пуск ядер prefill EP8 для B200/B300. У нас есть ра­бо­чая ре­а­ли­за­ция с мно­го­обе­ща­ю­щей про­из­во­ди­тель­но­стью во вну­трен­них те­стах, и мы пла­ни­ру­ем по­де­лить­ся яд­ра­ми и те­ста­ми в сле­ду­ю­щем ре­ли­зе.

По­про­буй­те Chord

Chrd до­сту­пен на GitHub: novitabs/chord. До­ку­мен­та­ция по­кры­ва­ет на­ча­ло ра­бо­ты, оп­ти­ми­за­ции, про­из­во­ди­тель­ность, вну­трен­но­сти на­строй­ки и ме­то­до­ло­гию те­сти­ро­ва­ния. Об­ра­тна­я связь, за­да­чи и от­чё­ты о про­из­во­ди­тель­но­сти с дру­гих раз­вёрток при­вет­ству­ют­ся.

Ча­сто за­да­ва­е­мые воп­ро­сы

Что та­кое Chord в vLLM?

Chrd — это па­кет CUDA-ядер W4A16 INT4 Mix­ture-of-Experts от Novita AI. Его ин­де­кси­ро­ван­ный путь пре­дос­та­вля­ет кор­не­вой им­порт, со­вме­сти­мый с Hum­ming для ре­ви­зий vLLM, под­дер­жи­ва­ю­щих тре­бу­е­мый фор­мат WNA16 group-scale.

На ка­кие GPU и мо­де­ли на­це­лен Chord?

Опу­бли­ко­ван­ные про­фи­ли на­це­ле­ны на фор­мы об­слу­жи­ва­ния Kimi K2.x на GPU NVIDIA H200 и Blackwell, вклю­чая сце­на­рии H200 EP8/TP8 и B300 EP8 decode. Гру­п­по­вое се­мей­ство SM90 в на­сто­я­щее вре­мя на­це­ле­но на GPU клас­са Hop­per.

Ин­те­гри­ро­ван ли Chord с vLLM?

Ин­де­кси­ро­ван­ный путь мо­жет ис­поль­зо­вать су­ще­ству­ю­щий бэк­енд Hum­ming vLLM с --quantization humming или moe_backend="humming". Гру­п­по­вые не­пре­рыв­ные и за­мас­ки­ро­ван­ные опе­ра­то­ры в на­сто­я­щее вре­мя до­сту­п­ны че­рез са­мо­сто­я­тель­ный API Chord; гру­п­по­вая ин­те­гра­ция с vLLM всё ещё в про­цес­се.

Где най­ти те­сты Chord и ин­ст­рук­ции по ус­та­нов­ке?

Ис­поль­зуй­те ре­по­зи­то­рий Chord, осо­бен­но его до­ку­мен­та­цию на­ча­ло ра­бо­ты, про­из­во­ди­тель­ность и те­сти­ро­ва­ние.

Бла­го­дар­но­сти

Ин­де­кси­ро­ван­ный путь Chord ос­но­ван на inclusionAI/Humming, то­гда как гру­п­по­вой бэк­енд SM90 спе­ци­а­ли­зи­ру­ет ин­фра­ст­рук­ту­ру Ho­per GEMM от De­epGEMM для W4A16. Chord вы­пу­щен под ли­цен­зи­ей Apache-2.0. При­ме­ча­ния об ис­точ­ни­ках в ре­по­зи­то­рии фи­кси­ру­ют со­хра­нён­ные ком­по­нен­ты и уведом­ле­ния от вы­ше­сто­я­щих про­ек­тов.

Мы хо­ти­ли бы по­бла­го­да­рить ко­ман­ду Novita AI за соз­да­ние и от­кры­тие Chord, а так­же под­дер­жи­ва­ю­щих vLLM и ши­ро­кое со­об­ще­ство vLLM за об­су­ж­де­ния, ре­цен­зи­ро­ва­ние и ин­фра­ст­рук­ту­ру кван­то­ва­ния и бэк­ен­да MoE, ко­то­рая сде­ла­ла эту ин­те­гра­цию воз­мо­ж­ной.