开发者将SIMT GPU编译进Qwen3.5-2B的45层权重中运行CUDA

qwenGPU 是一个实验性模型，它把一个 SIMT GPU 直接编译进 Qwen3.5-2B 的权重里。开发者 Vixhal（@TheVixhal）以 MIT 许可证在 GitHub 上发布了它。模型在自己的前向传播中运行真正的 CUDA，并编译为 NVIDIA PTX。对开发者而言，它验证了一个想法：在预训练 LLM 内部进行精确计算，而不是在它旁边。

概要

• 规模：Qwen3.5-2B（24 层）加上 45 个编译出的 GPU 层，共 69 层。上下文：262,144 token（基础模型）。
• 运行环境：Apple 芯片、NVIDIA GPU 或 CPU，约 16 GB 可用内存。没有量化版本。
• 性能：开发者运行中每次启动都是精确的；文本漂移的 KL 约为 7e-4。
• 最佳表现：参考解释器在 49 个 kernel 中的 49 个上与真实 NVIDIA Tesla T4 一致。
• 最差表现：在 Apple MPS 上约每秒 10 到 15 个 token；Qwen 编写的 CUDA 在 5 个中通过了 4 个。

qwenGPU 是什么？

qwenGPU 是 Qwen3.5-2B，在 Qwen 自有的 24 层前面放置了 45 个编译出的 transformer 层。这些层都没有经过训练。每个权重都由编译器计算得出。

GPU 如何嵌入模型权重？

这个 GPU 是一个 4 lane 的 SIMT 机器。它有每 lane 15 个寄存器、4 级流水线、32 块 ROM 和 126 个内存字。它只使用 Qwen3.5 自有的层类型：Gated DeltaNet 线性注意力和 SwiGLU MLP。 关键在于 Gated DeltaNet 的循环状态。在电路的头中，门保持打开，写入强度为 1。每个 key 都是 one-hot。每次更新恰好覆盖 1 行，one-hot 查询再把它读回来。因此每一行都充当一个存储单元。 该状态保存程序 ROM、寄存器和全局内存。README 列出每 lane 16 个寄存器行，其中 1 个始终读为零。这与 X 帖子中 15 个可用寄存器的说法一致。 该 GPU 支持整数运算、谓词、发散分支和循环。它还支持共享内存、atomicAdd、__syncthreads() 和 warp shuffle。

CUDA 启动时会发生什么？

主机扮演 NVRTC 和 CUDA 驱动的角色。它用 clang 把 CUDA 编译为 PTX，再把 PTX 降为机器块。它把程序和输入数组作为 token 喂给模型。 从那时起，模型自己写出启动过程。每个 EXEC(b, w) token 在 1 次前向传播中为 warp w 运行块 b。其后的 VAL token 携带寄存器结果。内存写入以 MEM 记录返回。主机从循环状态读取结果，自身不计算任何结果。 在演示视频中，[1, 2, 3] 和 [4, 5, 6] 的点积返回 32。模型用 47 个 token 写出了那次启动。 Meet qwenGPU: Qwen3.5-2B with a SIMT GPU inside its weights, 45 compiled layers in front of Qwen’s own, using its Gated DeltaNet layers as registers and memory. It runs real CUDA in its forward passes. The host compiles the kernel to PTX and loads it, then the model writes every… pic.twitter.com/iM5IETEbIv — vixhaℓ (@TheVixhal) November 1, 2026

GPU 会破坏 Qwen 的文本能力吗？

按开发者的测量，不会。电路在文本 token 上被静音，Qwen 自身的权重没有改变。合并后的模型的下一个 token 分布与原始模型保持在约 7e-4 的 KL 范围内。

qwenGPU 的准确度如何？

所有数字均为自报，在 Apple 芯片上用 MPS 以 bf16 测量：

• 在 --check replay 下，每次启动都是精确的。
• 参考解释器在 49 个测试 kernel 中的 49 个上与 NVIDIA Tesla T4 一致。
• 普通模式用 GPU 计算出的数字回答了 24 个抽样聊天问题中的 20 个。
• 高级模式在 5 个问题中的 4 个上首次尝试就写出了正确的 kernel。 在普通模式下，对于现成的 kernel，Qwen 写出类似 [1, 2, 3] @ [4, 5, 6] 的表达式。在高级模式下，Qwen 自己写出 CUDA kernel。README 警告说，2B 模型的 CUDA 仍然不可靠。

qwenGPU 与 Qwenputer 和 Percepta 相比如何？

qwenGPU 使用了 Divyansh Agrawal 的 Qwenputer 中的 transformer 编译器。Qwenputer 把一个 32 位 RISC-V CPU 编译进 Qwen3.5-4B。Percepta 在 2026 年 3 月把一个 WebAssembly 解释器构建进一个带有 2D 注意力头的 transformer。 特性qwenGPUQwenputer（RISC-V）Percepta 基础模型Qwen3.5-2B，24 层Qwen3.5-4B带有 2D 注意力头的自定义 transformer 编译进的机器4 lane SIMT GPU32 位 RISC-V（RV32IM + Zba + bignum）WebAssembly 解释器 运行的程序CUDA，从 PTX 降级而来C 和 MicroPythonC，编译为 WebAssembly 编译层数4511未披露 状态存放位置Gated DeltaNet 状态Gated DeltaNet 状态上下文中的执行轨迹 主机角色把 CUDA 编译为 PTX，加载 token，读取结果加载固件；不执行计算未披露 正确性检查解释器在 49 个 kernel 中的 49 个上与 Tesla T4 一致约 230,000 个重放位置中 0 个不匹配声称数百万个可靠步骤 报告速度每秒 10 到 15 个 token（Apple MPS）约每秒 3,200 个轨迹 token（L4，Triton）每秒 30,000+ token（CPU） 相对基础模型的文本漂移KL 约 7e-4KL 6e-4不适用 许可证MIT 代码；Apache-2.0 checkpointMIT 代码；Apache-2.0 权重未披露 速度数字使用不同的硬件和 token 类型，因此不能直接比较。

运行它需要什么？

快速开始需要 Python 3.11+、uv、just 和约 16 GB 可用内存。高级模式还需要带有 NVPTX 后端的 clang。just build 会下载 Qwen3.5-2B 并把 GPU 编译进去。没有发布预构建的 checkpoint。构建出的 checkpoint 归入 Qwen 的 Apache-2.0 许可证。

要点

• qwenGPU 把一个 4 lane SIMT GPU 作为 45 个未训练层编译进 Qwen3.5-2B。
• Gated DeltaNet 状态充当寄存器、程序 ROM 和全局内存。
• 模型写出每一次 CUDA 启动步骤；主机只负责编译和加载。
• 解释器在 49 个 kernel 中的 49 个上与 Tesla T4 一致。
• 在每秒 10 到 15 个 token 的速度下，它是一个研究演示，而不是运行时。 查看 GitHub 仓库、qgpu CUDA 参考、Qwen3.5-2B 模型卡和原始 X 帖子。另外，欢迎在 Twitter 上关注我们，别忘了加入我们 150k+ 的 ML SubReddit 并订阅我们的 Newsletter。等等！你在 Telegram 上吗？现在你也可以在 Telegram 上加入我们。

（原文较长，此处译出前 6000 字，完整内容请点「打开原文」。）