开发者 Vixhal 在 GitHub 上以 MIT 许可证开源了实验项目 qwenGPU,用编译器直接算出 45 层 Transformer 权重,前置拼接到原本 24 层的 Qwen3.5-2B 之前,让这个 69 层的合并模型在自身前向传播里直接运行编译为 NVIDIA PTX 的 CUDA 代码。这并不意味着大模型内部已经长出了一块能跑深度学习的显卡,翻开项目底层技术文档就会发现,这台塞进权重里的微型机器完全不支持浮点数(无 float 或 double),全局内存只有 126 个字,在 Apple silicon(MPS,bf16)上的生成速度只有 10 到 15 tokens/s。
它真正立住的价值,是像硬件工程师刻录 FPGA 一样,不靠任何梯度训练,直接把 Qwen3.5 的线性注意力状态改造成了可寻址、可擦写的确定性寄存器与内存。对于大模型架构和机制可解释性研究者,这是一份把复杂并行状态机硬编码进混合线性注意力模型的解剖样本;但对于追求工程效率的应用开发者,老老实实让模型调用外部 Python 解释器或物理 GPU,依然是唯一理性的选择。
劫持线性注意力状态:不靠训练怎么在权重里刻出一台 GPU
大语言模型做不好精确算术,行业通行解法是外部工具调用(Tool Use),即模型写代码、外部沙盒跑结果。但自 2026 年初以来,把确定性计算电路直接编译进 Transformer 权重矩阵的路线突然提速。3 月,Percepta 团队利用 2D 注意力头在 Transformer 里搭出了 WebAssembly 解释器;随后开发者 Divyansh Agrawal 发布 Qwenputer,把一颗 32 位 RISC-V CPU 编译进 Qwen3.5-4B 的 11 个附加层里;qwenGPU 则直接复用了 Qwenputer 的编译器框架,首次把单指令多线程(SIMT)的并行 GPU 架构塞进了模型权重。
为什么 Qwenputer 和 qwenGPU 接连选中阿里的 Qwen3.5 系列作为宿主?秘密不在它的语言能力,而在于 Qwen3.5 采用的 Gated DeltaNet 线性注意力与 SwiGLU MLP 混合架构。
在标准语言模型里,注意力状态会随上下文推进而混合衰减;但在 qwenGPU 编译出的那 45 层电路头中,开发者把衰减门(gate)恒定敞开,将写入强度死死锁在 1,并全部采用 one-hot(独热编码)作为键(Key)和查询(Query)。这样一来,每次状态更新都会精准覆写矩阵里的唯一一行,每次查询又能原封不动地把这一行读出来。神经网络里原本模糊的概率状态,就这样被硬生生改造成了可寻址的物理内存单元,用来存放程序 ROM、寄存器和全局内存。
在这个被劫持出来的状态空间里,Vixhal 塞进了一台 4 通道(4-lane)SIMT 机器,包含 4 级流水线、支持谓词执行、分支分化、循环、共享内存、atomicAdd 原子加、__syncthreads() 线程同步以及 warp 洗牌指令。当执行一次 CUDA 启动时,宿主端(Host)扮演 NVRTC 编译器和驱动的角色,用 clang 把 CUDA 编译成 PTX 再降级为机器块,当作普通 token 喂给模型。随后模型自己接管执行,每吐出一个 EXEC(b, w) token 就在单次前向传播里跑完对应线程块与 warp 的计算,接着吐出携带寄存器结果的 VAL token 和写回内存的 MEM 记录。演示视频里算一组 [1, 2, 3] 与 [4, 5, 6] 的向量点积得出 32,模型一共写了 47 个 token,宿主端全程只负责装载程序和从循环状态里读取最终答案,不参与任何算术运算。
剥离 100% 匹配 Tesla T4 的光环:无浮点数、百字内存与降速代价
看到跑通真实 CUDA 且在 49 个测试内核上与 NVIDIA Tesla T4 结果 100% 匹配的宣传,很容易让人产生大模型已经能在权重里自我加速矩阵运算的错觉。但如果仔细核查项目仓库的底层规范文档(docs/qgpu-cuda.md),就会发现这台权重 GPU 有着严苛的物理边界。
能跑通 CUDA 控制流,不等于能跑深度学习算力。
最大的限制在于数据类型。qwenGPU 仅支持 32 位整数(有符号 int 与无符号 unsigned,采用环绕溢出),底层电路完全没有编译任何浮点数算子(无 float 或 double),也不支持 64 位整数。一套不能算浮点数的 GPU,自然无法承担任何现代 AI 训练或科学计算任务,只能跑整数向量点积、前缀和这类基础算法。
其次是极度袖珍的硬件规格与两阶段验证口径。这台 4 通道机器的每个线程只有 15 个可用寄存器(外加 1 个恒为零的寄存器行,共 16 行),程序 ROM 上限是 32 个机器块,全局内存总共只有 126 个字(words)。至于 49 个测试内核全部匹配 Tesla T4 的成绩,实际上采用了两步验证:先由项目自带的 Python 参考解释器在 49 个受限整数内核上与真实 Tesla T4 比对一致,再通过 --check 重放模式确认模型的前向传播轨迹与参考解释器逐 token 吻合。
| 维度 | qwenGPU | Qwenputer | Percepta |
|---|---|---|---|
| 基础底座 | Qwen3.5-2B (24层) | Qwen3.5-4B | 自研 2D 注意力 Transformer |
| 编译进权重的机器 | 4 通道 SIMT GPU | 32 位 RISC-V CPU | WebAssembly 解释器 |
| 附加编译层数 | 45 层 (总计 69 层) | 11 层 | 未公开 |
| 状态存储位置 | Gated DeltaNet 循环状态 | Gated DeltaNet 循环状态 | 上下文执行轨迹 |
| 实测生成速度 | 10-15 tokens/s (Apple MPS) | 约 3,200 trace tokens/s (L4) | 30,000+ tokens/s (CPU) |
把并行 SIMT 架构塞进权重的代价,直接反映在推理耗时上。相比于只加了 11 层编译层、在 L4 GPU 上配合 Triton 内核能跑出 3,200 trace tokens/s 的 RISC-V 方案 Qwenputer,为了支持线程同步和共享内存而强行前置 45 层网络的 qwenGPU,模型总层数几乎翻了三倍。在需要约 16GB 可用内存且无量化版本的条件下,它在 Apple MPS 上每秒只能吐出 10 到 15 个 token。由于一次简单的 CUDA 内核启动往往需要展开几十到上百个执行轨迹 token,算一道简单的整数题就得花掉几秒到几十秒。
- 风险.目前所有精度数据(含常规模式下 20/24 的数学聊天准确率、高级模式下模型自己写 CUDA 内核 4/5 的首次通过率,以及合并后模型 7e-4 的 KL 散度文本漂移)均为开发者单方自测,且受限于 2B 小模型自身孱弱的代码生成能力,尚无第三方独立复现报告。
对于关注这一前沿方向的开源社区和底层研究者来说,接下来真正决定权重编译计算机能否走出玩具阶段的变量只有三个:有没有独立团队能在更大规模的语言基准上复现其 7e-4 KL 散度的无损声明;编译器能否以合理的层数开销攻克 IEEE 754 浮点数运算;以及能否像 Qwenputer 那样通过底层算子融合,把每秒十几 token 的龟速拉回实用水位。在此之前,它依然是一件精巧但昂贵的实验室展品。
