Vixhal released an experimental model on GitHub that compiles a SIMT GPU directly into the weights of Qwen3.5-2B. The project, named qwenGPU, runs real CUDA code inside the forward pass of the language model.
Technical specifications
The system combines Qwen3.5-2B, which contains 24 layers, with 45 additional compiled GPU layers. This results in a total of 69 layers. The base model supports a context window of 262,144 tokens.
It operates on Apple silicon, NVIDIA GPUs, or CPUs provided there is approximately 16 GB of free memory available. There is no quantized version of the build.
Performance testing showed exact execution for every launch. Text drift measured as a Kullback-Leibler divergence was about 7e-4. In a reference interpreter test, the system matched a real NVIDIA Tesla T4 on 49 of 49 kernels. Speed on Apple MPS ranged from 10 to 15 tokens per second. The Qwen-written CUDA code passed 4 of 5 tests in advanced mode.
How the architecture works
The implementation places 45 compiled transformer layers in front of Qwen’s original 24. None of these layers are trained. A compiler calculates every weight.
The GPU functions as a 4-lane SIMT machine. It features 15 registers per lane, a 4-stage pipeline, a 32-block ROM, and 126 memory words. The architecture relies solely on Qwen3.5’s existing layer types: Gated DeltaNet linear attention and SwiGLU MLPs.
The core mechanism uses the Gated DeltaNet recurrent state. In the circuit’s heads, the gate remains open and write strength is set to 1. Every key is one-hot. Each update overwrites exactly 1 row, and a one-hot query reads it back. This allows each row to function as a memory cell.
The state holds the program ROM, registers, and global memory. Documentation lists 16 register rows per lane, with one reserved to always read zero. This matches the 15 usable registers described in the developer’s post.
The GPU supports integer arithmetic, predicates, divergent branches, and loops. It also supports shared memory, atomicAdd, __syncthreads(), and warp shuffles.
Execution flow
The host system acts as NVRTC and the CUDA driver. It compiles CUDA to PTX with clang, then lowers PTX to machine blocks. The program and input arrays are fed to the model as tokens.
The model then writes the launch itself. Each EXEC(b, w) token runs block b for warp w in a single forward pass. The VAL tokens following it carry register results. Memory writes return as MEM records. The host reads results from the recurrent state and performs no computation on them.
In the demo video, a dot product of [1, 2, 3] and [4, 5, 6] returned 32. The model generated that launch in 47 tokens.
Meet qwenGPU: Qwen3.5-2B with a SIMT GPU inside its weights, 45 compiled layers in front of Qwen’s own, using its Gated DeltaNet layers as registers and memory.
It runs real CUDA in its forward passes. The host compiles the kernel to PTX and loads it, then the model writes every… pic.twitter.com/iM5IETEbIv
— vixhaℓ (@TheVixhal) November 1, 2026
Impact on language capabilities
The developer reports no degradation in text ability. The circuit is muted on text tokens, and Qwen’s own weights remain unchanged. The merged model’s next-token distribution stays within a KL of about 7e-4 of the original.
Accuracy results
Following numbers are self-reported, measured on Apple silicon with MPS in bf16:
- Every launch was exact under
--checkreplay. - The reference interpreter matched an NVIDIA Tesla T4 on 49 of 49 test kernels.
- Normal mode answered 20 of 24 sampled chat questions with GPU-computed numbers.
- Advanced mode wrote a correct kernel first try on 4 of 5 questions.
In normal mode, Qwen writes an expression like [1, 2, 3] @ [4, 5, 6] for ready-made kernels. In advanced mode, Qwen writes the CUDA kernel itself. The README warns that a 2B model’s CUDA is still unreliable.
Comparison with similar projects
qwenGPU uses the transformer compiler from Qwenputer by Divyansh Agrawal. Qwenputer compiles a 32-bit RISC-V CPU into Qwen3.5-4B. Percepta built a WebAssembly interpreter into a transformer with 2D attention heads in March 2026.
| Feature | qwenGPU | Qwenputer (RISC-V) | Percepta |
|---|---|---|---|
| Base model | Qwen3.5-2B, 24 layers | Qwen3.5-4B | Custom transformer with 2D attention heads |
| Machine compiled in | 4-lane SIMT GPU | 32-bit RISC-V (RV32IM + Zba + bignum) | WebAssembly interpreter |
| Programs it runs | CUDA, lowered from PTX | C and MicroPython | C, compiled to WebAssembly |
| Compiled layers | 45 | 11 | Not disclosed |
| Where state lives | Gated DeltaNet state | Gated DeltaNet state | Execution trace in context |
| Host role | Compiles CUDA to PTX, loads tokens, reads results | Loads firmware; runs no computation | Not disclosed |
| Correctness check | Interpreter matched Tesla T4 on 49 of 49 kernels | 0 mismatches in about 230,000 replayed positions | Claims millions of reliable steps |
| Reported speed | 10 to 15 tokens/s (Apple MPS) | About 3,200 trace tokens/s (L4, Triton) | 30,000+ tokens/s (CPU) |
| Text drift vs base | KL about 7e-4 | KL 6e-4 | Not applicable |
| License | MIT code; Apache-2.0 checkpoint | MIT code; Apache-2.0 weights | Not disclosed |
The speed figures use different hardware and token types, so they are not directly comparable.
Setup requirements
The quickstart needs Python 3.11+, uv, just, and about 16 GB of free memory. Advanced mode also requires clang with the NVPTX backend. Running just build downloads Qwen3.5-2B and compiles the GPU into it. No prebuilt checkpoint is published. A built checkpoint falls under Qwen’s Apache-2.0 license.
For those interested in the project, the GitHub repository is available at vixhal-baraiya/qwenGPU. The developer also posted the original announcement on X.
