Когда мы создавали Fast Apply, мы поняли, что в открытом стеке вывода не хватает ключевого компонента: специализированного предсказателя (speculator) для конкретной задачи. Универсальный черновик дает ускорение в 1,93 раза, а специализированный — в 100 раз. Опираясь на это, мы обучили WarpGrep и некоторые из лучших специализированных моделей. Этот механизм лежит в основе спекулятивного декодирования.
Малые модели не подчиняются тем же законам масштабирования, что и большие. Мы обучаем пользовательские предсказатели для генерации кода, чтобы они были лучшими в своем классе при выводе программного кода. Мы обслуживаем открытые модели — Qwen, GLM, DeepSeek, MiniMax — для одной задачи: агента для программирования. Чтобы сделать их быстрыми, нужно выполнить три вещи, которые открытый стек не сделает за вас.
- Специализированный предсказатель. Черновик, обученный на собственном коде модели, а не на данных из интернета. Универсальный черновик: 1,93x. Обученный на целевой модели: 3,07x.
- Ядра (kernels), оптимизированные для задержки или пропускной способности. Декодирование FP4, которое остается 4-битным вплоть до тензорного ядра, автоматически подбирается на недорогих GPU, которые никто другой не настраивает. От 97 до 162 токенов в секунду на карте за 7000 долларов.
- Межсоединение, оптимизированное для высокой частоты попаданий в кэш. Однопроходный all-reduce через PCIe и префиксный кэш, который работает через обычный TCP даже там, где нет NVLink.
Каждый из этих пунктов — это место, где стандартный стек останавливается, а мы идем дальше.
1. Вывод, который становится быстрее по мере использования
Спекулятивное декодирование: небольшая модель-черновик угадывает следующие несколько токенов, целевая модель проверяет их за один проход, и вы продолжаете выполнение до первой ошибки. Все решает один показатель: коэффициент принятия, то есть как часто целевая модель соглашается с догадкой.
Универсальный черновик — плохой предсказатель. Для Vicuna-13B готовый черновик на 68 млн параметров дает 1,93x; черновик, обученный на выводе целевой модели, дает 3,07x при той же целевой модели и настройках. Этот разрыв и есть суть.
Архитектуры общедоступны и эффективны. EAGLE-3 позволяет черновику обучаться на «сырых» данных, а не копировать признаки целевой модели, и длина принятия возрастает с 3,96 до 6,62. DFlash, спецификация V2 для SGLang с июня 2026 года, составляет черновик целого блока за один проход: более чем 6-кратное ускорение без потерь, 3,2x на HumanEval, где EAGLE-3 дает 2,2x.
Но архитектура — это лишь пустой разъем. Никто не даст вам готовый черновик, обученный под вашу целевую модель и вашу нагрузку. Вы либо обучаете его сами, либо используете универсальный и довольствуетесь 1,93x.
Обучение хорошего черновика — это обучение малой модели. Fast Apply и Compact сделали нас одной из лучших команд в мире в этой области. Вещь, которую вы узнаете при работе с моделями менее 30 млрд параметров: границы законов масштабирования перестают действовать. Chinchilla говорит, что около 20 токенов на параметр — это вычислительно оптимально, но это предполагает, что обучение — основная статья расходов. Для модели, которую вы обучаете один раз, а используете миллиарды раз, это не так, и оптимум сильно смещается в сторону малых и переобученных моделей.
Предсказатель живет именно там. Маленький, переобученный, настроенный под одно распределение.
Поэтому мы обучаем по одному предсказателю для каждой открытой модели, используя код вместо веб-текстов. Сгенерированный код повторно использует шаблоны и символы, уже имеющиеся на экране, а правка — это в основном копия файла, который она редактирует. Черновик, прочитавший миллион diff-файлов, предсказывает эти токены. Тот, что читал интернет — нет, поэтому код является задачей с самым высоким ускорением для любого метода спекуляции. Для Fast Apply этот специализированный предсказатель позволяет достичь 10 500 токенов/с при применении и 33 000 при сжатии. Те же веса Qwen, которые можно скачать. Наша версия быстрее, потому что предсказатель, работающий с ней, был обучен нами на этой работе.
2. Ядра, оптимизированные для задержки или пропускной способности
Декодирование ограничено памятью. При размере пакета 1 matmul — это GEMV: он считывает каждый вес один раз на токен и почти не выполняет арифметических операций на байт, поэтому токены в секунду — это пропускная способность, деленная на количество затронутых байтов. Единственный рычаг — это байты. Храните веса в FP4 (e2m1), и вы перемещаете 4 бита там, где bf16 перемещает 16, что дает 4-кратное сокращение трафика HBM. Поэтому вся задача заключается в том, чтобы держать их в 4-битном формате до тех пор, пока они не попадут в тензорное ядро.
Стандартные ядра останавливаются на этом этапе. Они деквантуют FP4 в bf16 в HBM, возвращают 4-кратный объем данных и вызывают обычный GEMM. Если сделать все правильно, вы вообще не будете материализовать bf16. Вы передаете упакованные веса через cp.async/TMA, распаковываете e2m1 и масштабирование блока 16 в регистрах, и подаете их прямо из общей памяти в тензорные ядра FP4 архитектуры Blackwell. Для 80B MoE это сгруппированный GEMM только по активным экспертам, плюс резидентное ядро, чтобы каждый SM оставался активным при пакете 1, плюс split-K для заполнения машины, которую GEMV не может насытить в одиночку. Пропустите что-то одно — и вы вернетесь к ограничению bf16.
Настройка этого для разных архитектур — тяжелый труд, а стандартные решения написаны для карт, которые покупают передовые лаборатории. Перенесите ядро без перенастройки, и оно будет работать на 7% от оптимальной производительности; достижение современного уровня на AMD MI250 потребовало переписывания 40% кода ядра flash-attention вручную. Поэтому мы не пишем пространство поиска вручную. Ядро проверяемо, оно либо корректно относительно эталонного вывода, либо нет, и это делает поиск автоматизируемым.
Система предлагает ядро, проверяет его на производственных трассах, тестирует на недорогих картах NVIDIA и AMD, которые никто другой не настраивает, и выпускает победителей. KernelBench показывает, почему это должно быть автоматизировано: при оценке по критериям «корректно и быстрее» передовые модели проходят менее 20% задач «на холодную». Большой объем через жесткий цикл проверки — единственный путь.
Один результат: наши ядра warp-decode запускают 80B MoE на скорости 162 токена/с на карте за 7000 долларов, что выше 97 и обходит 120 токенов/с на H100 за 25 000 долларов. Без потери точности, код открыт. Это окупается только потому, что вычислительные мощности в дефиците, что придало цену тем пробелам, которые проигнорировал общий стек.
3. Межсоединение, оптимизированное для высокой частоты попаданий в кэш
Трафик программирования разделяет 97% своих префиксных токенов, при этом промпты в 37–2494 раза длиннее выводов. Кэшируйте префикс, и следующий запрос оплатит только новые токены. Частота попаданий — это цена. Абстракция кэша открыта, и мы ее используем: RadixAttention хранит префиксы в дереве, кэш-ориентированный маршрутизатор повышает частоту попаданий с 20% до 75%, HiCache выгружает дерево в оперативную память хоста и удаленное хранилище, а на Qwen3-Coder-480B повышает частоту попаданий с 40% до 80% и удваивает пропускную способность.
У дешевых GPU есть подвох. Нет NVLink.
NVLink передает 900 ГБ/с между GPU. PCIe Gen5, шина на доступных серверах, передает 64 ГБ/с в каждом направлении. В 14 раз меньше. Это незаметно, пока вы не шардируете модель между GPU, а затем это становится всем: тензорный параллелизм запускает all-reduce после внимания и MLP каждого слоя, и этот all-reduce стоит 8–11% шага на NVLink и 40–75% на PCIe. Нет быстрой фабрики — и коммуникации съедают большую часть прямого прохода.
Стандартное решение — купить NVLink. Мы вместо этого написали ядро.
Кольцевой алгоритм all-reduce в NCCL предназначен для передачи больших объемов данных по высокоскоростной сети. При размере пакета 1 сообщения составляют несколько сотен КБ и ограничены задержкой, поэтому 2(N-1) переходов в кольце становятся определяющим фактором. Вместо этого мы используем однопроходный all-reduce: каждый GPU отправляет свой сегмент напрямую в память партнера через PCIe P2P и выполняет редукцию локально за один цикл обмена. Мы запускаем его в отдельном потоке, чтобы он перекрывался с операцией GEMM следующего слоя, объединяем редукцию с эпилогом нисходящей проекции вместо запуска второго ядра и квантуем данные до FP8, чтобы вдвое сократить нагрузку на шину. Графы CUDA устраняют накладные расходы на запуск для каждого слоя, которые в противном случае возникали бы при использовании коллективных операций. Это позволяет компенсировать большую часть 14-кратного разрыва. Другое ядро — это префиксный кэш, который передает данные между машинами по обычному протоколу TCP.
HiCache уже определяет удаленный уровень L3 за бэкендом, состоящим из трех функций: get, exist, set. Это работает через любой транспорт. Подвох в том, что заявленные преимущества достигаются при использовании RDMA, где передача занимает менее миллисекунды и составляет менее 0,1% от задержки запроса. Обычный TCP на порядок медленнее. На системе, оснащенной только PCIe, открытый стек незаметно выходит из строя, потому что приведенные цифры предполагали наличие оборудования, которого у вас нет.
Таким образом, преимущество TCP не может быть обеспечено самим транспортом. Оно достигается за счет частоты попаданий. Обученный спекулятор и автооптимизированные ядра обеспечивают достаточно высокий уровень попаданий, чтобы префикс, отсутствующий в GPU и оперативной памяти хоста, извлекался от соседа по TCP вместо повторного вычисления, а пропуск этапа префиллинга оказывается быстрее, чем медленная выборка. По сравнению с полным перевычислением, такая выборка сокращает время до получения первого токена на 84%.
Мы заменили высокоскоростную сеть, которую все покупают, чтобы избежать этой проблемы, на программные ядра. Мы используем GPU, списанные рынком, с частотой попаданий, для которой, как предполагается, требуется оборудование, которое мы не покупали.
Одна рабочая нагрузка
Три вещи, один цикл:
- Спекулятор составляет черновик собственного программного кода модели.
- Ядра поддерживают кэш в актуальном состоянии на оборудовании, которое никто другой не поддерживает.
- Сеть распределяет этот кэш между машинами, которые никогда не были предназначены для обмена данными.
Ничего из этого не является универсальным. Все это направлено на агента-программиста, самую массовую рабочую нагрузку в ИИ. Почему все агенты станут агентами-программистами:
Если вы выпускаете агента, который пишет код, стек находится на расстоянии одного импорта.
Самые быстрые развертывания всего вышеперечисленного являются частными. Спекулятор, обученный на вашем трафике, а не на нашем, кэш, настроенный под вашу нагрузку, выделенные мощности, цены ниже публичных тарифов за токен. Более 100 миллиардов токенов в день проходят через Morph таким образом. Если вы движетесь к таким масштабам, свяжитесь с нами.
Если эта работа вас интересует, подумайте о присоединении к нашей команде. https://www.workatastartup.com/jobs/92617