最近读到一篇很有冲击力的论文题目很直白AI 就是编译器——让 LLM 直接写 PTX绕开整个编译器后端。核心思想不复杂既然大语言模型已经能写 C、写 Python那为什么不干脆让它直接写 GPU 的底层指令 PTX这样一来传统编译器中那些玄之又玄的优化 pass、IR 变换、指令选择全都可以被一个训练好的神经网络替代。这篇论文的价值不在于马上替代 nvcc 或者 LLVM而在于它把一个原本被认为是“不可学习”的底层问题真正交给了 LLM并且给出了一套可验证的工程路径。我把论文反复读了几遍也顺着它的思路做了些实验今天把里面的关键设计、实操环节、以及真正动手时才会踩到的坑一次性讲清楚。1. 内容整体设计与思路拆解1.1 这个思路到底在解决什么问题传统 GPU 程序从高级语言到硬件要经历一条很长的流水线。以 CUDA 为例大致是CUDA C/C - NVVM IR - LLVM 后端NVPTX- PTX - ptxas - SASS - 硬件执行其中“编译器后端”通常指从 LLVM IR 生成 PTX 这一大段。它包含指令选择、寄存器分配、指令调度、内存访问优化、循环优化等一大堆逻辑。这些逻辑被固化成一个个编译 pass每一行都是编译器工程师用 C 手写的启发式规则。问题在于pass 之间相互作用复杂度极高很难调优每出来一代新硬件就需要写新的 pass 去匹配新的指令特性很多优化是局部启发式的做不到跨全局的协同。论文提出的思路是把“从高层描述到 PTX”的映射直接交给 LLM。给模型看一段 CUDA 代码甚至是一句功能描述让模型输出对应的 PTX。如果能做到等于把整个 NVPTX 后端的学习任务变成了一个语言建模任务。传统编译器后端里的优化逻辑不再是手写规则的集合而变成了模型参数的一部分。这个转变看着简单实际意义很大。传统编译器的优化规则是“人写的经验”它的上限就是编译器工程师对硬件的理解。而 LLM 如果从足够多的源码、PTX、SASS 文本里学习理论上可以学到比单个工程师更广、更细的模式。这也是论文题目里“AI 就是编译器”的底气所在。1.2 为什么选择 PTX而不是 SASS 或 LLVM IR选 PTX 是个很聪明的折中。最开始可能有人会想既然要绕开编译器为什么不直接让 LLM 生成 SASSSASS 是真正的机器指令但它是二进制格式反汇编出来的文本可读性差并且强绑定 GPU 架构A100 的 SASS 不能直接跑到 H100 上。让 LLM 直接生成 SASS对模型来说难度极高对应用来说也失去了跨架构的可移植性。反过来看 LLVM IR它的表达力更强但指令和类型系统太复杂。LLVM IR 有几千种 opcode还有大量 metadata、属性、模块结构让一个文本生成模型去学习这种庞然大物成本比 PTX 高很多。更重要的是LLVM IR 和 PTX 之间还需要一个 NVPTX 后端来转换如果 LLM 只生成 LLVM IR那“绕开编译器后端”的初衷就没了。PTX 处于中间位置它是文本格式指令集规模适中文档公开而且有大量现成语料。PTX 虽然是虚拟指令集但它已经很接近硬件模型——有明确的寄存器、地址空间、同步语义。模型学会写 PTX等于学会了一个精简的 GPU 汇编语言。PTX 还能通过 ptxas 再编译成 SASS支持不同架构实用性好很多。这里要澄清一个容易被标题误导的点论文说绕开整个编译器后端其实并没有连 ptxas 也绕开。PTX 到 SASS 这一步仍然存在因为硬件相关的寄存器分配和调度仍然需要底层工具完成。更准确地说绕开的是从高级语言/NVVM IR 到 PTX 的 LLVM 后端。PTX 生成之后还是要交给 ptxas 做最后的机器码生成。搞混这一点后面很多技术决策就没法理解了。1.3 这么选的好处与代价好处是明显的省掉了一整个庞大的编译器后端代码库不需要维护几千个手写优化 passLLM 可以吸收海量人类写的汇编片段有潜力生成比固定 pass 更聪明的指令序列可以通过自然语言指导编译方向比如“尽量用向量化加载”“减少 bank conflict”这种交互式编译是传统编译器做不到的。代价同样明显正确性没有保证。编译器经过几十年验证输出基本可预期LLM 是概率模型可能偶尔写出完全跑不动的 PTX性能没有保证。模型能生成语法正确的 PTX但性能可能远不如 nvcc -O2可维护性差。项目里全是模型生成的底层汇编后人基本没法改。这不是一篇站队“LLM 取代编译器”的论文更像是在说“你们编译器工程师天天吹自己的优化多牛我偏要让神经网络也试试。”对我这种做 AI Infra 的人来说这种挑战本身就很有价值。2. 核心细节解析与实操要点2.1 PTX 到底长什么样在深入论文之前先把 PTX 的底子补上。PTX 是 NVIDIA 提供的一种虚拟指令集架构它的设计目的是让上层编译器不直接面对每一代 GPU 的硬件差异。一个 PTX kernel 有明显的入口声明、寄存器声明、指令序列、同步原语等。一个很简单的向量加法 kernel用 CUDA 写是这样__global__ void vec_add(const float* a, const 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 做简化大概长这样.visible .entry vec_add( .param .u64 vec_add_param_0, .param .u64 vec_add_param_1, .param .u64 vec_add_param_2, .param .u32 vec_add_param_3 ) { .reg .b32 %r6; .reg .f32 %f3; .reg .u64 %rd10; .reg .pred %p1; ld.param.u64 %rd1, [vec_add_param_0]; ld.param.u64 %rd2, [vec_add_param_1]; ld.param.u64 %rd3, [vec_add_param_2]; ld.param.u32 %r1, [vec_add_param_3]; cvta.to.global.u64 %rd4, %rd1; cvta.to.global.u64 %rd5, %rd2; cvta.to.global.u64 %rd6, %rd3; mov.u32 %r2, %ctaid.x; mov.u32 %r3, %ntid.x; mov.u32 %r4, %tid.x; mad.lo.s32 %r5, %r2, %r3, %r4; setp.ge.s32 %p1, %r5, %r1; %p1 bra L_bb0; mul.wide.s32 %rd7, %r5, 4; add.s64 %rd8, %rd4, %rd7; ld.global.f32 %f1, [%rd8]; add.s64 %rd9, %rd5, %rd7; ld.global.f32 %f2, [%rd9]; add.f32 %f3, %f1, %f2; add.s64 %rd10, %rd6, %rd7; st.global.f32 [%rd10], %f3; L_bb0: ret; }这段代码为了展示核心语句省略了一些无关紧要的转换细节但结构是真实的。PTX 里的很多东西是编译器已经替你决定好的用哪个寄存器、怎么算地址、用标量加载还是向量加载。如果 LLM 想写高性能 PTX它就必须在这些细节上做出选择。对于第一次接触 PTX 的读者只需要抓住几点.reg .b32/.f32/.u64声明虚拟寄存器ld.global.f32表示从 global memory 读出 32 位浮点setp设置谓词寄存器配合%p1 bra实现分支kernel 参数通过.param传入地址计算通常需要 64 位运算索引要乘上元素字节数。2.2 编译器后端到底在做什么为什么值得绕开理解了 PTX 的样子再回头看编译器后端的工作会更清楚。从 LLVM IR 到 PTX 不是简单的一对一翻译它要干这些事指令选择同一个语义可以用不同指令实现。比如浮点乘加在 PTX 里是fma.rn.f32但 LLVM IR 里可能是muladd两条指令访存可以用ld.global.f32也可以合并成ld.global.v4.f32一次读 128 位。寄存器管理LLVM IR 里的虚拟寄存器是无限的但硬件物理寄存器有限。后端要为每个值选择合适寄存器并处理溢出。调度把访存指令提前、把无依赖的指令穿插在一起隐藏延迟。目标特化针对不同 GPU 架构选择不同的指令版本和寄存器数量。这些工作本质上是一套极其复杂的“规则系统”。规则系统的优点是可预测、可验证缺点是僵化——硬件更新时规则要跟着改优化空间也被限制在工程师预设的模板里。论文的赌注是LLM 可以从海量数据里学到比手写规则更普适的映射。人类写 PTX 时可以做到“我理解这个算法所以我用 shared memory 做 tile、用双缓冲隐藏延迟”LLM 如果训练得好理论上也能做到而且它看过的 kernel 可能比任何一个工程师都多。2.3 LLM 要写好 PTX必须理解哪些东西不是说会生成 PTX 语法就行。能通过 ptxas 编译的 PTX和能高效跑的 PTX完全是两个层次。一个好的 LLM 编译器至少需要理解执行模型GPU 有大量线程按 block 和 grid 组织线程之间有共享内存、有同步原语。忘掉bar.sync就会出数据竞争。内存层次global、shared、local、constant、texture 各有各的延迟和带宽。模型要会选不能总用 global。指令级能力fma、向量加载、ld.shared等指令的吞吐差异很大。边界条件kernel 的循环通常有边界判断漏掉一个if就可能越界。这也是为什么论文不可能做“从一句话直接到 PTX”这么野的事更合理的做法是给 LLM 输入一段已有 kernel让它在保持语义的前提下生成或优化 PTX。换句话说LLM 扮演的不是编译器前端的角色而是“后端的优化器 代码生成器”。3. 实操过程与核心环节实现这一节我按从论文和工程中能沉淀下来的路径来写。论文里未必会把每一个细节都写上但要做同类实验下面这些环节基本绕不开。3.1 训练数据怎么来源码-PTX 平行语料要让 LLM 直接写 PTX最稳的办法是用“CUDA 源码—对应 PTX”的平行语料做微调。你不需要手工写 PTX直接让 nvcc 编译就行nvcc -archsm_80 -ptx vec_add.cu -o vec_add.ptx厉害的话可以编译不同优化级别、不同 GPU 架构比如-O0和-O3或者sm_70和sm_86。这样一个 kernel 可以有多个 PTX 版本模型就能学到“同样的逻辑在不同架构上有什么不同写法”。更进阶的做法是把编译器优化前后的 PTX 都保留下来让 LLM 自己去对比。甚至可以把 SASS 反编译后的文本也混进去。但要注意SASS 文本噪音很大反汇编器的输出不一定能重新编译。所以最常见的还是nvcc -ptx产生的合法 PTX。训练样本的格式建议用类似指令微调的模板请将下面的 CUDA kernel 翻译成 PTX要求目标架构 sm_80使用向量化访存。 kernel __global__ void vec_add(...) { ... } /kernel然后输出就是对应 PTX。这里最容易被忽略的是数据清洗不能只收集能编译的 PTX还要确保它运行结果正确。很多 PTX 可能有编译错误或者语义错误混进训练集只会让模型学坏。我的建议是数据管线上至少跑两关第一关ptxas -c第二关在真 GPU 上用 CUDA driver 加载跑一遍和 reference 比对。能过这两关的样本才进训练集。3.2 推理时闭环生成-编译-运行-反馈训练完之后怎么用论文里的推理阶段通常不是一个“跑就完了”的过程而是一个带反馈的闭环。我复现过类似的流程非常简单但有效给定一个目标 kernel把它转成 prompt让 LLM 先生成 PTX 初稿用ptxas -archsm_80 -c检查语法把错误信息贴回去让 LLM 修改如果编译通过再用 CUDA driver 或者工具链加载 .cubin跑随机数据验证正确性如果输出错把错误信息或差异反馈给 LLM让它继续改正确通过后再做性能基准测试。这个闭环的关键是“反馈的颗粒度”。如果把一整页 ptxas 错误日志都塞给模型模型很容易迷失。更好的做法是先让脚本把错误归类比如“第 X 行找不到寄存器类型”“第 Y 行地址空间不匹配”然后只把前几个关键错误喂回去。实测下来这种方式能把修复成功率拉高不少。我写过一个最小验证脚本大致长这样# 假设 model_ptx.out 是 LLM 生成的 PTX ptxas -archsm_80 -c model_ptx.out -o model_ptx.cubin if [ $? -eq 0 ]; then echo compile ok else ptxas -archsm_80 -c model_ptx.out 21 | head -20 fi当然这只是第一步。编译通过不代表能跑能跑不代表性能好。所以我的习惯是把这个闭环做成一个自动化评估系统对每个 kernel批量生成 N 个候选 PTX逐一编译、运行、测评最后挑最好的那一个。这一步相当于把 LLM 当作候选取样器而不是一次定终身的生成器。3.3 评估指标怎么定做这类研究最忌讳只看“能不能生成 PTX”要有一套可复现的指标。我建议用下面这套指标说明怎么测编译成功率生成的 PTX 通过 ptxas 的比例统计ptxas -c退出码为 0 的比例执行正确率能拿到正确输出的比例CUDA driver 运行后与 CPU reference 比对性能归一化相对 nvcc -O2 生成版本的速度比反复运行 100 次取最短时间 / nvcc 时间单次生成耗时LLM 推理平均耗时记录 prompt 到完整 PTX 的时间修复轮数平均需要几轮反馈才能通过统计闭环中的迭代次数性能归一化是最能体现“AI 编译器”水平的核心指标。LLM 生成的 PTX 如果正确但比 nvcc 慢两倍那离落地还很远如果能打平甚至更快这就真的有意思了。据我看到的很多早期实验结果性能通常是劣势的——模型先保证语法正确就不错了优化是更难的坎。但这恰恰是论文的价值它把这个评估框架搭出来了后面的人可以在这个框架里去提升性能。4. 常见问题与排查技巧实录这一节是重头戏因为论文读起来很美但自己上手跑一遍全是坑。我把我踩过的、以及从同行那儿听到的高频问题整理一下。4.1 生成的 PTX 编译不过先降级再升级最常见的问题是ptxas报一堆错。比如ptxas error: File text of model_ptx.out is not valid PTX: Cannot find a register in class b32这类错误大多是操作数类型写错了。b32寄存器装的是 32 位整型你给它塞了一个 64 位值自然找不到合适寄存器。另一个高发问题是地址空间没写把ld.global.f32写成了ld.f32ptxas 根本不知道你从哪里 load。排查技巧先让 LLM 生成最朴素的 PTX不要一上来就向量化、不要用cp.async这种高级指令。朴素版本通常寄存器类型简单容易过把 ptxas 的报错信息截取前面几行原样返回给 LLM并明确要求它“只修改第 X 行附近不要动整体结构”如果反复失败检查 prompt 里是否明确指定了架构版本。同一个 PTX 在sm_80上成立放到sm_70上可能就不成立比如wgmma这种新指令在旧架构上根本不存在。4.2 编译过了但结果不对用 Sanitizer 抓妖怪编译成功只是万里长征第一步。我遇到过一个典型案例LLM 写了一个分块矩阵乘的 PTX看起来什么都有bar.sync也加了但输出全是乱的。后来用工具一查是有一个线程在bar.sync之前就把 shared memory 的值覆盖了典型的同步位置错误。排查工具随手可用compute-sanitizer --tool racecheck ./my_app compute-sanitizer --tool memcheck ./my_appracecheck能抓 shared memory racememcheck能抓越界访问。这两个工具比你肉眼 review PTX 快一百倍。如果都查不出问题那就做分步验证把 kernel 里一部分指令屏蔽掉跑中间结果对比。比如先只写一个线程关闭 grid看单线程的数据对不对再放大到 block逐个条件加回来。这个“增量正确性”排查法在汇编调试里非常管用。另外浮点问题不能只看 bitwise 相等。GPU 上abc和(ab)c的舍入结果可能不一样。比对时设置相对误差阈值比如 1e-5不要用memcmp。4.3 性能比不过编译器别急着放弃试试这三个方向第一次跑性能 benchmark 的时候LLM 生成的 PTX 大概率被 nvcc -O2 吊打。这时候最怕的是直接下结论“这条路走不通”。我试下来有三个方向能救指令选择引导在 prompt 里直接把想要的指令写死比如“计算内积时使用fma.rn.f32”“访存时使用ld.global.v4.f32”模型照做之后性能通常能上一大截多次采样做自动调优让 LLM 生成 20 个候选 PTX逐个跑 benchmark直接选最快的。这个思路的本质是“LLM 负责提供人的创造力benchmark 负责验证”把搜索交给外部工具人机协同让 LLM 生成框架 PTX人手动把最关键的循环改成优化版本再让 LLM 模拟你的写法去生成其他 kernel。这有点像“我给你一个 few-shot 示例你照着写”效果比纯零样本好很多。4.4 上下文窗口不够用拆而不是塞复杂 kernel 的 PTX 动辄几百行LLM 上下文窗口经常装不下。很多人第一反应是“换更大的模型”但更实际的办法是拆。比如矩阵乘 kernel可以拆成地址计算 helper 函数shared memory 加载函数核心计算循环写回函数。让 LLM 先分别生成这些局部的 PTX 片段每个片段都是完整 kernel 的一部分最后由脚本拼起来。拼接的时候要小心标签重复比如L_exit在不同片段里重复定义拼出来就编译不过。我的做法是让 LLM 在每个标签前加一个前缀比如$s0_L_exit、$s1_L_exit从源头上避免冲突。还有一个技巧使用更小的sm_XX目标。有些花哨的大指令只在高端架构上有而小架构的 PTX 更精简生成难度更低。等模型能稳定生成后再逐步提升架构。4.5 别忽略版本和工具链一致性这种底层汇编最怕版本漂移。ptxas版本、CUDA Toolkit 版本、GPU 驱动版本都会影响编译结果。同一段 PTX在 CUDA 11.8 上能过到 CUDA 12.4 可能因为指令名废弃而报错。我的经验是把所有实验固定在一个容器镜像里把nvcc --version和ptxas --version写进日志。这样哪天结果变了你知道是环境问题还是模型问题。5. 影响与后续可能5.1 编译器不会立刻消失但会被重新定义论文最大的贡献不是“证明 LLM 能写 PTX”而是把编译器后端这个看似固若金汤的黑盒打开了。传统编译器会以“pass 启发式”为核心继续存在很长时间但我相信未来会出现更多“LLM 生成候选 形式化验证器保底”的混合编译器。编译器工程师的角色也可能从写 pass 变成设计反馈回路和验证器。这和我看到的自动调优领域的发展路径很像早期 TVM 用模板 搜索后来有了 ML-based cost model现在 LLM 直接进汇编层每一步都是在扩大搜索空间的边界。5.2 对 AI Infra 的落地价值如果是做 AI Infra 的朋友这个方向值得长期跟进。我个人的判断是最有落地可能的不是“LLM 完整替代 nvcc”而是这几个场景手写 kernel 的自动优化用户写好一个朴素 CUDA kernelLLM 把它翻译成带优化技巧的 PTX再由专家 review快速原型验证在探索新算法时用 LLM 生成可运行的 PTX 验证想法而不是花几天手写汇编教学与反编译让 LLM 解释一段 PTX 在干什么或者把 PTX 还原成伪代码这个比直接生成 PTX 更容易在产品里落地。5.3 局限别把“能写汇编”当成“理解硬件”最后还是要泼一盆冷水。LLM 生成 PTX离真正的“硬件理解”还有很大距离。它可能学会了指令搭配但不懂微架构里的调度器行为它能复制常见 pattern但很难凭空发明一个全新的数据布局方案。更重要的是安全性和可验证性没解决之前没有任何生产系统敢让模型直接输出底层汇编并部署。所以这篇论文的意义更多在“探索可能性”而不是“发布一个可用产品”。把这篇论文反复读下来我最大的体会是让 LLM 直接写 PTX最大的价值其实不是“绕开编译器”而是逼我们把“什么是好的底层代码”这个问题想得更清楚。传统编译器用一个庞大的规则系统回答这个问题LLM 用概率模型回答两者都需要一个共同的裁判——硬件执行结果。所以如果你也想做类似方向我建议别一开始就追求“全自动生成高性能 PTX”先搭一个“生成-编译-运行-反馈”的闭环拿最简单的 kernel 跑通再慢慢加复杂度。这里面的坑远比论文里写的多但这种从底层验证 AI 能力边界的实验做一次比看一百篇综述都有用。