跳到正文
RadarAI· MarkTechPost(RSS)·· 6 小时前AI 评分58

开发者发布 qwenGPU:将 45 层编译 GPU 融入 Qwen3.5-2B 权重并运行 CUDA

AI 导读

开发者 Vixhal 以 MIT 许可在 GitHub 发布实验模型 qwenGPU,把 4 通道 SIMT GPU 编译成 45 层未训练权重,置于 Qwen3.5-2B 自身 24 层之前,使模型在前向传播中运行编译为 PTX 的真实 CUDA。

正文 · AI 翻译

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。

特性qwenGPUQwenputer (RISC-V)Percepta
基础模型Qwen3.5-2B,24 层Qwen3.5-4B带有 2D 注意力头的自定义 transformer
编译进的机器4 通道 SIMT GPU32 位 RISC-V(RV32IM + Zba + bignum)WebAssembly 解释器
运行的程序CUDA,从 PTX 降级而来C 和 MicroPythonC,编译为 WebAssembly
编译层数4511未披露
状态存放位置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-4KL 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