Автотюнер Helion перебирает тысячи конфигураций GPU-ядра и регулярно попадает в пару процентов от результата инженера, который подбирал параметры руками. Кажется, что после этого оптимизировать нечего. Инженеры PyTorch и NVIDIA показали, что это не так: под конфигурацией ядра лежит второй слой решений, который вообще не относится к языку описания ядер. Речь о настройках ассемблера ptxas, которые до сих пор оставались по умолчанию у всех продакшн-ядер, потому что их никто не искал. На ядре, уже снабжённом зафиксированным AOT-конфигом, поиск по этим настройкам дал 4.7% ускорения, а в многокритериальном режиме 6.9% и одновременно +4.5 процентных пункта пропускной способности кэша L1/TEX.

Что такое CompileIQ и зачем он нужен

CompileIQ это инструмент NVIDIA, который подбирает настройки компилятора ptxas так же, как автотюнер подбирает размеры тайлов: запускает эволюционный поиск по пространству управляющих параметров, проверяет корректность каждого кандидата и упаковывает победителя в переносимый файл Advanced Control File (ACF). Helion научился оценивать такие файлы прямо внутри своего автотюнера, наравне с конфигурациями ядра.

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

Что ищет автотюнер и где он замолкает

Ядро в Helion это обычный PyTorch плюс одна идея: тайлы. Разработчик пишет hl.tile(n) и не указывает размер блока, за него это решает автотюнер. Одно определение ядра порождает пространство из тысяч Triton-конфигураций: block_sizes, loop_orders, l2_grouping, num_warps, num_stages, способ индексации (арифметика указателей, блочные указатели или TMA-дескрипторы), pid_type, reduction_loops. Helion опускает код через TorchInductor в Triton, компилирует каждого кандидата, замеряет время и оставляет победителя.

Дальше победивший Triton-кернел уходит в ptxas, и всё, что происходит ниже этой точки, идёт на настройках по умолчанию. Пространство поиска автотюнера ограничено тем, что умеет выразить язык ядер, а ни один язык ядер не умеет выразить фразу «распредели регистры иначе». Это отдельный интерфейс, который ptxas выставляет наружу начиная с CUDA 13.3.

Что именно крутит ptxas

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

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

Как выглядит ACF и как его подключают

Advanced Control File это бинарный артефакт, который применяется на этапе компиляции. Он переносимый, версионируемый, и его можно поставлять вместе с ядром. Helion открывает этот слой двумя настройками.

Первая просто фиксирует известный хороший ACF в конфиге, который разработчик уже выбрал: helion.Config(advanced_controls_file="h100_example1.acf", block_sizes=[4], num_warps=8, num_stages=2, ...). Здесь ничего не ищется, файл применяется как есть.

Вторая интереснее. Параметр autotune_search_acf принимает список ACF и передаёт их автотюнеру Helion как ещё одно измерение поиска, рядом с размерами тайлов и числом warp-ов. Важная деталь: пустая строка, то есть вариант «без ACF, настройки компилятора по умолчанию», добавляется в список автоматически. Базовая конфигурация всегда остаётся кандидатом, и автотюнер выберет ACF только если тот реально выиграл по замеренному времени. Худший исход включения этой опции: вы потратите часть бюджета автотюнинга и получите то же ядро, что было.

Почему поиск обязан быть совместным

Самый содержательный аргумент в пользу одного цикла вместо двух касается порядка действий. Представьте, что вы сначала тюните ядро, получаете block_sizes=[128, 128] с num_stages=4, а потом отправляетесь искать настройки компилятора. Вы ищете их при условии конфигурации, которая сама была выбрана на дефолтных настройках ptxas.

Но эти измерения связаны через железо. Размер блока и число стадий определяют давление на регистры и объём разделяемой памяти, а политика распределения регистров как раз и решает, превратится ли это давление в spills. Набор настроек, который блеснул бы при block_sizes=[64, 64], просто не будет оценён в этой точке, потому что конфигурация ядра зафиксирована до начала поиска компиляторных параметров. Обратный порядок страдает тем же изъяном. Совместный поиск единственный, кто его лишён, поэтому autotune_search_acf живёт внутри автотюнера Helion, а не выполняется как постобработка.

Плата за это время поиска. Каждый ACF умножает число компиляций и замеров, так что держать опцию включённой для всех ядер кодовой базы бессмысленно. Её запускают один раз, офлайн, на ядре, которое действительно важно, а победителя фиксируют через advanced_controls_file, точно так же, как сегодня фиксируют оттюненные конфиги.

Откуда берутся сами файлы

