Знакомьтесь с qwenGPU: Qwen3.5-2B с 45 скомпилированными GPU-слоями, выполняющими CUDA
qwenGPU — экспериментальная модель, которая компилирует SIMT-GPU непосредственно в веса Qwen3.5-2B. Разработчик Vixhal (@TheVixhal) выпустил её на GitHub под лицензией MIT. Модель выполняет настоящий CUDA-код, скомпилированный в NVIDIA PTX, внутри собственных прямых проходов. Для разработчиков это проверка одной идеи: точные вычисления внутри предварительно обученной LLM, а не рядом с ней.
Кратко
- Размер: Qwen3.5-2B (24 слоя) вместе с 45 скомпилированными GPU-слоями — всего 69 слоёв. Контекст: 262 144 токена (базовая модель).
- Запускается на: Apple Silicon, GPU NVIDIA или CPU, примерно с 16 ГБ свободной памяти. Квантизированной сборки нет.
- Производительность: Каждый запуск в тестах разработчика был точным; отклонение текста составляет около 7e-4 по KL.
- Лучший результат: эталонный интерпретатор совпал с реальной NVIDIA Tesla T4 для 49 из 49 ядер.
- Худший результат: около 10–15 токенов в секунду на Apple MPS; CUDA-код, написанный Qwen, прошёл 4 из 5 тестов.
Что такое qwenGPU?
qwenGPU — это Qwen3.5-2B с 45 скомпилированными трансформерными слоями, размещёнными перед собственными 24 слоями Qwen. Ни один из этих слоёв не обучался. Компилятор вычисляет каждый вес.
Как GPU помещается в веса модели?
GPU представляет собой 4-канальную SIMT-машину. У неё 15 регистров на канал, 4-этапный конвейер, ПЗУ на 32 блока и 126 ячеек памяти. Она использует только собственные типы слоёв Qwen3.5: линейное внимание Gated DeltaNet и MLP SwiGLU.
Секрет заключается в рекуррентном состоянии Gated DeltaNet. В головах схемы вентиль удерживается открытым, а сила записи равна 1. Каждый ключ является one-hot-вектором. Каждое обновление перезаписывает ровно 1 строку, а one-hot-запрос считывает её обратно. Поэтому каждая строка работает как ячейка памяти.
В этом состоянии хранятся ПЗУ программы, регистры и глобальная память. В README указано 16 строк регистров на канал, одна из которых всегда считывается как ноль. Это соответствует 15 используемым регистрам из публикации в X.
GPU поддерживает целочисленную арифметику, предикаты, расходящиеся ветвления и циклы. Он также поддерживает общую память, atomicAdd, __syncthreads() и перестановки внутри варпа.
Что происходит во время запуска CUDA?
Хост выполняет роль NVRTC и драйвера CUDA. Он компилирует CUDA в PTX с помощью clang, а затем преобразует PTX в машинные блоки. Он передаёт программе и входным массивам модели в виде токенов.
После этого модель самостоятельно записывает запуск. Каждый токен EXEC(b, w) выполняет блок b для варпа w за 1 прямой проход. Следующие за ним токены VAL содержат результаты регистров. Записи в память возвращаются в виде записей MEM. Хост считывает результаты из рекуррентного состояния и не выполняет никаких вычислений.
В демо-видео скалярное произведение [1, 2, 3] и [4, 5, 6] возвращает 32. Модель записала этот запуск за 47 токенов.
Познакомьтесь с qwenGPU: Qwen3.5-2B с SIMT-GPU внутри своих весов, с 45 скомпилированными слоями перед собственными слоями Qwen, использующими слои Gated DeltaNet в качестве регистров и памяти.
— vixhaℓ (@TheVixhal) 1 ноября 2026 г.
Она выполняет настоящий CUDA-код в своих прямых проходах. Хост компилирует ядро в PTX и загружает его, после чего модель записывает каждый… pic.twitter.com/iM5IETEbIv
Нарушает ли GPU способность Qwen работать с текстом?
Согласно оценке разработчика, нет. Схема отключена для текстовых токенов, а собственные веса Qwen не изменены. Распределение следующего токена объединённой модели остаётся в пределах примерно 7e-4 по KL от исходного.
Насколько точен qwenGPU?
Все числа предоставлены разработчиком и измерены на Apple Silicon с MPS в bf16:
- Каждый запуск был точным при повторном воспроизведении с
--check. - Эталонный интерпретатор совпал с NVIDIA Tesla T4 для 49 из 49 тестовых ядер.
- В обычном режиме на 20 из 24 выбранных вопросов в чате были получены числа, вычисленные GPU.
- В расширенном режиме правильное ядро с первой попытки было написано для 4 из 5 вопросов.
В обычном режиме Qwen записывает выражение вроде [1, 2, 3] @ [4, 5, 6] для готовых ядер. В расширенном режиме Qwen сам пишет ядро CUDA. В README предупреждается, что CUDA-код модели с 2B параметров всё ещё ненадёжен.
Как qwenGPU сравнивается с Qwenputer и Percepta?
qwenGPU использует компилятор трансформера из проекта Qwenputer Divyansh Agrawal. Qwenputer компилирует 32-битный процессор RISC-V в Qwen3.5-4B. Percepta в марте 2026 года встроила интерпретатор WebAssembly в трансформер с двумерными головами внимания.
| Характеристика | qwenGPU | Qwenputer (RISC-V) | Percepta |
|---|---|---|---|
| Базовая модель | Qwen3.5-2B, 24 слоя | Qwen3.5-4B | Пользовательский трансформер с двумерными головами внимания |
| Встроенная машина | 4-канальный SIMT-GPU | 32-битный RISC-V (RV32IM + Zba + bignum) | Интерпретатор WebAssembly |
| Запускаемые программы | CUDA, преобразованная из PTX | C и MicroPython | C, скомпилированный в WebAssembly |
| Скомпилированные слои | 45 | 11 | Не раскрыто |
| Где хранится состояние | Состояние Gated DeltaNet | Состояние Gated DeltaNet | Трассировка выполнения в контексте |
| Роль хоста | Компилирует CUDA в PTX, загружает токены, считывает результаты | Загружает прошивку; вычислений не выполняет | Не раскрыто |
| Проверка корректности | Интерпретатор совпал с Tesla T4 для 49 из 49 ядер | 0 несовпадений примерно в 230 000 воспроизведённых позициях | Заявлены миллионы надёжных шагов |
| Заявленная скорость | 10–15 токенов/с (Apple MPS) | Около 3 200 токенов трассировки/с (L4, Triton) | Более 30 000 токенов/с (CPU) |
| Отклонение текста относительно базовой модели | KL около 7e-4 | KL 6e-4 | Неприменимо |
| Лицензия | Код MIT; контрольная точка Apache-2.0 | Код MIT; веса Apache-2.0 | Не раскрыто |
Показатели скорости получены на разном оборудовании и для разных типов токенов, поэтому напрямую сравнивать их нельзя.
Что нужно для запуска?
Для быстрого старта нужны Python 3.11+, uv, just и около 16 ГБ свободной памяти. Для расширенного режима также требуется clang с бэкендом NVPTX. Команда just build загружает Qwen3.5-2B и компилирует GPU в неё. Предварительно собранная контрольная точка не опубликована. Собранная контрольная точка распространяется по лицензии Apache-2.0 Qwen.
Основные выводы
- qwenGPU компилирует 4-канальный SIMT-GPU в Qwen3.5-2B в виде 45 необученных слоёв.
- Состояние Gated DeltaNet служит регистрами, ПЗУ программы и глобальной памятью.
- Модель записывает каждый этап запуска CUDA; хост только компилирует и загружает.
- Интерпретатор совпал с Tesla T4 для 49 из 49 ядер.
- При скорости 10–15 токенов в секунду это исследовательское демо, а не среда выполнения.
Посмотрите репозиторий на GitHub, справочник qgpu CUDA, карточку модели Qwen3.5-2B и оригинальную публикацию в X. Также подписывайтесь на нас в Twitter и не забудьте присоединиться к нашему сабреддиту о машинном обучении с более чем 150 тыс. участников и подписаться на нашу рассылку. Стоп! Вы пользуетесь Telegram? теперь вы также можете присоединиться к нам в Telegram.
Хотите сотрудничать с нами для продвижения своего репозитория на GitHub, страницы на Hugging Face, выпуска продукта, вебинара и т. д.? Свяжитесь с нами
Переведено автоматически с английского. Оригинал статьи — по ссылке ниже.