CUDA DSLs уходят на пенсию. В прошлом году разработчики взялись писать megakernel — и намучились: сложные структуры данных, синхронизация между потоками, SM и GPU, глубоко вложенный контроль потока. В голове это не удержать, поэтому они построили слой абстракции. Даже с ним несколько месяцев разбирались с гонками и дедлоками, пока модели Llamas не заработали быстро. В этом году они собрали MoE megakernel, но уже без абстракции. Агенты позволили работать со сложностью напрямую и писать код сразу под конкретное железо. Промежуточные C++-слои не понадобились.
За квартал проявился паттерн: задачи, которые и раньше решались без абстракций, например оптимизированное GEMM-ядро, теперь почти автоматизируются. Агенту всё ещё надо говорить, какие PTX-инструкции использовать и как устроена warp specialization, но до state-of-the-art он доходит очень быстро. А вот задачи, для которых абстракции были нужны, например megakernel, пока не автоматизированы: one-shot megakernel сегодня не получится. Но агент позволяет засунуть абстракцию прямо в промпт — неаккуратную, неполную, — вместо выверенных C++-шаблонов. Сложность, которая в прошлом году была неуправляемой, стала терпимой: появился «компилятор», который принимает расплывчатые инструкции и выдаёт код. Идеи из megakernel остались важны, но выражать их стало проще.
Работа абстракции как разгрузки для мозга начинает уходить. Агенты занимают эту должность. По индукции дальше на пенсию отправятся CUDA DSLs, включая любимый ThunderKittens. Возможно, в следующем году или раньше. Насколько это зайдёт — непонятно. Где-то внизу есть пол, где абстракция остаётся чистым контрактом, а не костылём для мышления, но отсюда его не видно.
Отсюда провокация: если промпт может сгенерировать любой слой стека, сколько стоит кодовая база? Её преимущество — точность. Это артефакт, который машина исполняет одинаково. Но точность дорого обходится: кодовая база привязана к языку, фреймворку, железу и договорённостям, которые понимает только своя команда. Она выглядит переносимой, но едва ли. Промпт наоборот: размытый, зато путешествует. Если исполнитель достаточно умён, чтобы правильно заполнить пробелы, результат будет верным без заранее стандартизированных деталей. DSL и фреймворки — это и есть спецификация в виде кодовой базы. Если исполнитель перестаёт быть тупым, DSL проигрывает.
Кодовая база — несовершенная среда. Огромная, барочная, с багами. Не потому что инженеры небрежны: «точно для машины, читабельно для человека, гибко для будущего» — почти невыполнимые требования к строке токенов. Мы уже принимали непрозрачные трансформации: мало кто пишет ассемблер. Компилятор переупорядочивает, инлайнит, векторизует. Интеллект уже стоит между намерением и машиной. Вопрос не в том, изменится ли это качественно, а в том, где пройдёт граница. Авторы считают, что качественного скачка тут нет.
Но просто так отпускать абстракции нельзя. Во-первых, абстракция — это ещё и общая поверхность, к которой крепятся приложение, повторное использование и ревью. Без неё — кембрийский взрыв проблем с верификацией. Десять команд на ThunderKittens проверяют одну и ту же семантику тайлов — это накапливается. Десять команд с уникальными megakernel имеют десять несвязанных множеств проблем, они не стратифицируются.
Во-вторых, нет общего оракула. Для megakernel тесты — это и контракт, и работа. Кто их пишет и кто проверяет? Когда фреймворк удалили, оракул остался: reference-реализации, допуски по численным результатам, интуиция о том, как должен выглядеть профилированный Gantt chart. Слой можно удалить, только если есть оракул, который его переживёт. В областях, где никто не знает оракула, ничего из этого не работает.
В-третьих, это один предвзятый пример. Авторы удалили абстракции в области, которую глубоко знают. Абстракции также передают знание новичкам. Неизвестно, заполнит ли агент пробелы для того, кто впервые слышит про warp specialization. Может, агенты сделают онбординг ненужным. А может, людям всё ещё нужны эти слои, и это просто удобная позиция для уже обученных.
Несмотря на всё, трансформация выглядит неизбежной. Возможно, придётся отпустить любимые фреймворки. Что остаётся? Намерение, инварианты, тесты и выстраданное знание предметной области, которое сейчас живёт внутри абстракций ThunderKittens и удерживает их на правильном железе. Библиотека может уйти, но знание не испарится от того, что CUDA или HIP перегенерируются, а не тюнятся вручную. Доверие поднимается на уровень выше: мы изучаем спецификацию и оракул, а не diff. Реализация становится одноразовой — кэшем конкретной компиляции, а не источником истины. Кодовые базы не исчезнут. Они просто перестают быть тем, что мы называем «проектом». Абстракции уходят, идеи остаются.