QwenGPU: What Does Building a CUDA-Executable SIMT Architecture Inside an LLM Imply?

Opening: The Latest Attempt at In-Model Computation Units

In the field of artificial intelligence, the arithmetic and deterministic reasoning capabilities of large language models (LLMs) have traditionally been their Achilles' heel. Conventionally, this issue is addressed by having LLMs call external tools or interpreters, outsourcing complex computational tasks. However, developer Vixhal recently released an experimental project, qwenGPU, which challenges this paradigm. Instead of having the model call an external GPU, this project compiles the complete architecture of a 4-lane SIMT (Single Instruction, Multiple Thread) GPU directly into the weights of the Qwen3.5-2B model. This allows it to execute real CUDA code within its own forward passes, performing precise numerical operations.

The core of this ambitious attempt lies in internalizing the computational unit within the model itself, rather than simply treating it as an external plugin or tool. qwenGPU follows the spirit of the earlier Qwenputer project, which compiled a RISC-V CPU into an LLM, but goes further by implementing a more general-purpose GPU architecture. For technical leads, the significance of this work is its attempt to address the inherent 'hallucinations' and instability LLMs face with complex, deterministic computations. It heralds a new paradigm for in-model computation. This not only enhances the LLM's computational power but also deeply explores the potential of LLMs as general-purpose computational substrates, aiming for models that can not only 'understand' the world but also 'execute precisely' within it.

Core Mechanism: The Clever Use of Gated DeltaNet

The core innovation of qwenGPU lies in not 'learning' a GPU through training, but rather by compiling 45 specialized Transformer layer weights. These are placed before Qwen3.5-2B's original 24 layers, bringing the total to 69 layers. These new layers are not designed for language tasks but serve as the hardware circuit emulating a SIMT GPU. To implement the registers, program ROM, and global memory required by the GPU, the developer ingeniously leverages the recurrent state of the Gated DeltaNet linear attention mechanism found in the Qwen3.5 model.

The uniqueness of Gated DeltaNet stems from its gating mechanism and update method. By keeping the gate perpetually open, setting the write strength to 1, and using one-hot encoded keys and queries, each update precisely overwrites one row in the state matrix, and each read retrieves exactly that corresponding row. This implies that each row of Gated DeltaNet's recurrent state can be precisely utilized as an independent memory cell. This allows it to store the GPU's program ROM, registers for each lane (15 usable registers per lane, plus one row that always reads zero), and 126 32-bit words of global memory. This design supports integer arithmetic, predicates, divergent branches, loops, and even advanced GPU features like shared memory, atomic additions (`atomicAdd`), thread synchronization (`__syncthreads()`), and warp shuffles, enabling it to operate within a 4-stage pipelined SIMT architecture.

Execution Flow: The Collaborative Boundary Between Host and Model

Although qwenGPU internalizes the computational unit within the model, its operation is not entirely decoupled from the external environment. The execution flow of CUDA code involves a collaborative process between the host and the model. First, the host program acts as an NVRTC and CUDA driver: it uses the `clang` compiler to compile CUDA source code into PTX (Parallel Thread Execution) assembly. Subsequently, it further lowers the PTX into instruction blocks recognizable by qwenGPU's internal SIMT machine. After compilation, the host feeds these program instructions and input data as token sequences to the Qwen model for processing.

The crucial steps then occur within the model. Based on the input instruction sequence, the model generates a series of execution record tokens, such as `EXEC(b, w)` to denote the execution of block `b`'s warp `w`, `VAL` tokens carrying register computation results, and `MEM` records reflecting memory write operations. Throughout this entire computational process, the host is solely responsible for reading these execution records and the final computed results (stored in the model's recurrent state), but it **does not participate in the actual kernel computations**. This means the real logical operations of the CUDA kernel are performed by the compiled circuit inside the model, not simulated on the host. It is important to clarify that 'running real CUDA' here refers to the execution of compiled and lowered CUDA instructions by the in-model circuit, not direct execution on NVIDIA GPU hardware. Furthermore, regarding the phrase 'model writes C