CompileIQ трактует выбор управляющих параметров компилятора как задачу оптимизации гиперпараметров. Он берёт подготовленное пространство поиска для конкретной версии компилятора и тулчейна, затем запускает эволюционный цикл: предлагает популяцию наборов параметров, компилирует каждый, сверяет результат с эталоном на корректность, замеряет выживших, скрещивает лучших и повторяет. На выходе получается ACF. В самом алгоритме поиска ничего GPU-специфичного нет, вся специфика остаётся в целевой функции, которую пишет пользователь, и обычно это несколько строк вокруг triton.testing.do_bench.

Для большинства задач дополнительный бюджет поиска не нужен, потому что поиск уже выполнили. CompileIQ публикует Booster Packs, готовые каталоги ACF, в своих GitHub-релизах. Текущий helion-booster-pack содержит 24 ACF, настроенных на диапазон ядер Helion и проверенных против ptxas из CUDA 13.4 с прицелом на B200. Реалистичный сценарий для пользователя выглядит так: скачать пакет, направить на него autotune_search_acf, один раз прогнать автотюнинг, зафиксировать победителя. Ни проектировать пространство поиска, ни писать целевую функцию, ни разбираться в компиляторе не требуется.

Эксперимент: последняя миля на NVFP4 GEMV

В качестве кейса авторы взяли nvfp4_gemv_fp4in_kernel, ядро декодного GEMV в формате W4A4 (веса NVFP4 на активации NVFP4). У него уже есть проверенный AOT-конфиг для целевой формы N=8192, K=28672, то есть для проекции down_proj: block_sizes=[4, 256], num_warps=2, num_stages=3. Это принципиальный момент. Сравнение идёт не с наивным ядром, а с ядром, чьё пространство конфигураций уже прочесали авторы на нужном железе. Если под этой планкой есть запас, значит запас настоящий.

Замеры шли на системе GB200 с sm_100 (Blackwell, compute capability 10.0), ptxas 13.3.73 из CUDA 13.3, CompileIQ 1.0.3, Triton 3.8.0 и устанавливаемым из исходников Helion. Все замеры выполнялись через triton.testing.do_bench с залоченной частотой GPU. Каждое улучшение перемеряли независимо, проверяли численно против эталонной реализации PyTorch и фильтровали статистикой: тест Уэлча при p < 0.05 и размер эффекта Коэна d > 0.2.

Первый вариант целевой функции самый очевидный: время выполнения. Кандидатный ACF применяется к фиксированному ядру, ядро запускается один раз, результат сверяется с эталоном, и только прошедшие проверку конфигурации отправляются в бенчмарк. Восемь поколений при размере популяции 15 дали набор, который снизил время с 0.0971 мс до 0.0928 мс, то есть 1.047x или 4.7%. Независимая валидация на больших числах итераций подтвердила результат: базовый прогон дал 0.09752 мс при n=1744, оптимизированный 0.09301 мс при n=2018, ускорение 1.0485 при p = 0.0 и d = 2.501. Повторный прогон воспроизвёл 1.0478 и d = 2.489. Размер эффекта почти в 12.5 раза превышает порог 0.2, так что это не шум.

Когда цель не время, а занятость

Второй эксперимент отвечает на вопрос, что делать, если известен конкретный архитектурный барьер. Профиль ядра на GB200 показал: пропускная способность памяти 78.99%, кэш L1/TEX 89.35%, вычислительные блоки SM всего 37.29%. Теоретическая занятость 25.0%, достигнутая 21.63%, а собственная модель стоимости Nsight Compute оценила возможность ускорения в 21.01% и прямо указала, что теоретическая занятость ограничена числом требуемых регистров. Иначе говоря, ядро упирается в регистры, а этим как раз управляют настройки распределения регистров в ptxas.

Чтобы развернуть поиск на эту цель, меняются две вещи: что именно оценивает целевая функция и направление оптимизации. Вместо do_bench читается метрика ncu sm__warps_active.avg.pct_of_peak_sustained_active, а problem_type переключается с минимизации на максимизацию. Четыре поколения при популяции 6 сдвинули достигнутую занятость на 0.16 процентного пункта, с 21.60% до 21.76%. Звучит как погрешность, но статистика говорит обратное: p = 9.2e-23, d = 13.17. Независимое перемеривание дало скромное, но устойчивое ускорение времени выполнения в районе 1.75–2.0% при p = 3.6e-116 и d = 0.78.

Мораль здесь практическая. Прокси-метрика из профайлера может вести поиск туда, куда прямой замер времени не доберётся, потому что на уровне шума в 1–2% время плохо отличает кандидатов, а счётчик занятости различает их уверенно.

Многокритериальный поиск и один файл на две формы

