qwenGPU is an experimental model that compiles a SIMT GPU directly into the weights of Qwen3.5-2B. Developer Vixhal (@TheVixhal) released it on GitHub under the MIT license. The model runs real CUDA, compiled to NVIDIA PTX, inside its own forward passes. For developers, it tests one idea: exact computation inside a pretrained LLM, not beside it.
TL;DR
- Size: Qwen3.5-2B (24 layers) along with 45 compiled GPU layers, 69 in total. Context: 262,144 tokens (base model).
- Runs on: Apple silicon, an NVIDIA GPU or a CPU, with about 16 GB of free memory. No quantized build.
- Performance: Every launch in the developer’s runs was exact; text drift is a KL of about 7e-4.
- Best: The reference interpreter matched a real NVIDIA Tesla T4 on 49 of 49 kernels.
- Worst: About 10 to 15 tokens per second on Apple MPS; Qwen-written CUDA passed 4 of 5.
What is qwenGPU?
qwenGPU is Qwen3.5-2B with 45 compiled transformer layers placed in front of Qwen’s own 24. None of those layers are trained. A compiler computes every weight.
How does a GPU fit inside model weights?
The GPU is a 4-lane SIMT machine. It has 15 registers per lane, a 4-stage pipeline, a 32-block ROM and 126 memory words. It uses only Qwen3.5’s own layer types: Gated DeltaNet linear attention and SwiGLU MLPs.
The trick is the Gated DeltaNet recurrent state. In the circuit’s heads, the gate is held open and the write strength is 1. Every key is one-hot. Each update overwrites exactly 1 row, and a one-hot query reads it back. So each row works as a memory cell.
That state holds the program ROM, registers and global memory. The README lists 16 register rows per lane, with 1 that always reads zero. That matches the 15 usable registers in the X post.
The GPU supports integer arithmetic, predicates, divergent branches and loops. It also supports shared memory, atomicAdd, __syncthreads() and warp shuffles.