QwenGPU：在语言模型内部构建一个可执行CUDA的SIMT架构意味着什么？

开场：模型内生计算单元的最新尝试

在人工智能领域，大型语言模型（LLM）的算术和确定性推理能力一直是其“短板”。传统上，解决这一问题的方法是让LLM调用外部工具或解释器，将复杂的计算任务外包出去。然而，开发者Vixhal最近发布了一个实验项目qwenGPU，挑战了这一范式。该项目并非让模型调用外部GPU，而是将一个4-lane SIMT（单指令多线程）GPU的完整架构直接编译到了Qwen3.5-2B模型的权重之中，使其能够在其自身的正向推理过程中运行真实的CUDA代码，执行精确的数值运算。

这一大胆尝试的核心在于，它将计算单元内化于模型本身，而非简单地作为模型之外的插件或工具。qwenGPU沿袭了此前Qwenputer将RISC-V CPU编译进LLM的思路，但更进一步，实现了更具通用性的GPU架构。对于技术负责人而言，这项工作的意义在于，它试图解决LLM在面对复杂、确定性计算时固有的“幻觉”和不稳定性问题，预示着一种全新的模型内部计算范式。这不仅仅是提升了LLM的计算能力，更是对LLM作为通用计算基座潜力的深度探索，试图让模型不仅“理解”世界，还能“精确执行”世界。

核心机制：Gated DeltaNet 的巧妙借用

qwenGPU的核心创新在于，它没有通过训练来“学习”一个GPU，而是通过编译器计算出了45个专门的Transformer层权重，并将其放置在Qwen3.5-2B原有的24层之前，使得总层数达到69层。这些新增层不是用来处理语言任务的，而是作为模拟SIMT GPU的硬件电路。为了实现GPU所需的寄存器、程序存储（ROM）和全局内存，开发者巧妙地利用了Qwen3.5模型中Gated DeltaNet线性注意力机制的循环状态。

Gated DeltaNet的独特之处在于其门控机制和更新方式。通过将门控（gate）始终保持打开、写入强度（write strength）设置为1，并使用one-hot编码的key和query，每一次更新都能精确地覆盖状态矩阵中的一行，而每一次读取也都能精确地检索到对应的那一行。这意味着，Gated DeltaNet的每一行循环状态都可以被精确地作为独立的内存单元来使用，从而存储GPU的程序ROM、各通道的寄存器（每lane 15个可用寄存器，外加一个永读零的行）以及126个32位字大小的全局内存。这种设计不仅支持整数算术、谓词、分支与循环控制，还实现了共享内存、原子加操作（`atomicAdd`）以及线程同步（`__syncthreads()`）和warp shuffle等高级GPU特性，使其能够在一个4级流水线的SIMT架构中运行。

执行流程：主机与模型的协同边界

尽管qwenGPU将计算单元内化于模型，但其运行并非完全脱离外部环境。整个CUDA代码的执行流程是一个主机与模型协同工作的过程。首先，主机程序扮演了NVRTC和CUDA驱动的角色：它使用`clang`编译器将CUDA源代码编译成PTX（Parallel Thread Execution）汇编代码，随后将PTX进一步降低为qwenGPU内部SIMT机器可识别的指令块。完成编译后，主机将这些程序指令和输入数据作为token序列馈送给Qwen模型，由模型进行处理。

接下来，关键步骤发生在模型内部。模型会根据输入的指令序列，生成一系列的执行记录token，例如`EXEC(b, w)`表示执行第`b`个block的第`w`个warp，`VAL`token携带寄存器运算结果，而`MEM`记录则反映内存写入操作。在整个计算过程中，主机仅负责读取这些执行记录和最终的计算结果（存储在模型的循环状态中），但**不参与具体的内核运算**。这意味着CUDA内核的实际逻辑运算是由模型内部的编译电路完成的，而非在主机上模拟。需要明确的是，这里“运行真实CUDA”指的是经编译和降低的CUDA指令由模型内部电路执行，而不是这些指令直接在NVIDIA GPU硬件上运行。此外，在“模型写CUDA”的表述上，也存在模式区分：在普通模式下，模型可能仅是写出算术表达式并调用仓库中预置的CUDA内核；而在高级模式下，模型才尝试从头开始“编写”自己的CUDA内核，展示更深层次的自主性。

性能与准确性：量化证据

开发者公布的自我报告数据显示，qwenGPU在概念验证层面表现出了惊人的准确性，尽管其性能指标尚不具备实际应用级的竞争力。以下是一些关键的量化证据：

**计算准确性**： * 在开发者测试中，参考解释器（reference interpreter）与真实的NVIDIA Tesla T4 GPU在**49个测试内核中的49个**都达到了精确匹配。在`--check replay`模式下，每一次CUDA launch的计算结果都被报告为精确无误。这表明，一旦指令被模型内部的电路执行，其数值结果是高度确定和准确的。

