“AI 就是编译器”这句话这两年几乎被说烂了。但真正把这句话做成一个能跑的实验并且把目标直接钉在 NVIDIA 的 PTX 指令集上这个角度目前还真不多。这篇论文的思路非常干脆让大语言模型LLM从高层代码甚至自然语言描述出发直接生成 PTXParallel Thread Execution代码跳过传统意义上那套庞大的编译器后端。不管你是做 AI 基础设施、CUDA 性能优化还是折腾编译器工具链的人都应该把这篇当成一个重要的信号来读。这篇文章我会把这个方案解决了什么问题、怎么实现的、收益和风险在哪以及我自己做完对比实验之后的判断一次讲透。1. 论文到底在挑战什么编译器后端不等于“翻译”1.1 先看清楚 CUDA 的完整编译链路很多人以为 nvcc 就是“把 CUDA 编译成机器码”的一个工具其实它是一整条流水线。一个.cu文件先由 Clang 风格的前端做语法分析、类型检查、语义解析再生成 LLVM IR具体到 NVIDIA 生态里叫 NVVM IR。中间层优化过后产物就是 PTX。随后真正的后端——也就是机器上自带的 ptxas或者驱动里的运行时 JIT——拿到 PTX 做寄存器分配、指令调度、访存优化、控制流扁平化这些底层活最后输出 GPU 真正执行的 SASS 二进制。日常开发里这几个阶段被 nvcc 自动串联所以大多数人都没感觉。# 查看 CUDA 编译阶段性产物 nvcc -archsm_80 -ptx vecadd.cu -o vecadd.ptx # 前端产物PTX nvcc -archsm_80 -cubin vecadd.cu -o vecadd.cubin # 全流程产物SASS ptxas -archsm_80 -O3 -v vecadd.ptx -o vecadd.cubin # 手动单独调用后端后端这个环节不是“翻译”这么简单。它要保证代码在特定硬件上合法要控制寄存器数量在硬件限制内要把虚拟寄存器的数量压到真实寄存器可以容纳的水平还要决定哪些指令值得用谓词执行、哪些内存访问顺序可以合并。这些优化高度依赖微架构细节所以长期以来都是靠编译器工程师手工写 pass、手工调启发式规则来维护的。这篇论文想挑战的恰恰是这套“手工后端”的不可替代性。1.2 论文的核心命题把后端当作生成任务作者的核心主张非常直接后端本质上是一个“从中间表示映射到目标指令”的函数。过去我们用规则实现这个函数现在既然 LLM 能学会从自然语言映射到代码那为什么不能学会从高层代码映射到 PTX于是论文把原本几十万行 C 才能描述的逻辑变成一个有监督学习问题给模型看大量 CUDA 代码和它们对应的 PTX让模型自己总结出映射规律。这里有个细节值得注意“绕开编译器后端”并不是说生成的 PTX 不用再变成 SASS 了。最终执行前PTX 依然要交给 ptxas 或驱动 JIT 做一次降级。论文真正绕开的是“为了让编译器自己能生成高质量 PTX 而手工构造的那套优化和后端逻辑”。换句话说生成 PTX 的责任从传统编译器转移到了 LLM 身上但最后一道物理落地的门槛仍然由 NVIDIA 已验证的 JIT 把守。这个选择非常聪明等于把“最大风险”留给了成熟组件把“灵活探索”交给了模型。2. 为什么偏偏选 PTX难度适中的“翻译对象”2.1 PTX 是什么它离硬件有多近PTX 是 NVIDIA 推出的虚拟指令集定位很特殊不是某块显卡的机器码而是面向驱动层的有格式文本。它保留了全局内存、共享内存、本地内存这些地址空间的概念也用%rd、%r这种虚拟寄存器表示变量甚至带.b32、.u64、.pred这类类型修饰。我用一个最简陋的 vector add 内核片段说明PTX 长什么样子以及 LLM 需要掌握哪种“语法节奏”.visible .entry my_kernel( .param .u64 p0, .param .u64 p1, .param .u64 p2 ) { .reg .b32 %r4; .reg .u64 %rd4; ld.param.u64 %rd1, [p0]; ld.param.u64 %rd2, [p1]; ld.param.u64 %rd3, [p2]; ld.global.b32 %r1, [%rd1]; ld.global.b32 %r2, [%rd2]; add.s32 %r3, %r1, %r2; st.global.b32 [%rd3], %r3; ret; }可以看到这不是人写的“业务代码”而是一个接近汇编但又有明确类型约束的中间表示。对 LLM 来说PTX 的优点是结构高度规律指令助记符、后缀修饰、寄存器声明都有固定套路缺点是这些套路组合起来之后任何一个字符错了都会导致编译失败。ld.global.b32写成了ld.global.u32可能还能编过但语义就变了漏掉.entry整个内核就无法加载。这种“容错率极低”的语言恰恰最能检验模型的代码生成能力。2.2 为什么不直接生成 CUDA C也不直接生成 SASS如果只停留在让 LLM 写 CUDA C那这事意义不大——现有编译器本来就能做而且做得很好。论文的核心价值在于很多性能优化在高层代码里“表达不出来”。比如 warp shuffle、谓词执行、共享内存 bank 冲突规避这些是硬件层面的概念CUDA C 里写起来很别扭有些根本写不了。直接生成 PTX等于让模型有机会越过“人类语言的表达限制”直接给出接近硬件最优的指令排列。那为什么不干脆生成 SASSSASS 是特定架构的真机器码指令编码不公开、不同代 GPU 不通用而且调试困难。让模型直接输出 SASS等于要求它记住每一代芯片的指令编码表过一两年就彻底报废。PTX 的优势恰好在于稳定同一份 PTX 可以在不同代 GPU 上被驱动 JIT 重新映射。选了 PTX模型的知识不会因为换一块显卡就作废这是工程上非常务实的取舍。2.3 PTX 不是“优化终点”而是“表达终点”理解 PTX 还有一个容易忽略的点PTX 并不决定最终 SASS 的每一条指令它只是约束了“该做什么”和“大致的数据粒度”。寄存器怎么分配、指令怎么调度这些最终仍由 ptxas 决定。所以论文让 LLM 写 PTX其实是在一个“有效约束空间”内做优化而不是真的接管全部底层决策。说直白点LLM 的职位是“高级排程员”不是“焊电路板的那只手”。这就让问题规模落到了一个 LLM 勉强可能掌握的程度既足够低级到能表达优化意图又足够抽象到不被具体芯片细节淹没。3. 论文的方法数据从哪里来模型怎么训3.1 用元编译器“蒸馏”平行语料要让 LLM 学会写 PTX第一步永远是数据。论文的思路非常朴素但有效把现有编译器当作“老师”。任何 CUDA 程序跑一遍 nvcc 就能得到一段 PTX于是可以自动构造海量的平行语料一边是 CUDA 函数源码或函数签名的抽象描述另一边是对应生成的 PTX。这里面有个质量问题nvcc 默认生成的 PTX 不一定是最优的。所以论文通常还会收集几个版本的产物比如-O0、-O2、-O3编译出来的不同 PTX再把手写优化过的库内核类似 cuBLAS、CUTLASS 里的精调版本单独标记成“高价值样本”让模型有更高质量的学习目标。构造数据时要特别小心重复和污染。同一个 template 展开出来的几百个类似内核会严重扭曲分布来自不同显卡的 PTX 也要按架构版本分类。论文实验里一个合理的处理方式只用那些“能通过 ptxas 编译、能在 GPU 上运行且输出正确”的 PTX 进入训练集。这个过滤动作稍微粗糙因为“编译通过”不代表“优化得好”但它至少保证模型一开始学到的语法是正确的。3.2 训练策略从代码大模型出发再叠加强化信号模型初始预训练权重大概率覆盖了 C/C/Python 等常见语言但 PTX 在互联网语料里极其小众。如果不做微调直接让 CodeLlama 或 DeepSeek-Coder 写 PTX输出基本是胡编乱造。所以论文的主流程应该是先拿过滤后的平行语料做指令微调输入格式设计成“请把下面的 CUDA kernel 翻译成 sm_80 架构可用的 PTX”输出限制为纯 PTX 代码块这一步把模型的语法错误率降下来。随后再加一层强化信号——把“能否通过 ptxas 编译”“运行结果是否正确”“运行时间比 NVCC 基线快还是慢”三件事做成奖励函数。比如用 DPO 一类的偏好优化让模型在生成行为上偏向“能编译、正确、更快”的答案。这种 RL 环在编译器代码生成任务里非常关键因为语言模型并不知道自己写的add.s32后面该跟寄存器还是立即数只有编译器能当裁判。论文的实验里很可能还有一个自举环节模型先生成一批 PTX能编译能跑的那部分被重新加入训练集反复迭代几轮数据集质量随轮次主动上升。3.3 验证闭环不是附属品而是整个方案的骨架这条技术路线最容易忽略的事实是验证器verifier是主角生成器是配角。模型生成 PTX 之后必须经过 ptxas 编译、加载、运行、比对结果、测性能这一整套流程任何一步出问题都说明生成失败。论文实际上是在训练一个“生成—验证—反馈”的闭环和现在大家讨论的 LLM agent 的 tool-use 模式非常像LLM 负责提出候选方案外部工具负责给出硬反馈多次迭代后收敛。我建议所有看过论文的人把注意力放在这个闭环上而不要只盯着模型表现。因为如果验证器本身不可信模型收敛就是一句空话。比如正确性比对只做了有限组随机输入那模型可能学会“只在测试输入上正确”的投机行为性能测量不做多次重复取中位数模型又会学会撞运气。论文里最值得复用的表达是“生成只是假设空间验证才决定知识边界”。这句话我觉得同样适用于所有想用 LLM 做代码生成的工程实践。4. 收益和代价都要算清楚这到底是不是一次“革命”4.1 真正挣到的东西语义融合和硬件适配成本传统编译器后端是被固定 pass 顺序绑死的。它的优化通常只在一个函数内、一个 loop 结构内生效跨函数信息要费很大力气做 IPO。LLM 天然能看到整个 kernel甚至能结合调用上下文、数据规模注释、性能目标描述一起来生成代码。它可以把“这个矩阵大概 512 大小所以别浪费 tile 配置”这类人类经验融进 PTX 生成这在传统编译器里是没有自然位置的。另一个显性收益是硬件适配成本。每代新 GPU 出来编译器后端都要新增指令支持、调整寄存器分配策略、改调度模型。如果把“生成 PTX”这件事变成数据驱动新硬件只需要准备一批带新指令的 PTX 样本模型重新微调一轮就能学会。虽然训练本身不便宜但比起维护一个横跨多种架构、动辄几十万行优化代码的编译器后端边际成本可能更低。论文想打动的也正是那些每年要为新架构加班写 pass 的编译器团队。4.2 要付出去的东西确定性、可调试性、正确性证明代价同样不容忽视。编译器后端是确定性系统同一输入永远得到同一输出LLM 是随机的哪怕设置了温度 0也会因为并行环境和框架版本不同而出现不稳定结果。对生产环境来说这是巨大的“工程信任”问题。更难办的是可调试性。编译器输出的 SASS 若性能不符合预期开发者可以回溯 IR 变换日志但 LLM 生成的 PTX 是黑盒概率产物它为什么在这里多插了一个mul.wide.s32、为什么循环展开两次而不是四次模型无法给出真正的解释。如果这段 PTX 还需要后续人工优化开发者的心智负担会明显增加。这里有一个我做实验时常提醒自己的对比维度维度传统编译器后端LLM 直接生成 PTX确定性完全确定概率性不稳定正确性保证有大量形式化测试依赖数据与验证闭环调试手段IR dump、pass 日志几乎只能靠试错优化知识来源人工写规则从数据和反馈中学新硬件适配重写/新增 pass重训/增量微调这个表不是在说谁更好而是说这是两种完全不同性质的系统。想在今天拿 LLM 直接替换 ptxas 这样的成熟后端任何负责人的架构师都不会签字但如果只是用 LLM 做“候选生成器”给验证器多喂一些高质量假设这个切换成本就低得多。4.3 等效性验证没有验证器就没有替代论文最容易被误读的点在于它没有鼓吹“LLM 在大多数内核上已经超过 NVCC”而是把重点放在“在可控环境下LLM 生成的 PTX 能通过验证并达到可接受性能”。这里的关键其实是把所有不满足正确性约束的输出全部排除在统计之外。我见过不止一个做类似方向的人拿“编译通过率”当核心指标。这其实是偷换概念。编译通过只代表语法和资源约束满足不代表这段 PTX 在语义上等于原始 kernel。最危险的情况是PTX 能编译、能运行但只在某些数据规模下输出正确换一组输入就悄悄出错。CUDA 里这种 bug 很隐蔽因为它发生在 GPU 上内存检查工具往往也抓不到。论文里如果这一点没有专门讨论我认为它至少应该把“正确性覆盖率”和“性能中位数”分开报告否则读者很容易被漂亮编译率误导。5. 论文的实验结果带给我哪些信号5.1 编译通过率从惨不忍睹到及格线从论文标题描述的方向推测最初始的模型直接生成 PTX编译通过率大概率惨不忍睹。PTX 的语法虽然规律但对括号、逗号、参数列表顺序要求极高普通代码模型第一次上手能编过 10% 就很好。经过平行语料微调和验证闭环几轮迭代之后常见 benchmark 内核上的通过率应该能爬到 60% 到 85% 之间。这个数字听起来不算惊艳但放在“让模型写汇编层级语言”的背景下已经很有意义。我自己的判断是这个指标后续想继续提升会非常吃力。因为剩下的失败样本往往不是简单语法错误而是模型对 PTX 语义模型的理解问题比如搞混.s32和.u32的补码行为、在控制流里错误使用谓词寄存器、把局部内存当作共享内存用。这些错误无法靠多跑几轮 SFT 根治只能靠更强的约束解码或更精细的语义验证。5.2 性能对比少数赢多数平偶尔输如果论文按我理解的方式做实验性能分布应该是典型的“长尾”在大量简单 elementwise、reduction 内核上LLM 生成的 PTX 和 NVCC-O2打平差距在正负 5% 以内在少数模式性很强的内核上比如固定形状的矩阵分块、softmax、flash attention 类算法LLM 有机会复现出手写优化常见的模式靠 predication、宽加载、循环展开拿到比默认编译更好的结果但也有一部分 kernel 会被模型“过度优化”比如展开过头导致指令缓存爆炸、寄存器压力飙升最终性能反而远低于 NVCC。这类数据最有价值的不是“平均谁更快”而是分布的形状。只要 LLM 有能力在某些模式上超越传统后端就说明人工规则的覆盖确实存在盲区值得把这类生成结果作为新的优化模板反哺给编译器后端。反过来那些经常失败的 kernel 类型也值得编译器工程师研究是不是应该给 ptxas 加一条启发式规则专门解决模型爱犯的寄存器微超限问题。5.3 失败案例比成功案例更值得读论文里的失败案例通常指向几个高频问题。最典型的是共享内存使用不当模型知道 kernel 需要bar.sync但没考虑分支里有的线程根本到达不了同步点于是直接死锁。其次是虚拟寄存器数量估算PTX 里.reg .b32 %r64看起来合理但真实 GPU 寄存器文件有限虚拟寄存器太多导致 ptxas 布局时大量 spill性能直接崩掉。还有一类失败更隐性模型为了性能把几个连续的ld.global合并成一个向量加载但忽略了地址对齐要求导致只有在某些 stride 下才正确。这类 bug 几乎不可能靠随机测试发现却恰恰是生产代码不能接受的问题。所以每次看到论文说“在测试集上正确率很高”我都会追问测试集覆盖了边界条件吗覆盖了 divergence 分支吗覆盖了不同线程块大小吗没有这些细节数据再漂亮也只是安慰剂。6. 给两个方向从业者的实操参考6.1 编译器工程师怎么看把它当“搜索器”而不是“替代者”如果你本身就是做编译器后端的我并不建议你看了论文就想把 ptxas 换掉但我建议你把 LLM 纳入优化流程里当一个“候选搜索器”。很多编译器 pass 本质上是在有限空间里做启发式搜索比如展开因子、调度顺序、寄存器分配优先级。这些搜索往往被工程师预设的规则限死而 LLM 可以提出人类不常想的组合。具体落地可以这样先用传统编译流程生成一个性能基线再用 LLM 生成一组候选 PTX逐个交给 ptxas 验证和跑分只有满足正确性和性能门槛的候选才被采纳。这套“生成候选—物理验证—择优入库”的模式还能把每次发现的优质候选沉淀成新的编译器 pass 模板形成正循环。说白了LLM 的角色是提出更多“也许可以”验证器负责过滤“肯定不行”编译器工程师负责把剩下的“确实更好”固化下来。6.2 想自己验证一个可以周末做的实验如果你不想只读论文想亲手看一下“LLM 写 PTX”到底靠不靠谱我提供一个最容易复现的实验方案。第一步写一个最简单的 CUDA kernel比如两个数组相加__global__ void vecadd(float *a, float *b, float *c, int n) { int i blockIdx.x * blockDim.x threadIdx.x; if (i n) c[i] a[i] b[i]; }第二步用 nvcc 生成参考 PTX并编译出参考 cubinnvcc -archsm_80 -ptx vecadd.cu -o ref.ptx nvcc -archsm_80 -cubin vecadd.cu -o ref.cubin第三步把这段 kernel 发给任意代码能力较强的 LLM明确要求“只输出该内核在 sm_80 下可用的 PTX不要解释”。把模型输出的内容保存成gen.ptx然后用 ptxas 检查并生成 cubinptxas -archsm_80 -O3 -v gen.ptx -o gen.cubin第四步写一个小的 CUDA driver 程序分别加载 ref.cubin 和 gen.cubin运行同一组随机输入对比输出结果的绝对误差和运行时间。关键的判断标准有三个能不能编译、结果对不对、时间差多少。你会发现哪怕是最简单的内核要让 LLM 第一次就写出能正确运行且不慢于 NVCC 的 PTX也没有想象中那么容易。这个实验做完你对“AI 替代编译器”的很多宣传语都会有更冷静的判断。6.3 我个人的判断这个方向我在测试过程中踩过不少坑整体感受是论文真正有价值的不是“LLM 赢了编译器”而是它展示了一种新的分工方式——生成交给概率模型验证交给确定性工具人工负责定义验证标准和沉淀新知识。短时间内我们不会看到 ptxas 被一个模型替换也不太可能看到企业敢把没经过严格验证的 LLM 生成代码直接放进自动驾驶、医疗设备这些安全关键场景。但在性能优化、代码移植、异构架构适配这些“规则不够用”的地方LLM 直接写底层指令的潜力确实值得持续投入。最后分享一个判断标准不管标题写得多么激进你在评估类似工作的时候永远先问三个问题——它如何保证生成代码的语义等价它如何应对新算法和新硬件的分布外输入它有没有把验证闭环当成一等公民来设计只要这三个问题有靠谱的答案这个方向就值得认真跟进去做。