CompileIQ умеет оптимизировать несколько целей сразу: целевая функция возвращает кортеж вместо одного числа, и поиск строит границу Парето вместо единственного победителя. В этом кейсе поколений было 8, популяция 15, а цели две: время выполнения и пропускная способность L1TEX. Поиск сошёлся к одной доминирующей точке, которая обошла всех остальных кандидатов сразу по обеим осям. Валидация на свежем кэше дала ускорение 1.0689x, то есть 6.9%, при p = 3.516e-295 и d = 87.024, и одновременно рост пропускной способности L1TEX с 81.51% до 85.96%, то есть +4.5 процентного пункта при p = 7.123e-19 и d = 3.686. Ни одна из метрик не была принесена в жертву другой.

Отдельный сценарий касается нескольких форм входа. Реальная модель редко работает с одной формой, и у nvfp4_gemv уже есть AOT-конфиги как для N=8192, K=28672, так и для N=15360, K=5120. Вместо отдельного файла под каждую форму целевая функция возвращает по времени на форму, и поиск ищет единый ACF, который хорош на обеих. Восемь поколений при популяции 15 дали границу Парето из четырёх кандидатов, и pareto_0 показал 0.0927 мс на down_proj и 0.0309 мс на второй форме против дефолтных 0.0972 и 0.0315 мс. Валидация подтвердила 1.05x на первой форме (p = 1.1e-224, d = 67.71) и 1.01x на второй (p = 6.2e-73, d = 4.54).

Что это значит для продакшна

Авторы честно называют ожидаемый диапазон. На ядрах, которые уже хорошо оптимизированы, настройка компиляторных параметров даёт 2–7% времени выполнения. Для критичных ядер топовых моделей это может быть много. На менее оптимизированных ядрах выигрыш доходит до двузначных процентов, и там потенциал для сквозного ускорения выше.

Сквозной рецепт выглядит так. Сначала профайлируем модель и находим самые горячие ядра. Каждое из них становится целью оптимизации CompileIQ и получает собственный ACF. Применяются эти файлы либо через зафиксированные конфиги Helion, либо через API ptxas_options в Triton, либо прямо через командную строку ptxas, в зависимости от того, как реализовано ядро. Ключевое свойство схемы: результат находится автоматически, почти без экспертизы в компиляторах, проверяется собственной проверкой корректности и сохраняется в файл, который можно закоммитить рядом с ядром и переиспользовать на каждой сборке.

Экономика здесь сходится именно там, где нужно. В цикле инференса есть ядра, которые выполняются непрерывно на больших парках ускорителей. Пара процентов на дне стека, применённая к операциям, которые доминируют в шаге декодирования, на таком масштабе превращается в реальные деньги.

Границы применимости

ACF привязан к железу и версии тулчейна. Файл, настроенный под sm_90, невалиден на sm_100, и применение несовпадающего ACF может дать не просто медленный код, а неверные результаты или ошибку времени выполнения. Манифест booster pack прямо указывает целевую архитектуру и версию тулчейна, и именно поэтому проверка корректности внутри цикла поиска не опциональна.

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

Лицензионная картина смешанная и её стоит знать заранее. Helion открыт и живёт под эгидой PyTorch Foundation. CompileIQ бесплатен, но распространяется по проприетарному NVIDIA Software License Agreement. Сами файлы ACF при этом свободно распространяемы, поэтому booster packs можно публиковать, а ACF класть в открытый репозиторий рядом с ядром.

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

Нужно ли быть экспертом по ptxas, чтобы это использовать?

Нет. Готовый helion-booster-pack из 24 ACF покрывает типовые ядра Helion, а подключение сводится к одному параметру autotune_search_acf и одному прогону автотюнинга. Собственный поиск через pip install compileiq нужен только для нагрузки, которой нет в пакете.

Может ли включение поиска ACF ухудшить результат?

Нет. Helion автоматически добавляет в список кандидатов пустой ACF, то есть настройки компилятора по умолчанию. Автотюнер выберет файл только при выигрыше по замеренному времени, так что худший исход это потраченный бюджет поиска и то же самое ядро на выходе.

Почему нельзя сначала оттюнить ядро, а потом отдельно поискать настройки компилятора?

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

Итог

Главный вывод статьи не в самом инструменте. Он в том, что управляющая поверхность компилятора это полноценное пространство оптимизации, оно поддаётся поиску, и почти никто его пока не ищет. Это верно для любого компилятора с настраиваемыми параметрами и для любого бэкенда, не только для ptxas и NVIDIA.

CompileIQ и Helion просто сделали первый шаг доступным: 4.7% на уже оттюненном NVFP4 GEMV, 6.9% в паре с ростом пропускной способности L1TEX, один ACF на две формы входа. Если у вас есть ядро, которое вы считали выжатым до предела, начните с pretuned_kernels в репозитории Helion, добавьте autotune_search_acf и сравните время на своём железе. Разница в пару процентов на дне стека обычно и есть та разница, которую видно в счёте за инференс.