Компания Cursor выложила в опенсорс Mixture-of-Kittens (MoK) — «мегаядро» для обучения MoE-моделей на стойках Nvidia GB300 NVL72. Вся пересылка данных между 72 видеокартами и все вычисления MoE-слоя в нем выполняются одним гигантским GPU-ядром. По замерам самой компании, MoK работает до 2.37 раза быстрее лучших публичных решений, а в продакшене поднял сквозную скорость обучения на 41%: с 761 до 1070 токенов в секунду на видеокарту. На этом ядре Cursor уже обучает свою модель Composer на десятках тысяч GPU.
MoE-слой — та часть трансформера, где токены разлетаются по сотням «экспертов», — в зависимости от конфигурации мог съедать больше половины всего времени обучения Composer. Cursor целый год боролся с этим по частям: написал собственные ядра для обучения в форматах MXFP8 и NVFP4, придумал технику warp decode для инференса. Но все эти оптимизации касались только вычислений, а узким местом к тому моменту стала пересылка данных между видеокартами. Пришлось переписать весь MoE-слой с нуля — так, чтобы коммуникация была встроена прямо в ядро.
Свою роль сыграло и железо. NVL72 — это стойка, в которой 72 GPU объединены в единый NVLink-домен и могут очень быстро обмениваться данными напрямую. Но у нее нашлась и слабость: встроенные ARM-процессоры Grace (буква «G» в GB300) оказались настолько медленными относительно видеокарт, что GPU регулярно простаивали, дожидаясь, пока CPU закончит свою часть работы — вплоть до банальной отправки метрик обучения на сервер. На машинах с процессорами Intel тот же стек таких проблем не создавал. Вывод инженеры сделали радикальный: выкинуть CPU из цикла вообще и держать всю работу на видеокартах.
Чтобы понять, что именно ускорял Cursor, нужен минимальный ликбез. В MoE-модели для каждого токена роутер выбирает несколько «экспертов» из сотен доступных. Веса экспертов не помещаются на одну видеокарту, поэтому их раскладывают по десяткам GPU — и каждый токен приходится сначала переслать на карты с нужными экспертами (операция dispatch), там обсчитать, а потом вернуть результаты на исходную карту (combine). Пересылка занимает сравнимое с вычислениями время, поэтому весь фокус в том, чтобы счет и передача данных шли параллельно и не ждали друг друга.
Первое оригинальное решение MoK — направление пересылки. Индустриальный стандарт, включая библиотеку DeepEP от DeepSeek, использует push: видеокарта-владелец токенов сама «заталкивает» их получателям. Cursor показал, что на NVL72 выгоднее наоборот — pull, когда карта-получатель сама «вытягивает» нужные токены. Расписание пересылок при этом упрощается до двух колонок вместо трех и строится без сортировок, а главное — исчезают сигналы о завершении между картами: не нужно ждать подтверждений от 71 соседа. В микробенчмарках такие сигналы при push-подходе оказались в 5.8 раза медленнее (103 микросекунды против 18), а pull дал до 29% прироста утилизации NVLink при неравномерной нагрузке на экспертов.
Второй трюк — кольцевой буфер токенов. Проблема MoE в том, что заранее неизвестно, сколько токенов прилетит на конкретную карту: либо синхронизируйся с медленным CPU, чтобы под каждый шаг выделять буферы точного размера, либо выбрасывай не поместившиеся токены, жертвуя качеством обучения. MoK вместо гигантского буфера «на всякий случай» гоняет токены через фиксированное кольцо в несколько сотен мегабайт: пересылка заполняет слоты, вычисления их освобождают, и все это без единого обращения к процессору. Есть и изящная деталь: при прямом проходе кольцо обходится в обратном порядке — так к его концу в буфере остаются как раз те активации, которые понадобятся для расчета градиентов, и пересчитывать приходится меньше.
Все это упаковано в «мегаядро»: вместо череды отдельных GPU-ядер с накладными расходами на каждый запуск работает одно, а задачи внутри него программно раскидываются по вычислительным блокам GPU — потоковым мультипроцессорам. Часть блоков занимается вычислениями, часть — пересылкой, и они сигналят друг другу через локальный счетчик. Бонусом ядро дает полный детерминизм: одинаковый вход дает бит-в-бит одинаковый выход независимо от того, как железо распланирует исполнение инструкций, — критичное свойство для RL-дообучения и воспроизводимых экспериментов.
Теперь о цифрах. Cursor сравнил MoK с четырьмя публичными стеками, включая HybridEP + Megatron — связку, которую Nvidia рекомендует для NVL72 вместо варианта с DeepEP. Замеры шли на конфигурациях реальных открытых моделей: Kimi K2.7 Code, GLM-5.2, Qwen3.5-397B-A17B и DeepSeek-V4-Pro. Максимальный отрыв — 2.37 раза на прямом проходе в MXFP8, на обратном — до 1.78. А в примере расчета параметров мимоходом названа база Composer 2.5 — китайская Kimi 2.5. Весной зависимость Composer от Kimi была скандалом: Cursor не раскрыл базу при запуске, пользователи вычислили ее по идентификатору в API, и сооснователю компании пришлось признавать ошибку. Теперь это просто строчка в техническом посте — и заодно подтверждение, что на китайском опенсорсе построено и новое поколение модели.
Самая же симпатичная деталь поста — о том, кто все это писал. Название Mixture-of-Kittens — реверанс фреймворку ThunderKittens из стэнфордской лаборатории Hazy Research Криса Ре: ведущий автор MoK Стюарт Сул пишет у Ре PhD и параллельно работает в Cursor. Так вот, по словам авторов, простые однооператорные ядра у них теперь почти целиком пишет ИИ, а агенты, которым задали верные архитектурные ориентиры, выдают ядра уровня state-of-the-art с первой попытки. Еще год назад для написания мегаядра требовался специальный фреймворк-прослойка, упрощающий работу человеку, — в этот раз команда обошлась без него и разобрала всю сложность напрямую, вместе с агентами. Получается ядро для обучения ИИ, в значительной части написанное ИИ.
Все бенчмарки пока — от самого Cursor, независимых прогонов за первые часы не появилось. Ядро жестко заточено под GB300 NVL72 — стойку стоимостью в миллионы долларов, которая есть у считаных компаний, так что декларируемое «снижение барьера для ИИ-исследований» звучит несколько иронично. И все же жест сильный: код, алгоритмы планирования и бенчмарк-скрипты открыты полностью, проверить заявленные цифры может любой обладатель подходящего железа, а идеи вроде pull-пересылки и кольцевых буферов переносимы и на другие платформы — авторы прямо пишут, что писали код так, чтобы его было легко перенести на другое железо. С помощью агентов, разумеется.
P.S. Поддержать меня можно подпиской на канал «сбежавшая нейросеть», где я рассказываю про ИИ с творческой стороны.
Загрузка...
