Когда мы создавали Fast Apply, мы поняли, что в открытом стеке для инференса не хватает ключевого компонента: специализированного спекулятора для конкретной задачи. Обычный черновик дает ускорение в 1,93 раза, а специализированный — в 100 раз. Опираясь на это, мы обучили WarpGrep и некоторые из лучших специализированных моделей. Этот механизм является тем же ядром, что и в спекулятивном декодировании.
Малые модели не подчиняются тем же законам масштабирования, что и большие. Мы обучаем кастомные спекуляторы для генерации кода, чтобы они были лучшими в своем классе. Мы обслуживаем открытые модели — Qwen, GLM, DeepSeek, MiniMax — для одной рабочей нагрузки: агента-программиста. Чтобы сделать их быстрыми, нужно выполнить три вещи, которые открытый стек не сделает за вас.
- Кастомный спекулятор. Черновик, обученный на собственном коде модели, а не на данных из интернета. Обычный черновик: 1,93x. Обученный на целевой модели: 3,07x.
- Ядра, оптимизированные для задержки или пропускной способности. Декодирование в FP4, которое остается 4-битным вплоть до тензорного ядра, автоматически подбирается на дешевых GPU, которые никто другой не настраивает. От 97 до 162 токенов/с на карте за 7000 долларов.
- Межсоединение, оптимизированное для высокого коэффициента попаданий в кэш. Однопроходная операция all-reduce через PCIe и префиксный кэш, который работает через обычный TCP там, где нет NVLink.
Каждый из этих пунктов — это место, где общий стек останавливается, а мы идем дальше.
1. Инференс, который становится быстрее по мере использования
Спекулятивное декодирование: небольшая модель-черновик угадывает следующие несколько токенов, целевая модель проверяет их за один проход, и вы продолжаете выполнение до первого промаха. Одно число решает всё: коэффициент принятия, то есть как часто целевая модель соглашается с догадкой.
Обычный черновик — плохой предсказатель. На Vicuna-13B стандартный черновик на 68 млн параметров дает 1.93x; a draft trained on the target's own output gets 3.07x, при той же целевой модели и тех же настройках. Этот разрыв и есть суть проблемы.
Архитектуры публичны и хороши. позволяет черновику обучаться на «сырых» данных, вместо того чтобы копировать признаки целевой модели, и длина принятия вырастает с 3,96 до 6,62. DFlash, Spec V2 от SGLang с июня 2026 года, составляет черновик целого блока за один проход: более чем 6-кратное ускорение без потерь, 3,2-кратное на HumanEval, где EAGLE-3 дает 2,2-кратное.
Но архитектура — это пустой разъем. Никто не даст вам черновик, обученный именно на вашей целевой модели для вашей нагрузки. Вы либо обучаете его сами, либо запускаете стандартный и получаете ускорение в 1,93 раза.
Обучение хорошего черновика — это обучение малых моделей. Fast Apply и Compact сделали нас одной из лучших команд в мире в этой области. Вещь, которую узнаешь при работе с моделями менее 30 млрд параметров: законы фронтирного масштабирования перестают работать. Chinchilla говорит, что ~20 токенов на параметр — это вычислительно оптимально, но это предполагает, что обучение — основная статья расходов. Для модели, которую вы обучаете один раз и используете миллиарды раз, это не так, и оптимум сильно смещается в сторону маленьких и переобученных моделей.
Спекулятор живет именно там. Маленький, переобученный, заточенный под одно распределение.
Поэтому мы обучаем по одному спекулятору на каждую открытую модель, используя выходные данные кода, а не веб-тексты. Сгенерированный код повторно использует шаблоны и символы, уже имеющиеся на экране, а правка — это в основном копия файла, который она редактирует. Черновик, прочитавший миллион diff-файлов, предсказывает эти токены. Тот, что читал интернет — нет, вот почему code is the highest-speedup task for every speculation method. Для Fast Apply этот специализированный спекулятор позволяет достичь 10 500 токенов/с при применении и 33 000 при компактизации. Те же веса Qwen, которые можно скачать. Наша версия быстрее, потому что спекулятор, работающий с ней, был обучен нами на реальных задачах.
2. Ядра, оптимизированные для задержки или пропускной способности
Декодирование ограничено памятью. При размере батча 1 матричное умножение превращается в GEMV: оно считывает каждый вес один раз на токен и выполняет почти нулевые вычисления на байт, поэтому токены в секунду — это пропускная способность, деленная на количество затронутых байтов. Единственный рычаг — это байты. Храните веса в FP4 (e2m1), и вы перемещаете 4 бита там, где bf16 перемещает 16, что дает 4-кратное сокращение трафика HBM. Так что вся задача сводится к тому, чтобы держать их в 4-битном формате до самого тензорного ядра.
Стандартные ядра сдаются прямо здесь. Они деквантуют FP4 в bf16 в HBM, возвращают 4-кратный объем данных и вызывают обычный GEMM. Сделайте это правильно, и вы вообще не будете материализовать bf16. Вы передаете упакованные веса через cp.async/TMA, распаковываете e2m1 и масштаб блока 16 в регистрах и подаете их на тензорные ядра Blackwell FP4 прямо из общей памяти. На 80B MoE это сгруппированный GEMM только по активным экспертам плюс постоянное ядро, чтобы каждый SM оставался резидентным при батче 1, плюс split-K, чтобы заполнить машину, которую GEMV не может насытить в одиночку. Упустите хоть что-то из этого, и вы вернетесь к ограничению bf16.
Настройка этого для разных архитектур — жестокая задача, а стандартные настройки написаны для карт, которые покупают фронтирные лаборатории. Перенесите ядро без перенастройки, и оно будет работать на 7% of optimal; достижение современного уровня на AMD MI250 потребовало переписывания 40% ядра flash-attention вручную. Поэтому мы не пишем пространство поиска вручную. Ядро проверяемо, оно либо корректно относительно эталонного вывода, либо нет, и это делает его поиском, который можно автоматизировать.
Система предлагает ядро, проверяет его на производственных трассах, тестирует на низкопроизводительных картах NVIDIA и AMD, которые никто другой не настраивает, и выпускает победителей. KernelBench показывает, почему это должно быть автоматизировано: при оценке по критерию «корректно и быстрее» фронтирные модели проходят менее 20% задач «на холодную». Большой объем через жесткий цикл проверки — единственный путь.
Один результат: наши ядра warp-decode запускают 80B MoE на скорости 162 токена/с на карте за 7000 долларов, по сравнению с 97, обгоняя 120 токенов/с на H100 за 25 000 долларов. Никакой потери точности, код открыт. Это окупается только потому, что вычислительные мощности в дефиците, что создало цену на те пробелы, которые проигнорировал общий стек.
3. Межсоединение, оптимизированное для высокого коэффициента попаданий в кэш
Программный трафик shares 97% of its prefix tokens, with prompts 37x to 2,494x longer than the outputs. Кэшируйте префикс, и следующий запрос будет платить только за новые токены. Коэффициент попаданий — это цена. Абстракция кэша открыта, и мы ее используем: RadixAttention хранит префиксы в дереве, маршрутизатор с учетом кэша повышает коэффициент попаданий с 20% до 75%, HiCache выгружает дерево в оперативную память хоста и удаленное хранилище, а на Qwen3-Coder-480B повышает коэффициент попаданий с 40% до 80% и удваивает пропускную способность.
У дешевых GPU есть подвох. Нет NVLink.
NVLink перемещает 900 GB/s between GPUs. PCIe Gen5, шина на доступных серверах, перемещает 64 ГБ/с в одном направлении. В 14 раз меньше. Это незаметно, пока вы не шардируете модель между GPU, а потом это становится всем: тензорный параллелизм запускает all-reduce после внимания и MLP каждого слоя, и этот all-reduce стоит 8-11% of the step on NVLink and 40-75% on PCIe. Нет быстрой фабрики, и коммуникация съедает большую часть прямого прохода.
Стандартное решение — купить NVLink. Мы написали ядро вместо этого.
Кольцевой all-reduce в NCCL создан для больших полезных нагрузок на быстрой фабрике. При батче 1 сообщения составляют несколько сотен КБ и ограничены задержкой, поэтому 2(N-1) переходов кольца доминируют. Мы запускаем однопроходный all-reduce: каждый GPU отправляет свой шард прямо в память соседа через PCIe P2P и выполняет редукцию локально, за один круговой путь. Мы запускаем это на отдельном потоке, чтобы оно перекрывалось с GEMM следующего слоя, объединяем редукцию в эпилог проекции вниз, а не запускаем второе ядро, и квантуем полезную нагрузку до FP8, чтобы вдвое уменьшить объем данных, проходящих через шину. Графы CUDA стирают накладные расходы на запуск для каждого слоя, которые в противном случае платил бы коллектив. Это скрывает большую часть 14-кратного разрыва. Другое ядро — это префиксный кэш, который работает между машинами через обычный TCP.
HiCache уже определяет удаленный уровень L3 за бэкендом, который является three functions: get, exist, set. Он работает поверх любого транспорта. Подвох в том, что заявленные преимущества достигаются через RDMA, где передача занимает менее миллисекунды и стоит under 0.1% of request latency. Обычный TCP на порядок медленнее. На системе, использующей только PCIe, открытый стек тихо выходит из строя, потому что приведенные цифры предполагали наличие оборудования, которого у вас нет.
Таким образом, преимущество TCP не может быть обеспечено транспортом. Оно достигается за счет коэффициента попаданий. Обученный спекулятор и автоисследованные ядра повышают этот коэффициент настолько, что префикс, который отсутствует в GPU и оперативной памяти хоста, извлекается от соседа по TCP вместо повторного вычисления, а пропуск предварительного заполнения оказывается быстрее медленной выборки. По сравнению с полным пересчетом, такая выборка сокращает время до получения первого токена .
Быструю фабрику, которую все покупают, чтобы избежать этого, мы заменили ядрами. Мы запускаем GPU, от которых рынок отказался, с коэффициентами попаданий, которые, как предполагается, требуют оборудования, которое мы не покупали.
Одна рабочая нагрузка
Три вещи, один цикл:
- Спекулятор составляет собственный программный код модели.
- Ядра поддерживают кэш в горячем состоянии на оборудовании, которое никто другой не поддерживает.
- Сеть распределяет этот кэш между узлами, которые никогда не были объединены для обмена данными.
Ничего из этого не является универсальным. Все это направлено на агента по написанию кода — рабочую нагрузку с самым высоким объемом в ИИ. Почему все агенты станут агентами по написанию кода:
Если вы выпускаете агента, который пишет код, стек находится на расстоянии одного импорта.
Самые быстрые развертывания всего вышеперечисленного являются частными. Спекулятор, обученный на вашем трафике, а не на нашем, кэш, настроенный под вашу рабочую нагрузку, выделенные мощности, цены ниже публичных тарифов за токен. Более 100 миллиардов токенов в день проходят через Morph таким образом. Если вы движетесь к такому масштабу, свяжитесь с нами.
Если эта работа вас интересует, подумайте о присоединении к нам. https://www.workatastartup.com/jobs/92617