Performance and Accuracy: Quantified Evidence

Self-reported data released by the developer shows that qwenGPU exhibits impressive accuracy at the proof-of-concept level, although its performance metrics are not yet competitive for practical applications. Here is some key quantified evidence:

**Computational Accuracy**: * In developer tests, the reference interpreter precisely matched a real NVIDIA Tesla T4 GPU on **49 out of 49 test kernels**. Every CUDA launch's computational result was reported as exact under the `--check replay` mode. This indicates that once instructions are executed by the model's internal circuit, their numerical results are highly deterministic and accurate.

**Impact on Language Ability**: * After adding 45 compiled layers to Qwen3.5-2B, the merged model's next-token prediction distribution showed a KL divergence of only about **7×10⁻⁴** compared to the original model. This suggests that despite the addition of numerous computational layers, the model's original text generation and language understanding capabilities were not significantly compromised.

**Computational Performance**: * When tested on Apple silicon (MPS) with bf16 precision, the model generated tokens at approximately **10–15 tokens/s**. A single CUDA launch typically required tens to hundreds of generated tokens to complete. This is very slow compared to the computational speed of modern GPUs, highlighting the efficiency bottleneck of the current in-model computation approach.

**Advanced Mode Performance**: * In advanced mode tests, where the model first attempted to write CUDA kernels, it successfully answered **4 out of 5 q

Engineering Considerations and Practical Implications

The emergence of qwenGPU forces us to reconsider the role of large language models in future system architectures. It presents an enticing prospect: if LLMs could directly perform precise, deterministic computations internally, they would be better equipped to handle numerical tasks, scientific computing, code execution verification, and other domains, thereby reducing reliance on external tools and potentially improving response speed and result reliability. In the current context where LLMs are prone to 'confidently hallucinating,' this intrinsic computational capability holds promise as a critical step toward resolving model hallucinations and enhancing factual accuracy.

However, this path of in-model computation also introduces significant engineering challenges and tradeoffs. First, there's the drastic increase in model size—a base 2B parameter model nearly doubles its total layers after adding 45 untrained layers. This directly leads to higher memory requirements (around 16GB) and potentially greater inference latency and computational costs. Second, current performance metrics, such as 10-15 tokens per second execution speed, are negligible compared to the powerful parallel processing capabilities of modern GPUs. For most practical, compute-intensive tasks, qwenGPU at this stage clearly cannot replace traditional GPUs. Furthermore, while computation occurs within the model, a host program is still required to compile CUDA to PTX, lower PTX to machine instructions, feed the program and inputs as tokens to the model, and ultimately read results from the model's state.

Boundaries and Future Outlook

The qwenGPU project has achieved a significant conceptual breakthrough, but its current form is not intended to directly replace traditional GPUs or serve as a general-purpose computing platform. Its limitations are apparent: relatively low throughput, a large model size (relative to the tasks it handles), limited memory and register space, and ongoing reliance on a host program for compilation and I/O. These factors collectively indicate that qwenGPU is currently a powerful experimental tool and a proof-of-concept, rather than a production-ready solution. It clearly defines the current boundaries of Transformer architecture capabilities when integrating general computing hardware.

However, this does not diminish the technology's importance. On the contrary, it provides invaluable experience for understanding how to embed complex, deterministic computational logic into the core architecture of LLMs. Future research may focus on optimizing these compiled layers, perhaps through more efficient architectural designs, more compact instruction sets, or by incorporating techniques like sparse activation to reduce computational costs and memory footprint. Another critical direction is exploring how to lessen reliance on host compilation and I/O, achieving deeper autonomy, allowing the model to more spontaneously break down high-level tasks into executable low-level computational instructions. This work suggests that the Transformer might not just be a language processor, but a more versatile, programmable intelligent entity capable of performing a full range of tasks from abstrac