Знайомтеся з qwenGPU: Qwen3.5-2B із 45 скомпільованими GPU-шарами, що виконують CUDA
qwenGPU — це експериментальна модель, яка компілює SIMT-графічний процесор безпосередньо у ваги Qwen3.5-2B. Розробник Vixhal (@TheVixhal) опублікував її на GitHub за ліцензією MIT. Модель запускає справжній CUDA, скомпільований у NVIDIA PTX, усередині власних прямих проходів. Для розробників це перевірка однієї ідеї: точні обчислення всередині попередньо навченої LLM, а не поруч із нею.
Коротко
- Розмір: Qwen3.5-2B (24 шари) разом із 45 скомпільованими GPU-шарами, загалом 69. Контекст: 262 144 токени (базова модель).
- Працює на: Apple silicon, графічному процесорі 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 — це 4-канальна SIMT-машина. Вона має 15 регістрів на канал, 4-етапний конвеєр, ROM на 32 блоки та 126 комірок пам’яті. Вона використовує лише власні типи шарів Qwen3.5: лінійну увагу Gated DeltaNet і MLP-блоки SwiGLU.
Секрет полягає в рекурентному стані Gated DeltaNet. У головках схеми вентиль утримується відкритим, а сила запису дорівнює 1. Кожен ключ є one-hot. Кожне оновлення перезаписує рівно 1 рядок, а one-hot-запит зчитує його назад. Тож кожен рядок працює як комірка пам’яті.
У цьому стані зберігаються ROM програми, регістри та глобальна пам’ять. У 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-графічним процесором усередині ваг, 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 авторства Дів’янша Агравала. Qwenputer компілює 32-бітний процесор RISC-V у Qwen3.5-4B. Percepta у березні 2026 року вбудувала інтерпретатор WebAssembly у трансформер із 2D-головками уваги.
| Функція | qwenGPU | Qwenputer (RISC-V) | Percepta |
|---|---|---|---|
| Базова модель | Qwen3.5-2B, 24 шари | Qwen3.5-4B | Спеціальний трансформер із 2D-головками уваги |
| Вбудована машина | 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) | Приблизно 3200 токенів трасування/с (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 виконує роль регістрів, ROM програми та глобальної пам’яті.
- Модель записує кожен крок запуску CUDA; хост лише компілює та завантажує.
- Інтерпретатор збігся з Tesla T4 у 49 із 49 ядер.
- За швидкості 10–15 токенів за секунду це дослідницька демонстрація, а не середовище виконання.
Перегляньте репозиторій на GitHub, довідник qgpu CUDA, картку моделі Qwen3.5-2B та оригінальний допис у X. Також підписуйтеся на нас у Twitter і не забудьте приєднатися до нашого сабреддіту про машинне навчання з понад 150 тисячами учасників та підписатися на нашу розсилку. Стривайте! Ви є в Telegram? тепер ви також можете приєднатися до нас у Telegram.
Хочете стати нашим партнером для просування вашого репозиторію на GitHub, сторінки на Hugging Face, випуску продукту, вебінару тощо? Зв’яжіться з нами
Перекладено автоматично з англійської. Оригінал статті — за посиланням нижче.