Одно ядро вместо трёх, автоматический подбор конфигурации под каждую форму входа и более 10% прироста пропускной способности на отдельных нагрузках. Команды OCTO и vLLM из Red Hat вместе с командой Helion из Meta встроили Helion, kernel DSL из экосистемы PyTorch, в linear backend vLLM. Результат, который обычно достигается ручной специализацией ядер под каждую модель и форму, здесь получен через автоматический поиск конфигураций. Разбираем, как устроена интеграция, что она даёт на NVIDIA Hopper и какие компромиссы у этой схемы.

Что такое Helion

Helion это PyTorch-нативный kernel DSL для написания высокопроизводительных ядер в tile-модели программирования. Разработчик описывает ядро один раз на лаконичном Python-подобном коде, а система генерирует специализированный вариант под конкретную нагрузку и железо. Вместо ручной подгонки работает ahead-of-time (AOT) автотюнер: он системно исследует пространство вариантов, от раскладки памяти и планирования ядра до алгоритмических решений, и выбирает лучший конфиг. Есть и режим LLM-guided search, где языковая модель предлагает перспективные точки старта для поиска.

В vLLM Helion уже упоминался как один из портируемых DSL, на которых строятся hardware-agnostic слои движка. Теперь он добрался до одного из самых горячих участков: линейных слоёв, которые в dense-моделях съедают заметную долю времени инференса.

Зачем трогать линейные слои

Для квантованных линейных слоёв, а это форматы FP8, INT8, INT4 и NVFP4, vLLM предоставляет специализированные linear backends, которые подключают аппаратно оптимизированные ядра из библиотек CUTLASS, DeepGEMM и FlashInfer. Схема работает, но упирается в экономику кернел-разработки. Разные формы входа выигрывают от разных алгоритмических вариантов GEMM, и каждый такой вариант приходится реализовывать, бенчмаркать и подбирать к нему эвристику диспатча.

Два главных варианта выглядят так. Split-K делит измерение K между несколькими thread-блоками и добавляет параллелизма, когда размерности M и N слишком малы для полной загрузки GPU. Swap-AB переписывает умножение A@B как (B.T@A.T).T и помогает на формах с маленьким M за счёт более выгодного тайлинга. Дальше начинается ручная работа: в текущем Block_FP8 бэкенде vLLM, например, прописано правило переключаться на Swap-AB при M меньше 32 на Hopper. Такие эвристики хороши в среднем, но почти никогда не оптимальны для каждой конкретной модели и нагрузки.

Одно ядро вместо трёх вариантов

Helion предлагает другой подход: вместо отдельных реализаций все алгоритмические решения выносятся в tunable-параметры одного унифицированного ядра. В примере из статьи split_k регистрируется как tunable со значениями степени двойки до 256, а swap_ab как булев параметр. Автотюнер перебирает комбинации в заданном пространстве, бенчмаркает их и сам выбирает лучший вариант вместе с низкоуровневым конфигом под каждую форму.

То есть вместо вопроса «какое ядро вызвать» остаётся вопрос «какой конфиг выбрал автотюнер», и вторая часть полностью автоматизирована. Для квантованных ядер структура та же, добавляется лишь логика квантизации и скейлинга. Форматы, на которых сосредоточились авторы: FP8_Dynamic с per-token активацией и per-channel весами, W8A8_INT8 с той же схемой скейлинга в INT8 и Block_FP8 с блоками 1×128 для активаций и 128×128 для весов.

Гибридный диспатч: Helion только там, где он выигрывает

Интеграция не пытается заменить дефолтные ядра везде. Работает гибридная схема, привязанная к runtime-значению num_tokens и покрытию CUDA Graph. Для маленьких форм, до порога max_helion_size (в этой работе он равен 32), linear backend направляет вычисления в Helion-ядро, выполняемое через CUDA Graph replay. Для форм побольше, а это в основном prefill, происходит откат на дефолтные ядра бэкенда: CUTLASS, DeepGEMM или FlashInfer в зависимости от формата.

