一句话结论：qgpu 把一台 GPU 直接编译进了 Qwen3.5-2B 的权重，模型自己写出每一次 CUDA 启动，49 个测试 kernel 与真实的 NVIDIA Tesla T4 逐条对齐。
11 月 1 日，GitHub 用户 vixhal-baraiya 放出开源项目 qgpu，两天内拿到 55 个 star。项目的起点是 Qwenputer，一个把 RISC-V 处理器编译进 Qwen3.5-4B 权重的实验；qgpu 把被编译的对象从 CPU 换成 GPU，算子也从加法乘法扩展到了能写完整的 CUDA kernel。
速度不是这个项目的看点。一个 2B 模型能跑多快这件事早有定论，真正有意思的是它把模型做算术这件事拆到了可以逐条核对的程度。
被塞进权重的这台 GPU 规格很小
qgpu 里的 GPU 不是真实硬件的复刻，规模小到可以完整写进一张表格。它是一个 4 通道 SIMT 机器，每个通道 15 个可用寄存器，4 级流水线，32 个块的程序 ROM，126 个内存字，全部用 Qwen3.5 自带的两种层搭出来，Gated DeltaNet 线性注意力与 SwiGLU MLP。
整机被编译成 45 层 Transformer，排在 Qwen3.5-2B 自带的 24 层之前，合计 69 层。这 45 层的权重一个都没有训练过，全部由编译器算出。
| 项目 | 数值 |
|---|---|
| 通道数 | 4（SIMT 结构） |
| 每通道可用寄存器 | 15（README 列 16 行，其中 1 行恒读零） |
| 流水线级数 | 4 |
| 程序 ROM | 32 个机器块 |
| 全局内存 | 126 个字 |
| 编译出的 Transformer 层 | 45 |
| 基座自有层 | 24 |
| 合并后合计层数 | 69 |
| 上下文长度 | 262,144 词元 |
| 基座隐藏维度与词表 | 2048 / 248,320 |
指令集覆盖整数运算、比较与谓词、发散分支、循环、全局与共享内存、atomicAdd、__syncthreads() 以及 warp shuffle。对照真实硬件，这一节的每个数字都小得不像话，但功能清单是完整的。
寄存器与内存怎么住进 Transformer 层里
难点在于 Transformer 层本身没有寄存器这种东西。qgpu 的做法是借 Gated DeltaNet 的循环状态：门控保持打开、写入强度设为 1、每个键都是独热向量，于是每次状态更新恰好覆盖一行，独热查询再把它读回来。每一行就成了一格存储单元。
程序 ROM 放在各头自己的头维度里，寄存器按 warp 与编号分行，全局内存每行一个 32 位字。README 明确列出每通道每 warp 16 行寄存器，其中 1 行恒读零，和前面说的 15 个可用寄存器对得上。
基座是 Qwen3.5-2B，它的配置里写明每 4 层才有一层全注意力，其余是线性注意力，隐藏维度 2048，词表 248,320。这些数字能对上，说明电路层与 Qwen 的维度确实做了共享，而不是简单拼接一段网络。
启动一次 kernel，模型要吐多少 token
宿主程序扮演 NVRTC 与 CUDA 驱动的角色。它用 clang 把 CUDA 源码编译成 PTX，再降到机器块，把程序镜像和输入数组当作词元喂给模型，之后的启动过程由模型自己写。
token 流有固定格式。PROG 与 LIVE 记录负责装载程序，MEM 记录装载输入，EXEC 令牌从 block 0 的 warp 0 开始。之后的每个 EXEC(b, w) 令牌在前向过程中执行一个块，后续的 VAL 令牌带回寄存器结果，每个被后续块读取的寄存器按每通道两个 16 位令牌返回，内存写入以 MEM 记录返回，EXEC(0, warps) 结束这次启动。
宿主只从循环状态里读结果，不做任何计算。演示里对向量 [1, 2, 3] 与 [4, 5, 6] 做点积，结果是 32，模型写出这一次启动用了 47 个词元。
正确性怎么验，49 个 kernel 与真实 T4 对齐
项目提供了 –check 选项，把每次启动在参考解释器上回放检查，模型的令牌必须逐条匹配。开发者自测称，每次启动都是精确的；参考解释器与真实的 NVIDIA Tesla T4 在 49 个测试 kernel 上全部匹配。
文本能力基本没被破坏。合并后的模型下一词元分布与原模型的 KL 散度约 7e-4，这是一个相当小的量级。普通模式下 24 道抽样问答题答对 20 道，高级模式下 5 道向量与矩阵题里有 4 道一次写对 kernel，涵盖点积、矩阵乘、平方长度与向量加法。
作者也提醒，2B 模型写出的 CUDA 依然靠不住，需要可靠答案时该用普通模式。
速度、门槛与合并方式
一次启动要几十到几百个生成词元，Apple MPS 环境下每秒 10 到 15 个词元，跑起来需要约 16 GB 空闲内存。运行要求 Python 3.11 以上、uv、just，高级模式还需要带 NVPTX 后端的 clang。
qgpu 沿用 Qwenputer 的合并方案：拓宽 Qwen 的残差流，把电路层放在前面，把 GPU 的令牌加进词表。电路中对文本无意义的维度与 Qwen 的维度重叠，因此合并后的检查点窄得多。检查点属于改过的 Qwen3.5，权重仍受 Apache-2.0 约束，代码本身是 MIT。
它和 Qwenputer、和专用推理卡的差别
同一体系下的 Qwenputer 把 32 位 RISC-V 处理器编译进 Qwen3.5-4B，仓库积累 524 个 star。qgpu 换了三件事，目标从 CPU 变成 GPU，基座从 4B 换成 2B，算子从通用计算扩展到能写 CUDA kernel。
与专用推理卡的差别更本质。专用卡靠的是硬件里的张量核心与大规模并行；qgpu 里的 GPU 是 4 通道 SIMT、15 个寄存器、126 个内存字，真实 GPU 上的 kernel 到了这里几乎全部需要重写。这条路线拿到的不是吞吐，是可验证性。
常见问题
qgpu 能替代真实 GPU 吗
不能。目前这 4 通道 SIMT 机器只够跑通演示级别的 kernel，一两个输入向量、几十个词元的一次启动。它验证的是模型能否自己写出并执行 kernel 这件事，本身不是替代方案。
自己跑一遍需要什么硬件
Apple 芯片走 MPS 路径就能跑，约 16 GB 空闲内存；NVIDIA GPU 与 CPU 也在支持范围内。高级模式要写 CUDA kernel，需要额外装好带 NVPTX 后端的 clang。
它会破坏 Qwen 原本的能力吗
文本分布几乎不受影响，KL 散度约 7e-4，24 道抽样题答对 20 道。真正的限制在模型规模上，2B 写出高级 kernel 的成功率是 4 比 5，作者本人也把高级模式定位为演示。
小结
把 qgpu 当成 GPU 的替代方案并不准确，它更像一份可以逐行核对的技术演示。它证明了一件此前少有人证明的事，模型确实能靠自身权重执行自己写出的 GPU kernel，而验证方式是 49 个 kernel 对着真实 Tesla T4 逐条比。
更长的看点在于这条路线与专用加速卡之间的分工。目前的推理加速思路集中在把模型压小或者把请求批量化，qgpu 提供了第三种，把一部分计算挪进权重本身，让模型既是推理引擎也是被调用的程序。
相关阅读
- Ternary Bonsai 2：Qwen 3.5 基座，手机端每秒 120 词元
- FlashMap 投机解码：Qwen3-8B 加速 6.14 倍与 16 token 草稿
- 阿里玄铁 RISC-V 芯片出货超 60 亿颗
- 谷歌开源 TorchTPU：CUDA 护城河首现裂痕
- 天数智芯天垓 300 实机跑通：单卡吞吐达 H200 的 1.05 倍
延伸阅读
- 英伟达 RTX 5090 美国全面下架：消费旗舰停产背后的 AI 算力争夺战3 周前
- Qwen-Image-2.1 开源：7B 参数统一文生图与编辑，原生透明图 + 10 张参考图1 月前
- OpenMobile-2 开源手机智能体：110 个应用，79.9 分1 周前
- Qwen3.8-LiveTranslate：Interleave架构同传延迟2.3秒3 周前
继续读
同一批更新里这几篇值得一看：
- SOL-ExecBench 235 题拆解：无问芯穹 71 个单题第一2026-10-30
- 微软单季从 OpenAI 收 100 亿美元：24% 持股与 400 亿年化2026-10-29
- 天数智芯天垓300实机跑通：单卡吞吐达H200的1.05倍2026-10-29
想快速查定义？对应的词条在这里：
- Qwen词条 · AI 自学笔记