当前位置:首页>python>5000行纯Python,手搓Qwen2.5-7B的CUDA编译器

5000行纯Python,手搓Qwen2.5-7B的CUDA编译器

  • 2026-10-11 08:26:22
5000行纯Python,手搓Qwen2.5-7B的CUDA编译器

TVM是50万行C++,PyTorch堆着Dynamo、Inductor和Triton。现在有人用约5000行纯Python写了编译器,直接输出原生CUDA。

  • 纯Python写编译器:约5000行代码,告别C++巨兽,易读易改
  • 支持主流小模型:能把TinyLlama和Qwen2.5-7B直接转成CUDA内核
  • 极致算力优化:6层IR降级,自动融合算子、消除中间显存占用

为什么我们要关注底层编译器

接活做AIGC,算力成本是大头。想省显存提速,就得看底层。但现在的编译器栈太折磨了。PyTorch那一套Dynamo、Inductor、Triton层层套娃,还有XLA、MLIR、Halide、Mojo,根本没教程讲清楚高层设计,直接把你扔进源码里。出了问题没法调,只能干瞪眼加钱买卡。

这哥们从零写了约5000行纯Python的参考编译器。它把模型通过6个IR(中间表示)一步步降级成原生CUDA。目标不是干翻Triton,而是做一个容易跟进、随便改的编译器。对我们来说,这意味着能真正看懂并控制模型是怎么跑在显卡上的,从哪省显存,从哪挤速度。

6层IR降级到底干了什么

拿torch.relu(torch.matmul(x + bias, w))举例,看它怎么把代码一步步压进显卡。

Torch IR:抓取FX图,和PyTorch算子一一对应,原汁原味。

Tensor IR:把算子拆成Elementwise、Reduction、IndexMap。统一算子面,未来ONNX、JAX前端接入不用动下游。那个(16, 64, 16)的中间变量看着吓人,但下一阶段直接融合掉,永远不会真正生成。

Loop IR:算子融合发生地。广播乘法、归约、输出布局、relu收尾,全压进一个循环嵌套,中间缓冲区全部干掉。

Tile IR:开始感知GPU。循环轴分配给线程/块,Stage把共享输入提入共享内存,2x2寄存器分块让一个线程算4个输出。K轴切成2次32宽的外部迭代。三大优化:buffers=2@a2双缓冲让加载和计算重叠;async用cp.async.ca.shared.global不阻塞warp;pad=(0,1,0)加1个元素填充消除bank冲突。

Kernel IR:变成硬件原语。THREAD/BLOCK变threadIdx/blockIdx,异步Stage变Smem加cp.async填充和commit/wait围栏,同步Stage变跨步填充循环。框架无关,也能降级到Metal或HIP。

CUDA:一对一转换,读x、bias、w各一次,写relu一次,全在一个内核完成,零中间变量。

自己动手改编译器

项目名叫deplodock,代码量极小,非常适合想搞懂编译器原理的创作者自己拉下来改。想给TinyLlama或Qwen2.5-7B写专属优化,或者研究怎么把显存占用压到极致,这是最好的切入点。

详细文章:

https://medium.com/data-science-collective/a-principled-ml-compiler-stack-in-5-000-lines-of-python-17f2db9549d4

代码仓库:

https://github.com/cloudrift-ai/deplodock

留言聊聊
你平时本地跑图,最头疼哪个环节的显存爆炸?

来源:Reddit MachineLearning|原文:A Hackable ML Compiler Stack in 5,000 Lines of Python

最新文章

随机文章