У такой схемы три практических плюса. Helion-ядра исполняются только под CUDA Graph, поэтому исчезает CPU-overhead на диспатч и запуск ядер. Тонкий тюнинг ограничивается диапазоном малых num_tokens, который доминирует в декодинге, а значит тюнить нужно меньше форм. И конфигов становится меньше, поэтому предварительно настроенный набор реально поддерживать и валидировать. Конкретный список форм: num_tokens = [1, 2, 4, 8, 16, 24, 32], каждая форма тюнится индивидуально под GEMM-размеры конкретной модели.

Как автотюнер ищет конфиги

Тюнинг запускается утилитой из vLLM: scripts/autotune_helion_kernels.py с ключами --kernels scaled_mm block_scaled_mm и --autotune-effort full. Два момента важны для качества результата.

Первый: бенчмаркинг идёт с включённым CUDA Graph (переменная HELION_BENCHMARK_CUDAGRAPH=1). Поскольку Helion-ядра в проде работают именно под графом, автотюнер должен оценивать конфиги в тех же условиях, а не в идеализированном одиночном запуске.

Второй: поиск стартует не с нуля. Автотюнер LLMSeededLFBOTreeSearch сначала просит языковую модель, в этой работе использовался Claude Opus 4.8, предложить набор перспективных конфигов-сидов, а затем LFBOTreeSearch оптимизирует их дерево. Качественные стартовые точки помогают найти лучшие конфиги и сокращают общее время поиска.

Результаты: уровень ядра

Замеры шли на NVIDIA H100 80GB HBM3 на пяти dense-моделях Qwen3: 1.7B, 4B, 8B, 14B и 32B, плюс свежая Qwen3.8-27B, в трёх восьмибитных форматах квантования. Каждое Helion-ядро сравнивалось с ядром, которое использует дефолтный linear backend vLLM на Hopper: CUTLASS для FP8_Dynamic и W8A8_INT8, FlashInfer и DeepGEMM для Block_FP8.

Геометрическое среднее ускорение по всем формам входа: 1,110× для FP8_Dynamic против CUTLASS, 1,178× для W8A8_INT8 против CUTLASS, 1,149× для Block_FP8 против FlashInfer и 1,177× против DeepGEMM. Разброс по отдельным формам заметный, и это как раз аргумент в пользу тонкого тюнинга: единый конфиг или общая эвристика не могут быть оптимальны сразу везде.

Результаты: end-to-end

Отдельный вопрос, конвертируются ли улучшения ядер в скорость сервинга. Сервер поднимался командой vllm serve с флагом --linear-backend helion, ограничением --max-num-seqs 32, тензорным параллелизмом 1 и намеренно выключенным prefix caching. Нагрузка: ShareGPT-датасет через vllm bench serve, сравнение с дефолтным бэкендом на тех же моделях и форматах.

Итог: прирост стабильный по всем проверенным моделям и форматам, а на отдельных нагрузках пропускная способность растёт более чем на 10%. Для сервинга, где линейные слои это заметная доля latency, такой прирост получается без единой строки нового CUDA-кода от команды.

Треугольник компромиссов

Авторы не прячут цену. Тонкий тюнинг Helion балансирует между производительностью, удобством и поддерживаемостью, и улучшение двух пунктов почти всегда ухудшает третий. Четыре конкретных издержки: AOT-тюнинг может занимать часы, старт vLLM дорожает из-за JIT-компиляции при захвате CUDA Graph (на прогретых запусках спасает кеш скомпилированных артефактов), вне зоны CUDA Graph диспатч ядер добавляет CPU-overhead, а большие наборы предварительно настроенных конфигов сложно поддерживать и невозможно исчерпывающе проверять юнит-тестами и CI.

Именно поэтому в работе и появился порог в 32 токена: он срезает все четыре издержки разом, оставляя выигрыш там, где он окупается. Это инженерный компромисс, описанный авторами прямо, а не обнаруженный читателями в комментариях.

Как это будут внедрять