**语言能力影响**： * 将45个编译层添加到Qwen3.5-2B之后，合并模型的下一词预测分布与原始模型相比，KL散度仅为约**7×10⁻⁴**。这说明，尽管模型新增了大量用于计算的层，但其原有的文本生成和语言理解能力并未受到显著破坏。

**计算性能**： * 在Apple silicon（MPS）平台使用bf16精度进行测试时，模型生成token的速度约为**10–15 tokens/s**。一次CUDA launch通常需要几十到几百个生成token来完成。这与现代GPU的计算速度相比非常缓慢，也凸显了当前这种内生计算方法的效率瓶颈。

**高级模式表现**： * 在模型首次尝试编写CUDA内核的高级模式测试中，它在**5个问题中成功答对4个**。而在普通模式下（即模型写算术表达式并调用预置内核），在抽样的24个问题中，有20个问题给出了由GPU计算的正确数字。这些结果表明，模型不仅能执行，甚至在一定程度上能生成或适应所需的计算逻辑。

需要注意的是，以上所有数据均为开发者Vixhal的自我报告，测试环境为Apple silicon上的MPS后端，并未使用量化版本。虽然这些数据强有力地支持了qwenGPU在理论和概念上的可行性，但在将此技术应用于生产环境之前，仍需要独立的性能评估和更广泛的硬件平台测试。

工程考量与实际意义

qwenGPU的出现，迫使我们重新思考大型语言模型在未来系统架构中的角色。它提出了一个诱人的前景：如果LLM能够直接在内部执行精确的、确定性的计算，那么它将能够更好地处理数值任务、科学计算、代码执行验证等领域，从而减少对外部工具的依赖，并可能提升响应速度和结果可靠性。在当前LLM容易“一本正经地胡说八道”的背景下，这种内生计算能力有望成为解决模型幻觉、提高事实准确性的关键一步。

然而，这种内生计算路径也带来了显著的工程挑战和权衡。首先是模型规模的急剧增加——一个2B参数的基础模型，在添加45个非训练层后，总层数几乎增加了两倍。这直接导致了更高的内存需求（约16GB），并可能带来更大的推理延迟和计算成本。其次，目前的性能指标，例如每秒10-15 token的执行速度，与现代GPU的强大并行处理能力相比微不足道。对于大多数实际的、计算密集型任务，现阶段的qwenGPU显然无法替代传统GPU。此外，虽然计算发生在模型内部，但仍需要一个主机程序来完成CUDA到PTX的编译、PTX到机器指令的降低，以及将程序和输入作为token馈送给模型，并最终从模型状态中读取结果。这并非一个完全自治的系统，而是一个紧密耦合的混合架构。

尽管存在这些限制，qwenGPU仍为将来的AI代理提供了新的想象空间。如果模型能够将核心的、可验证的计算逻辑作为其“思维”的一部分，而不是作为外部API调用，那么这些代理将能够进行更深层次的规划、自我纠正和验证。它为模型内集成领域特定语言（DSL）解释器或小型虚拟机提供了蓝图，从而实现更强大的可编程性和智能行为。对于需要高度精确性和可追溯性的场景，如金融建模、药物发现或工业自动化，这种内生计算能力可能成为未来AI系统的重要组成部分。

边界与未来展望

qwenGPU项目在概念上取得了显著突破，但其当前形态并非用于直接替代传统GPU或作为通用计算平台。其限制是显而易见的：相对较低的吞吐量、庞大的模型体积（相对于其所处理的任务）、有限的内存与寄存器空间，以及对主机程序在编译和I/O环节的持续依赖。这些因素共同表明，qwenGPU目前是一个强大的实验工具和概念验证，而非一个生产就绪的解决方案。它清晰地划定了当前Transformer架构在集成通用计算硬件时的能力边界。

然而，这并不意味着该技术不重要。相反，它为我们理解如何将复杂、确定性的计算逻辑融入LLM的核心架构提供了宝贵的经验。未来的研究可能会专注于如何优化这些编译层，例如通过更高效的架构设计、更紧凑的指令集，或者结合稀疏激活等技术来降低计算成本和内存占用。另一个重要的方向是探索如何减少对主机编译和I/O的依赖，实现更深层次的自主性，使模型能够更自发地将高层任务分解为可执行的低层计算指令。这项工作预示着，Transformer可能不仅仅是一个语言处理器，而是一个更具通用性、可编程性的智能实体，能够执行从抽象推理到精确数值计算的全范围任务。对于致力于构建更鲁棒、更智能、更少幻觉的下一代AI系统的技术团队来说，qwenGPU指明了一条值得长期关注和投入的研究路径，它在根本上改变了我们对“模型能力”的定义和期待。