qwenGPU用模型内置层模拟GPU执行
qwenGPU将四通道SIMT GPU编译为Qwen3.5-2B中的45个无需训练的层,并利用Gated DeltaNet状态表示寄存器和内存。据介绍,模型可自行推演CUDA启动的各个执行步骤。
AI 深度解读
qwenGPU 把一台微型 GPU 的执行逻辑编译进了 Qwen3.5-2B 的权重,让模型在前向计算中执行 CUDA。它想验证的是:大模型能否在内部完成确定性计算,而不只是生成答案,或者调用外部计算工具。
原文给出的结果有吸引力:开发者报告执行回放全部精确,文本输出分布变化也很小。但它目前仍是研究演示。读这项工作的价值,是看懂「模型写程序」和「模型执行程序」如何分开,以及为什么执行正确不等于回答可靠。
45 层没训练的网络,怎么变成 GPU?
开发者 Vixhal 在 Qwen3.5-2B 原有的 24 层前面,加入了 45 层编译生成的 transformer,总计 69 层。新增层没有经过训练,每个权重都由编译器计算出来。
这里的 GPU 是一台用网络实现的 SIMT 机器,也就是让多个执行通道按照同一套指令工作。它只有 4 个通道,每个通道有 15 个可用寄存器,配备 4 级流水线、32 块 ROM 和 126 个内存字。README 写的是每通道 16 个寄存器行,其中一个恒为零,与 15 个可用寄存器并不矛盾。
关键在 Gated DeltaNet 的循环状态。电路使用的注意力头保持门控开启,将写入强度设为 1,再用 one-hot key 精确覆盖某一行,用 one-hot query 读回那一行。可以把它理解为:给状态中的每一行分配固定地址,让网络按地址存取数据。
程序 ROM、寄存器和全局内存就放在这些状态里。实现只使用 Qwen3.5 自身已有的 Gated DeltaNet 线性注意力和 SwiGLU MLP 层,却支持整数运算、谓词、分支分歧和循环,以及 shared memory、atomicAdd、__syncthreads() 和 warp shuffles。
这里的确定性来自编译出的电路,而不是模型学会了可靠心算。
CUDA 到底由谁执行,宿主又做了什么?
宿主仍然负责重要的准备工作:扮演 NVRTC 和 CUDA driver 的角色,用 clang 把 CUDA 编译成 NVIDIA PTX,再把 PTX 降成机器块,将程序和输入数组编码为 token。
随后才轮到模型执行。每个 EXEC(b, w) token 在一次前向计算中,为 warp w 执行 block b;后续 VAL token 携带寄存器结果,内存写入则以 MEM 记录返回。宿主读取循环状态里的结果,不负责计算这些结果。
因此,「在模型内部运行 CUDA」有明确边界:编译和加载仍在外部,执行逻辑进入了模型。演示中,[1, 2, 3] 与 [4, 5, 6] 的点积返回 32,模型用了 47 个 token 写出这次执行过程。
普通模式与高级模式的区别也在这里。普通模式让 Qwen 写出类似 [1, 2, 3] @ [4, 5, 6] 的表达式,交给预制 kernel;高级模式让 Qwen 自己编写 CUDA kernel。前者限制了程序生成空间,后者给模型更多自由,也引入更多出错机会。
原文没有提供完整代码片段,不能仅凭这些执行标记还原接口实现。
执行精确,为什么回答仍然会错?
原文中的测试均为开发者自报,环境是 Apple silicon、MPS 和 bf16,需要分开看:
- 在开发者运行的测试中,每次 launch 都通过了
--check replay精确回放检查。 - 参考解释器在 49 个测试 kernel 上,与 NVIDIA Tesla T4 的结果全部一致。
- 普通模式抽样的 24 个聊天问题中,20 个回答使用了 GPU 计算的数字。
- 高级模式的 5 个问题中,4 个首次生成的 kernel 正确。
前两项支持执行机制的正确性,后两项涉及模型如何理解问题、选择计算和编写程序。49/49 的对照对象是参考解释器,也不能直接扩写成整个模型在所有 CUDA 程序上都经过硬件验证。
执行器可以准确运行程序,但程序本身仍可能写错。
文本能力的证据也有范围。新增电路在文本 token 上静默,Qwen 原有权重保持不变;合并模型与原模型的下一 token 分布,KL 散度约为 7e-4。这个数字表示输出概率分布接近,并不能单独证明所有问答、编码和长上下文能力都没有变化。基础模型的上下文长度为 262,144 tokens,原文没有给出合并后对应的完整能力评测。
运行门槛不算低。项目支持 Apple silicon、NVIDIA GPU 或 CPU,需要约 16 GB 空闲内存,没有量化构建,也没有发布预构建 checkpoint。快速开始需要 Python 3.11+、uv 和 just,执行 just build 下载 Qwen3.5-2B 并编译 GPU;高级模式还需要带 NVPTX backend 的 clang。代码采用 MIT,构建后的 checkpoint 遵循 Qwen 的 Apache-2.0。
我的判断
我认为最有价值的是它把语言生成与确定性执行接在了一起,而且没有改动原模型权重。它展示了一种可检查的计算结构,比单纯要求模型「认真算一遍」更有工程解释力。
但它目前不能替代常见的工具调用。日常 AI 编程中,让模型生成表达式或代码,再交给外部执行器,通常更容易接入、调试和隔离故障。qwenGPU 是否能在可靠性、成本或延迟上胜过这条路线,原文没有对照数据,不能替它下结论。
速度尤其需要克制解读。qwenGPU 在 Apple MPS 上约为 10 到 15 tokens/s;它借用编译器的 Qwenputer 将 32-bit RISC-V 编入 Qwen3.5-4B,报告在 L4、Triton 上约 3,200 trace tokens/s;2026 年 3 月的 Percepta 则使用带 2D attention heads 的定制 transformer,在 CPU 上报告 30,000+ tokens/s。
这些数字的硬件和 token 含义不同,无法直接排名,也不能当成同一种计算吞吐量。qwenGPU 的现阶段定位,仍是验证思路的研究演示。
能拿来做什么
- 做 AI 编程助手时,把「程序是否生成正确」和「程序是否执行正确」分开记录。只看最终答案,会掩盖错误究竟出在哪一步。
- 处理可枚举的计算任务时,可以借鉴普通模式:先提供有限的预制 kernel 或函数,让模型填写表达式,再评估是否需要开放自由代码生成。
- 想复现实验,可以先确认约 16 GB 空闲内存和 Python 3.11+、uv、just 环境,运行
just build;准备测试高级模式时,再检查 clang 是否包含 NVPTX backend。 - 阅读类似项目的跑分时,先核对测试对象、硬件与 token 定义。参考解释器的 49/49、生成 kernel 的 4/5 和聊天样本的 20/24,分别回答不同问题。