Сейчас Helion linear backend живёт в форке vLLM у Red Hat и готов к продакшену: форк включает и тулинг автотюнинга, чтобы пользователи могли генерировать конфиги под свои модели. Главный вопрос для upstream: кто поддерживает гору предварительно настроенных конфигов. Обсуждаемая модель такая: ядра и фреймворк интеграции живут в основном репозитории с дефолтным конфигом для функциональных тестов и CI, а тюнинг под конкретные нагрузки делегируется пользователям. Для latency-критичных ядер это разумно: команда один раз прогоняет автоматический тюнинг перед деплоем и получает оптимизацию под своё железо и модель.

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

Почему это важно за пределами одной интеграции

Кернел-инженерия переживает сдвиг в сторону DSL и автоматизации. Triton сделал написание ядер доступным без CUDA-экспертизы, ThunderKittens вернул контроль над железом тем, кому нужен максимум, а Helion добавляет к этому автоматический поиск конфигураций. На этом фоне интеграция в vLLM это тест на зрелость: kernel DSL должен не просто генерировать быстрые ядра в лаборатории, а переживать продакшен-требования сервинга, включая CUDA Graph, динамические формы и квантованные форматы.

Второй контекст это кодинг-агенты: они уже умеют писать оптимизации под конкретное железо, и вопрос ближайших лет в том, останется ли оптимизация ядер уделом людей или превратится в автоматический шаг пайплайна. Автотюнинг с LLM-подсказками в этой работе как раз такой шаг, только осторожный: модель предлагает стартовые точки, а численная оптимизация остаётся за математикой.

Что дальше

Планы у команды широкие. По железу: созревание CuteDSL-бэкенда для NVIDIA Blackwell (первые результаты на нём уже конкурентны), работа над AMD GPU и TPU. По моделям: текущая работа сфокусирована на dense-моделях, а для MoE-архитектур важнее MoE-бэкенд, и он в разработке. Расширение на эти направления превратит точечную интеграцию в полноценную альтернативу ручной кернел-инженерии для всего инференс-стека.

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

Что такое Helion простыми словами?

Это язык описания ядер для GPU из экосистемы PyTorch. Вы пишете ядро один раз на Python-подобном коде, а Helion сам подбирает конфигурацию под конкретную форму данных и железо через автоматический поиск. Вместо ручной оптимизации под каждый случай работает автотюнер, который перебирает варианты и выбирает лучший.

Почему Helion работает только до 32 токенов?

Это гибридный диспатч. Малые num_tokens доминируют в декодинге и выигрывают от тонкого тюнинга, а на больших формах prefill дешевле и надёжнее дефолтные CUTLASS и DeepGEMM. Порог в 32 токена ограничивает и время тюнинга, и набор конфигов, и CPU-overhead, оставляя выигрыш там, где он окупается.

Заменяет ли Helion CUTLASS и DeepGEMM?

Нет, он работает в паре с ними. Для форм до порога max_helion_size вычисления идут через Helion-ядро под CUDA Graph, для крупных форм backend откатывается на дефолтные библиотеки. Это дополнение к существующим ядрам, а не замена: каждый участок получает лучший из доступных вариантов.

Нужно ли тюнить конфиги самостоятельно?

Зависит от сценария. В форке Red Hat уже лежат предварительно настроенные конфиги для популярных моделей, и для старта этого достаточно. Если у вас своя модель или особое железо, автотюнинг можно прогнать самому: скрипт и инструкции входят в тот же форк.

Итог

История с Helion в vLLM это про смену экономики кернел-разработки. Раньше за единицы процентов пропускной способности платили ручной специализацией ядер под каждую модель, теперь эту работу забирает автотюнер: одно ядро, tunable-параметры вместо форков, LLM-подсказки вместо перебора вслепую. На Hopper это даёт от 1,11× до 1,18× на уровне ядер и больше 10% end-to-end на отдельных нагрузках, а цена измерена в часах тюнинга и обслуживании конфигов.

Практический шаг простой: если вы сервите dense-модели на Hopper, попробуйте форк с флагом --linear-backend helion и прогоните автотюнинг под свою модель. Замеры и фидбек авторы собирают в RFC по Helion linear backend в vLLM, и именно от реальных деплоев зависит, станет ли эта схема мейнстримом.