开发者发布 qwenGPU:将 45 层编译 GPU 融入 Qwen3.5-2B 权重并运行 CUDA
开发者 Vixhal 以 MIT 许可在 GitHub 发布实验模型 qwenGPU,把 4 通道 SIMT GPU 编译成 45 层未训练权重,置于 Qwen3.5-2B 自身 24 层之前,使模型在前向传播中运行编译为 PTX 的真实 CUDA。
MarkTechPost· Michal Sutter·2026-11-02 15:28· 4 小时前AI 评分51
认识 qwenGPU:Qwen3.5-2B 搭载 45 个编译 GPU 层,可运行 CUDA
AI 导读
开发者 Vixhal 在 GitHub 以 MIT 许可发布实验模型 qwenGPU,把一个 4 通道 SIMT GPU 编译为 45 层未训练权重,放在 Qwen3.5-2B 自身 24 层之前,使模型在前向传播中运行真实的 CUDA(编译为 PTX)。
正文 · 原文
qwenGPU 是一个实验模型,它将一个 SIMT GPU 直接编译进 Qwen3.5-2B 的权重中。开发者 Vixhal(@TheVixhal)以 MIT 许可在 GitHub 上发布了它。该模型在自身的前向传播中运行真实的 CUDA(编译为 NVIDIA PTX)。对开发者而言,它检验了一个想法:在预训练 LLM 内部进行精确计算,而非在其旁边。
TL;DR
- 规模:Qwen3.5-2B(24 层)加上 45 个编译 GPU 层,共 69 层。上下文:262,144 个 token(基础模型)。
- 运行于:Apple 芯片、NVIDIA GPU 或 CPU,需约 16 GB 可用内存。无量化版本。
- 性能:开发者运行中的每次启动都是精确的;文本漂移的 KL 约为 7e-4。
- 最佳:参考解释器在 49 个内核中的 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 通道 SIMT 机器。它每通道有 15 个寄存器、一个 4 级流水线、一个 32 块 ROM 和 126 个内存字。它仅使用 Qwen3.5 自身的层类型:Gated DeltaNet 线性注意力和 SwiGLU MLP。
诀窍在于 Gated DeltaNet 的循环状态。在电路的头中,门保持打开,写入强度为 1。每个键都是 one-hot。每次更新恰好覆盖 1 行,one-hot 查询将其读回。因此每一行都充当一个存储单元。
该状态保存程序 ROM、寄存器和全局内存。README 列出每通道 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 写入了该启动。
认识 qwenGPU:Qwen3.5-2B 权重内嵌 SIMT GPU,在 Qwen 自身层之前有 45 个编译层,使用其 Gated DeltaNet 层作为寄存器和内存。
它在前向传播中运行真实 CUDA。主机将内核编译为 PTX 并加载,然后模型写入每个…… pic.twitter.com/iM5IETEbIv
— vixhaℓ (@TheVixhal) November 1, 2026
GPU 会破坏 Qwen 的文本能力吗?
不是按开发者的标准。电路在文本 token 上被静音,Qwen 自身的权重没有改变。合并模型的下一 token 分布与原始模型的 KL 散度保持在约 7e-4 以内。
qwenGPU 的准确度如何?
所有数字均为自报,在 Apple 芯片上使用 MPS 以 bf16 测量:
- 在
--check重放中,每次启动都精确无误。 - 参考解释器在 49 个测试内核中的 49 个上与 NVIDIA Tesla T4 匹配。
- 普通模式用 GPU 计算出的数字回答了 24 个抽样聊天问题中的 20 个。
- 高级模式在 5 个问题中的 4 个上首次尝试就写出了正确的内核。
在普通模式下,Qwen 为现成内核写出类似 [1, 2, 3] @ [4, 5, 6] 的表达式。在高级模式下,Qwen 自己编写 CUDA 内核。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。
| 特性 | qwenGPU | Qwenputer (RISC-V) | Percepta |
|---|---|---|---|
| 基础模型 | Qwen3.5-2B,24 层 | Qwen3.5-4B | 带有 2D 注意力头的自定义 transformer |
| 编译进的机器 | 4 通道 SIMT GPU | 32 位 RISC-V(RV32IM + Zba + bignum) | WebAssembly 解释器 |
| 运行的程序 | CUDA,从 PTX 降级而来 | C 和 MicroPython | C,编译为 WebAssembly |
| 编译层数 | 45 | 11 | 未披露 |
| 状态存放位置 | Gated DeltaNet 状态 | Gated DeltaNet 状态 | 上下文中的执行轨迹 |
| 主机角色 | 将 CUDA 编译为 PTX,加载 token,读取结果 | 加载固件;不运行计算 | 未披露 |
| 正确性检查 | 解释器在 49 个内核中的 49 个上与 Tesla T4 匹配 | 在约 230,000 个重放位置中 0 处不匹配 | 声称数百万个可靠步骤 |
| 报告速度 | 10 到 15 tokens/s(Apple MPS) | 约 3,200 trace tokens/s(L4,Triton) | 30,000+ tokens/s(CPU) |
| 相对基础模型的文本漂移 | KL 约 7e-4 | KL 6e-4 | 不适用 |
| 许可证 | MIT 代码;Apache-2.0 检查点 | MIT 代码;Apache-2.0 权重 | 未披露 |
速度数据使用不同的硬件和 token 类型,因此不能直接比较。
运行它需要什么?
快速入门需要 Python 3.11+、uv、just 和约 16 GB 可用内存。高级模式还需要带有 NVPTX 后端的 clang。just build 下载 Qwen3.5-2B 并将 GPU 编译进去。未发布预构建检查点。构建的检查点受 Qwen 的 Apache-2.0 许可证约束。
关键要点
- qwenGPU 将 4 通道 SIMT GPU 作为 45 个未训练层编译进 Qwen3.5-2B。
- Gated DeltaNet 状态充当寄存器、程序 ROM 和全局内存。
- 模型编写每个 CUDA 启动步骤;主机只负责编译和加载。
- 解释器在 49 个内核中的 49 个上与 Tesla T4 匹配。
- 以每秒 10 到 15 个 token 的速度,它是一个研究演示,而不是运行时。
查看 GitHub 仓库、qgpu CUDA 参考、Qwen3.5-2B 模型卡和原始 X 帖子。另外,欢迎在 Twitter 上关注我们,别忘了加入我们的 150k+ ML SubReddit 并订阅我们的新闻通讯。等等!你在 Telegram 上吗?现在你也可以在 Telegram 上加入我们。
需要与我们合作推广您的 GitHub 仓库、Hugging Face 页面、产品发布或网络研讨会等? 联系我们
文章 Meet qwenGPU: Qwen3.5-2B With 45 Compiled GPU Layers That Run CUDA 首次发布于 MarkTechPost。
来源:MarkTechPost · marktechpost.com
来源:RadarAI · aihot.news