Meet qwenGPU: Qwen3.5-2B With 45 Compiled GPU Layers That Run CUDA

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.

What happens during a CUDA launch?

The host plays the role of NVRTC and the CUDA driver. It compiles CUDA to PTX with clang, then lowers PTX to machine blocks. It feeds the program and input arrays to the model as tokens.

From there, the model writes the launch itself. Each EXEC(b, w) token runs block b for warp w in 1 forward pass. The VAL tokens after it carry register results. Memory writes return as MEM records. The host reads results from the recurrent state and computes none of them.

In the demo video, a dot product of [1, 2, 3] and [4, 5, 6] returns 32. The model wrote that launch in 47 tokens.

Does the GPU break Qwen’s text ability?

Not by the developer’s measure. The circuit is muted on text tokens, and Qwen’s own weights are unchanged. The merged model’s next-token distribution stays within a KL of about 7e-4 of the original.

How accurate is qwenGPU?

All numbers are self-reported, measured on Apple silicon with MPS in bf16:

• Every launch was exact under --check replay.
• 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.

How does qwenGPU compare with Qwenputer and Percepta?

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.

The speed figures use different hardware and token types, so they are not directly comparable.

What do you need to run it?

The quickstart needs Python 3.11+, uv, just and about 16 GB of free memory. Advanced mode also needs clang with the NVPTX backend. 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.

Key Takeaways

• qwenGPU compiles a 4-lane SIMT GPU into Qwen3.5-2B as 45 untrained layers.
• Gated DeltaNet state serves as registers, program ROM and global memory.
• The model writes every CUDA launch step; the host only compiles and loads.
• The interpreter matched a Tesla T4 on 49 of 49 kernels.
• At 10 to 15 tokens per second, it is a research demo, not a runtime.

Need to partner with us for promoting your GitHub Repo OR Hugging Face Page OR Product Release OR Webinar etc.? Connect with us