qwenGPU: autor wkompilował w wagi Qwena „procesor graficzny”, który wykonuje kod CUDA

Projekt qwenGPU (w kodzie: qgpu) to model Qwen3.5-2B z 45 dodatkowymi warstwami transformera, ustawionymi przed jego własnymi 24. Żadna z tych warstw nie jest trenowana. Każdą wagę wylicza kompilator, a razem warstwy tworzą maszynę SIMT, która wykonuje prawdziwy kod CUDA skompilowany do PTX. Autor, Vixhal, zapowiadał go na X. Repozytorium jest na licencji MIT. W chwili sprawdzania repozytorium miało 63 gwiazdki i 9 forków.

Mały GPU zbudowany z warstw Qwena

Według opisu w repozytorium „GPU” jest bardzo skromny: 4 ścieżki (lane), 15 rejestrów na ścieżkę, czterostopniowy potok, 32-blokowa pamięć programu i 126 słów pamięci globalnej. Obsługuje arytmetykę całkowitoliczbową, predykaty, rozgałęzienia rozbieżne, pętle, pamięć współdzieloną, atomicAdd, __syncthreads() i warp shuffle.

Całość zbudowano wyłącznie z typów warstw, które Qwen3.5 już ma: liniowej uwagi Gated DeltaNet i MLP typu SwiGLU. Stan maszyny trafia do stanu rekurencyjnego warstw DeltaNet. W głowicach układu bramka jest stale otwarta, siła zapisu wynosi 1, a każdy klucz jest „one-hot”. Każda aktualizacja nadpisuje więc dokładnie jeden wiersz stanu, a zapytanie one-hot odczytuje go z powrotem. Wiersz działa jak komórka pamięci. W takich komórkach siedzą pamięć programu, rejestry i pamięć globalna.

Kto co robi podczas uruchomienia

Host pełni rolę NVRTC i sterownika CUDA. Kompiluje źródło CUDA do PTX za pomocą clanga, zamienia PTX na bloki maszynowe i podaje modelowi program oraz dane wejściowe jako tokeny. Potem model sam zapisuje przebieg: wygenerowane tokeny wskazują m.in., który blok wykonał dany warp i jakie wartości powstały, aż do tokenu końca uruchomienia. Według autora host niczego nie liczy, a wyniki tylko odczytuje ze stanu modelu.

Dodanie tak obcego obwodu nie powinno psuć zwykłego tekstu. Obwód jest wyciszany na tokenach tekstowych, a wagi Qwena pozostają nietknięte. Rozkład następnego tokenu scalonego modelu mieści się w KL rzędu 7e-4 względem oryginału.

Do budowy autor użył kompilatora transformerów z projektu Qwenputer Divyansha Agrawala, który wkompilowuje procesor RISC-V w Qwen3.5-4B. qwenGPU korzysta z niego jako z biblioteki i dokłada własny GPU, sterownik oraz format tokenów. W odróżnieniu od Qwenputera host reaguje tu na wywołanie narzędzia, tak jak strona hosta w prawdziwym GPU.

Dwa tryby i wyniki

Czat w terminalu działa w dwóch trybach. W normalnym Qwen zapisuje działanie jako wyrażenie, np. [1, 2, 3] @ [4, 5, 6], a każda operacja idzie na gotowy kernel z katalogu kernels/. W zaawansowanym model sam pisze kernel CUDA, korzystając z opisu w prompcie systemowym. Obok czatu widać panel pokazujący na żywo pracę harmonogramu warpów, pamięci programu, potoku i rejestrów.

• Autor podaje, że każde uruchomienie w jego testach było dokładne. Z opcją --check tokeny modelu porównuje się 1:1 z interpreterem referencyjnym obwodu.
• Ten interpreter zgodził się z prawdziwym NVIDIA Tesla T4 w 49 z 49 kernelów testowych.
• Tryb normalny odpowiedział liczbami policzonymi na GPU w 20 z 24 próbkowanych pytań.
• Tryb zaawansowany napisał poprawny kernel za pierwszym razem w 4 z 5 pytań o wektory i macierze (iloczyn skalarny, iloczyn macierzy, kwadrat długości, dodawanie wektorów).

Wszystkie liczby pochodzą od autora i zmierzono je na Apple silicon (MPS, bf16). Sam README zastrzega, że CUDA pisana przez model 2B nadal bywa zawodna i do wyników, na których ktoś polega, zaleca tryb normalny. Jedno uruchomienie to od kilkudziesięciu do kilkuset wygenerowanych tokenów, przy tempie około 10–15 tokenów na sekundę. W demonstracji autora iloczyn skalarny [1, 2, 3] i [4, 5, 6] daje 32, a model zapisuje przebieg w 47 tokenach.

Do uruchomienia potrzeba Pythona 3.11+, narzędzi uv i just oraz około 16 GB wolnej pamięci na model 2B. Działa na Apple silicon, GPU NVIDII lub CPU. Tryb zaawansowany wymaga dodatkowo clanga z backendem NVPTX. Komenda just build pobiera Qwen3.5-2B i wkompilowuje w niego GPU. Kod jest na MIT, ale checkpoint zbudowany przez qgpu.merge to zmodyfikowany Qwen3.5, więc jego wagi podlegają licencji Apache-2.0 Qwena.