LLM绕过编译器后端直接生成PTX：一场GPU编程的范式实验

先说一个也许有点反直觉的结论：编译器的后端优化，可能是整个软件栈里最像“魔法”、又最不该被神化的部分。最近我读到一篇很有意思的论文，标题直白得不像学术论文——AI 就是编译器。它要做的事也很简单粗暴：让大语言模型直接生成 PTX，绕开传统编译器后端那一整套指令选择、寄存器分配、指令调度的流程。换句话说，以前我们写 CUDA 代码，要经过 nvcc 前端生成 PTX，再经过 ptxas 做底层优化变成 SASS；现在有人尝试让 LLM 直接输出 PTX，相当于把“编译器后端”整段跳过去。

这篇文章我前前后后读了三遍，又自己动手跑了一些小实验。今天想把里面最核心的思路、技术细节、以及我作为编译器方向从业者的一些真实看法拆开来讲。适合谁看？如果你在做 CUDA 优化、在搞 LLM 代码生成、或者单纯好奇“AI 到底能不能替代编译器”，这篇文章都能给你一些不一样的参考。我尽量不堆术语，但该讲清楚的底层原理一个都不会少。

1. 先搞明白：绕开的“编译器后端”到底是谁

要理解这篇论文为什么敢打“绕开整个编译器后端”的旗号，得先把一条 CUDA 代码从源码到真正跑在 GPU 上要经历什么流程捋清楚。很多人写 kernel 的时候只关心 nvcc 能不能编译通过、跑得快不快，很少去想中间那个“黑盒”里到底发生了哪些事。

1.1 一条 CUDA 代码要经历什么才能上 GPU

正常情况下，你写好的.cu文件会经过编译器的前端，把 CUDA C/C++ 语法解析、类型检查、语义分析之后，生成一种中间表示叫 PTX。PTX 全称是 Parallel Thread Execution，是 NVIDIA 定义的虚拟指令集架构（ISA），看起来像低层伪汇编，但并不是最终驱动硬件的机器码。

随之而来的才是真正的编译器后端环节。ptxas 会把 PTX 读进去，做指令选择、寄存器分配、指令调度、分支优化、循环优化等几十个 pass，最终生成当前 GPU 架构对应的 SASS 机器码，打包成 cubin 文件。程序运行时，再交给 GPU 驱动加载执行。如果是 JIT 模式的 CUDA 程序，这个过程可能会推迟到运行期，由驱动里的 JIT 编译器接过 PTX 继续编译。

nvcc -arch=sm_80 example.cu -ptx -o example.ptx nvcc -arch=sm_80 example.cu -cubin -o example.cubin

上面这两条命令差不多能说明问题：-ptx 拿到的是中间语言，-cubin 拿到的是经过后端处理完的二进制。日常开发里 99% 的人只会用默认方式整条编译链跑完，根本不关心 PTX 长什么样。

1.2 后端里装着的都是“硬骨头”

编译器后端不像前端那样——“看懂语法、生成中间码”就完事了。后端要解决的是一堆极其琐碎又互相牵制的组合优化问题。比如寄存器分配：GPU 上寄存器文件有限，spill 到 local memory 性能会雪崩；指令调度：不同的调度顺序会影响指令流水线利用率和访存延迟隐藏；还有向量化、合并访存、分支消除、peephole 优化等等。

这些优化里面最典型的是寄存器分配。你可以把 GPU 的寄存器文件想成一个酒店，一个 warp 的线程是住客，编译器是前台经理。好的经理能在满房和空房之间精确调配，既不浪费房间，又不会让人挤到走廊。编译器后端几十年来的工作，就是把这个分配方案做到极致，每差一个百分点的寄存器压力，可能直接影响 kernel 的 occupancy。

所以“绕过整个编译器后端”这句话听起来才那么唬人。它不是一个简单环节，而是编译器研究大几十年积累下来的核心壁垒。

1.3 为什么要盯上 PTX 而不是更底层的目标

论文的方案不是直接让 LLM 输出 SASS 机器码，而是输出 PTX。这个选择非常关键。PTX 本身是 NVIDIA 公开发布、有完整文档的 ISA，所有 CUDA 开发者都能读，NVIDIA 也允许第三方工具生成和修改 PTX；而 SASS 的编码细节跟具体 GPU 微架构强绑定，不同世代的 SM 之间差异很大，公开资料也不完整，让 LLM 直接去预测二进制并不现实。

PTX 的好处还在于可移植性。同样一份 PTX，可以直接跑在安培、图灵、Ada 等多代架构上，由驱动根据架构生成对应 SASS。这意味着 LLM 只需学习一个相对稳定的中间表示，不需要跟踪每个新 GPU 的机器码变化。

更重要的是，PTX 里带着比较丰富的语义信息，比如类型信息、谓词寄存器、操作数修饰符。LLM 学这种半结构化文本比学纯二进制容易得多，把 PTX 当成一个“目标语言”来做生成模型，跟训练一个代码翻译模型在形式上是类似的。

2. 论文核心方案拆解：LLM 怎么“写”PTX

大方向清楚了，接下来看具体怎么实现。论文里基本思路是端到端建模：输入是高级语言代码或者伪代码，输出是可直接交给 ptxas 的 PTX 文本。这一步相当于让 LLM 同时扮演编译器的前端和后端，把“C/CUDA → PTX”甚至“自然语言 → PTX”一步做完。

2.1 整体架构：从高层代码到 PTX 的一步到位

典型做法是用一个 7B 或 13B 的开源代码模型作为底座，比如 Code Llama 或者 StarCoder 这一类，然后在专门构造的 PTX 语料上做监督微调。输入序列可以是这样的：

输入： __global__ void vec_add(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]; }

输出：

.version 8.0 .target sm_80 .address_size 64 .visible .entry vec_add( .param .u64 .ptr .global .rdptr _param_0, .param .u64 .ptr .global .rdptr _param_1, .param .u64 .ptr .global .rdptr _param_2, .param .u32 _param_3 ) { .reg .b32 %r<6>; ... }

从任务形式上看，这更像是代码翻译任务、而不是传统的代码补全任务。模型必须学会把高层控制流、访存模式、内建函数一一映射到 PTX 指令和寄存器规则里去。难度比“生成普通 C 代码”高一个量级，因为 PTX 本身是不适合人读的，它对操作数类型、寄存器分配、控制流约束都非常敏感。

这类工作的创新点更多集中在数据构造和训练目标上，而不是网络结构。LLM 结构不需要做什么革命性改动，关键是训练数据能不能让模型学会“编译”这门手艺。

2.2 训练数据怎么造：这是整个项目的核心难点

没有现成的海量 (CUDA, PTX) 平行语料可以下载，这是所有此类项目里第一个要解决的问题。我看过的这类工作通常会从几个渠道拼数据。

第一个渠道是反向挖掘。把 GitHub 上开源的高性能 CUDA 项目里已经编出的 cubin 文件反汇编，用cuobjdump提取回 PTX，再和仓库里的源码配对。这样做的好处是数据全部来自真实场景，覆盖各种 SASS 级优化的结果；缺点是只有那些作者调过的 kernel 才有对应源码，而且 GitHub 上大量 CUDA 代码质量差异巨大，很多是教学示例或者过时代码，直接拿来训练会把模型带偏。

第二个渠道是编译器批量生成。拿 nvcc 对大量 CUDA 源码做不同优化等级、不同架构目标、不同编译选项的编译，自动得到对应的 PTX 变体。这样能产出非常大规模的对齐数据，缺点是同质化严重，模型容易学到的只是 nvcc 在常规代码上的行为模式，而不是真正的底层手工优化经验。

第三个渠道是人工增广。手工写一些有代表性的低层次 kernel，或者对已有 PTX 做寄存器分配扰动、指令重排、循环展开调整，制造出编译器通常不会自动生成的“人味”变体。这条路数据质量最高，但也最贵，论文里一般只用于小规模验证集。

高质量混合比例很关键。我接触过的经验是，纯编译器生成的语料会让模型变成“nvcc 模仿器”，泛化能力停留在流水线编译的范畴；加入手工优化数据之后，模型才可能学会一些弱优化技巧，比如利用向量化访存指令、提前预取 shared memory 等。

2.3 生成过程与约束处理：如何保证“能编译”

LLM 做代码生成有个老问题：模型不保证产出语法正确、类型合法、语义可编译的代码。这个问题在生成普通 Python 时还可以靠运行时报错一点点调，但生成 PTX 时错误成本高很多——PTX 的语法检查几乎零容错，一个寄存器的符号写错、一个.b64写成.b32，同一个 kernel 可能压根就编不出来。

所以这套系统里，模型的生成只是第一步，后面必须接一层验证与修复循环。常见做法是先生成一段 PTX，交给 ptxas 编译，如果不通过就把编译器的报错信息拼回 prompt，让 LLM 自己迭代修改。这个过程其实很像人类程序员用编译器报错来调试代码。一篇论文里如果写了自纠错机制，基本都会强调它能显著把编译通过率提上来，尤其是在-arch=sm_80这类严格模式下的修复能力。

注意：这里有个容易被低估的细节——PTX 的目标架构参数。LLM 生成的 PTX 里如果硬编码.target sm_80，在更老或者更新的 GPU 上就会出现兼容问题。所以生成阶段通常要强制模型输出里不带具体 target，留到实际部署时由驱动或者 ptxas 按需指定。

约束处理这块做得好不好，直接决定这个思路能不能从“实验室玩具”走向“真实工具”。我看过一些复位到工程模式的实现，会在生成器外面套一层 PTX 语法检查器，先砍掉明显非法指令，再用规则补全.version和.target头，最后才送 ptxas。这一步虽然听起来很“不 AI”，但往往是最能提升实际可用性的部分。

3. 效果评估：它到底行不行

论文效果不能只听作者自己说。我这篇文章里给你提供一套能自己复现的评估思路，以及我实测下来的感受。这里不写具体数字，因为模型、数据、训练步数、GPU 型号都会影响结果，但评估维度和方法是可以通用的。

3.1 正确性：先不聊性能，聊“能不能跑”

生成 PTX 的第一个门槛是能不能编译通过。你可以统计“生成后 ptxas 接受的比例”，在工程里叫 compile rate。一个称职的模型在这个指标上至少要做到 90% 以上，否则后续全流程没有意义。

第二步是 runtime correctness。把编译通过的 PTX 用驱动 API 加载（CUDA driver API 可以直接加载 PTX 字符串），在 GPU 上跑一遍，和用常规 CUDA 编译出来的参考 kernel 比对输出结果。这一步要注意：对浮点计算，两组结果不能要求 bitwise 完全相等，因为 PTX 里的指令顺序、FMA 合并策略可能不同，浮点累加顺序不一样会导致最后几位不同。通常做法是设一个很小的相对误差阈值，比如 1e-5。

我在自己的小实验里观察到的现象是：简单 kernel、比如向量加法和归约，LLM 生成 PTX 的编译率蛮高，正确性也容易通过。但一旦 kernel 涉及 shared memory 动态大小、barrier 同步、原子操作、更多样的访存模式，模型就开始频繁翻车，生成的 PTX 经常缺寄存器声明或者控制流标签对应不上。

3.2 性能：绕开编译器是否真的值得

性能才是这个方向最关键也最尴尬的指标。论文标题既然敢说“绕开整个编译器后端”，潜台词是 LLM 生成的 PTX 至少不能比传统后端差太多。但实测里事情没那么乐观。

传统编译器的寄存器分配和指令调度，是经历了数十年改进的成熟算法。LLM 要在这类组合优化问题上超越它是极其困难的。我在实验里看到的结果是：对于大部分常规 kernel，nvcc 生成的默认代码几乎总是优于或相当于 LLM 生成的 PTX。LLM 想赢，往往只能在少数场景，比如模型自己学到某种手工优化技巧，像循环展开配合共享内存平铺、或者巧妙地利用向量化加载，而这些技巧编译器在特定场景下也可能做得到。

不过反过来说，编译器生成的代码都是“安全的通用解”，它不会为了某个特定 GPU 架构做极端激进的优化。这里 LLM 其实存在理论优势：它可以学习海量人工调优过的 PTX 样本，把人类高手的“手艺活”直接学进去。只是目前基座模型的能力还不足以稳定复现这种高度上下文相关的优化决策。

实践建议：如果你想复现性能评估，别只看 kernel 运行时间。还要统计寄存器溢出次数（spill count）、共享内存占用、指令条数。这几个指标能帮你定位性能瓶颈到底在生成阶段还是调度阶段。用 NVIDIA Nsight Compute 做细节分析会清楚很多。

3.3 泛化能力：换一个 GPU 架构还行不行

另一个值得关注的问题是泛化。模型如果在 shader model 80（Ampere）的上训练，让它直接生成针对 Hopper 或 Blackwell 的 PTX，会不会崩盘？

得益于 PTX 本身的虚拟化设计，这个问题被淡化了一部分。同一份 PTX 可以运行在不同架构上，驱动负责映射到具体 SASS，所以 LLM 不一定需要学习新架构语法。但 PTX 版本之间是有差异的，新硬件往往引入新指令和新的语义，比如 Tensor Core 相关指令、warp 级别的协作指令等。模型没见过，自然生成不出来。

我建议在看这类论文时重点追问几个问题：

• 测试集的架构分布是什么？只是 Ampere 还是覆盖三代以上架构？
• 有没有做跨架构的零样本测试？
• 如果生成 PTX 依赖了某个新指令，如何保证该指令在老架构上能优雅降级？

如果这些问题在论文实验部分没有正面回答，那泛化性就是未定状态。

4. 值得玩味的点与可能的坑

这个方向真正触动我的地方，不是“LLM 替代编译器后端”这个宏大叙事，而是它逼着我们去重新思考编译器的边界在哪里。一个语言模型真的理解了“编译”吗？还是说它只是学会了做高频 pattern 的机械映射？这中间的差别，决定了这个方向的最终价值。

4.1 “绕开后端”不等于“消灭编译器”

论文标题是一个极简宣扬，但实际系统里 ptxas 依然不可缺失。LLM 生成的 PTX 要跑起来，还是要靠驱动里自带的 JIT 编译器把它翻译成 SASS。所以说到底，传统编译器后端只是“被绕开”了，并没有“被消灭”。

更准确的说法是：LLM 试图替代的是编译流水线里“高级语义 → 中间表示”这一整段决策过程，至于“中间表示 → 机器码”这部分，依然依赖成熟的底层工具链。这种分层思想恰恰是 PTX 设计的初衷——它是一个“合同层”，上面谁生成不重要，下面必须由厂商自有后端兜底。

这个设计其实挺聪明的。LLM 不需要面对真实硬件的千奇百怪，它只需要学会一种稳定的中间语言，剩下交给硬件后端。工程上这是“用 AI 换掉容易出错的映射层，保留不可绕过的验证层”的思路，比端到端生成机器码务实得多。

4.2 真正会取代的是什么

就现状而言，LLM 最有可能先落地的场景，不是替代 nvcc，而是作为编译器流水线里的“生成前端”。比如从一个高度领域化的 DSL（领域特定语言）自动生成 PTX，或者从用户自然语言描述出发生成高效 kernel 的雏形。在这些场景里，传统编译器反而拿不到好的输入，因为根本没有明确的 C/CUDA 源码可供编译，LLM 恰好补上这个空缺。

另一个很现实的场景是“帮编译器做生意”。传统编译器对未知代码做优化比较保守，LLM 可以在生成阶段就应用类似手工优化的策略，比如提前在算法层面选择更优的访存模式。这更像是在和编译器“合作”，而不是“替代”。我见过一些团队尝试把 LLM 生成的 PTX 跟 ptxas 的输出做融合，保留 LLM 生成的高层策略，同时用 ptxas 做底层寄存器分配，效果比纯 LLM 输出稳定很多。

关注 AI 编译器生态的话可以看到，MLIR、TVM、XLA 这些基础设施长期以来都在做“多层级中间表示”的事情。LLM 的介入不会推翻这些框架，反而可能成为这些框架中的一个新 pass。只是这些框架里的验证和回滚机制非常成熟，LLM 要想在里面站稳，必须先解决“生成结果可解释、可回归”的问题。

4.3 常见理解误区快查

这个方向讨论度一高，各种误读就跟来了。我整理了几个经常在评论区看到的问题：

这个表格我在很多场合发过，每次都有朋友说“早点看到就好了”。尤其是第一行，很多搞 AI 的同学把 PTX 和 SASS 混为一谈，导致后续很多分析都建立在错误的理解上。

4.4 从 GPU 延展到嵌入式与 CPU 编译器

虽然这篇论文聊的是 NVIDIA 的 PTX，但这类“端到端生成目标 ISA”的思路，在更宽泛的编译器生态里也在悄然生长。我看到社区里已经有人尝试用 LLM 直接生成嵌入式目标平台的汇编，比如针对 RISC-V 或者 CH32V 这类 MCU，让模型直接生成寄存器初始化序列和中断函数模板。

这类尝试的套路非常相似：先用编译器批量生成对齐语料，再用微调后的模型输出目标平台的汇编文本。有趣的是，嵌入式场景里开发者对性能的敏感度不像 GPU 那么极端，大家更在意“省事”“正确”“快”，所以 LLM 反而可能更快地落地成辅助工具。

不过嵌入式圈子的社区文化比 GPU 圈子保守得多。很多人对自动生成的汇编信任度很低，更愿意用一个带严格模板校验的“半自动化”方案，让 LLM 生成 C 代码或者寄存器配置片段，再用编译器作为最终验证关卡。这也是我很认同的落地节奏：让真实编译工具链当守门员，LLM 只当进攻手。

5. 想动手复现？给你几条实操建议

如果你读到这里还没被劝退，那说明你确实对这个方向感兴趣。下面给一套我自己跑过的、最小可行的复现路径。不需要 A100，不需要大规模训练集群，只要有一张能跑 CUDA 的 NVIDIA 显卡就行。

5.1 准备工具链和数据集

先确保你的环境里有 nvcc、ptxas、cuobjdump、nvdisasm 这套工具链。CUDA Toolkit 装好就自带这些。建一个文件夹专门放实验产物。

想快速体验“编译器后端到底做了什么”，可以用 cuobjdump 把现成二进制反汇编出中间表示：

cuobjdump -ptx example.exe -o example_exported.ptx cuobjdump -sass example.exe -o example_exported.sass nvdisasm -c example_exported.cubin -o example_exported.nvdisasm

其中example.exe是你用 nvcc 编译出来的可执行文件。这三条命令分别能拿到 nvcc 前端生成并内嵌的 PTX、最终机器码 SASS、以及带指令地址的反汇编文本。对比着看，你很快就能体会到 LLM 要学和要绕开的到底是什么东西。

把一批开源 CUDA 项目编译后用这种方式批量导出，就可以形成最初的训练/评估语料。清洗数据时建议至少做这几步：

• 去掉编译不通过的源码样本
• 过滤掉源码与 PTX 无法对应的大杂烩文件
• 按架构标记数据，比如 Ampere 的 sm_80 单独存放
• 对同一份源码保留多个优化等级生成的 PTX，增加多样性

5.2 搭建最小验证闭环

现在的关键需求是：能把 LLM 生成的 PTX 文本，快速变成 GPU 上能跑的结果。我最推荐的路径是用 CUDA Driver API，直接把 PTX 作为字符串加载，运行时让驱动做 JIT 编译。这比生成 cubin 再写 loader 简单得多。

下面是一个最小验证脚本的伪结构：

# 1. 模型生成 PTX python generate_ptx.py --input "vec_add kernel" --output gen.ptx # 2. ptxas 编译检查 ptxas -arch=sm_80 -c gen.ptx -o gen.cubin # 3. 用 driver API 加载 PTX 运行（编译 embedding 到测试程序里） nvcc test_driver.cu -o test_driver ./test_driver gen.ptx

test_driver.cu 里的核心逻辑是调用 cuModuleLoadData 把 PTX 字符串加载进来，再用 cuLaunchKernel 启动。这个流程不需要经过 nvcc 前端，所以能真正检验“绕过编译器前端”之后代码能不能跑。

陷阱提示：Driver API 加载 PTX 时，.target版本要和当前 GPU 兼容。如果你机器是 RTX 4090（sm_89），强行加载 sm_80 的 PTX 通常没问题，反过来就不行。为了省事，可以去掉 PTX 里的.target行，让驱动自动选择兼容目标。

5.3 评测指标与注意点

我自己评测生成结果时会保存三个数：compile rate、runtime pass rate、kernel time。还有一个容易被忽视的指标是“peek register count”和“spill count”，这俩能很好反映生成 PTX 的底层质量。

用 Nsight Compute 做 profile 时，重点关注 stall 原因分布。如果大量 stall 是 long scoreboard 或者 barrier，说明模型生成的访存调度不够好；如果大量 stall 集中在 wait，说明寄存器依赖链太长。这些细粒度指标比单纯跑一遍 time 更能定位问题。

追加评测建议：不要只跑一个输入规模。同一个 kernel 在不同 block size、grid size、数据量下，表现差异极大。LLM 生成的 PTX 经常只“适配”训练时常见的规模，一换参数就性能崩盘。这其实就是泛化能力不足的直接体现。

5.4 微调时的小技巧

如果你真想从零微调一个模型来生成 PTX，我有几个实际踩过坑之后的心得：

• 学习率要比常规代码生成任务低一个量级。PTX 语法太严格，任何一步震荡都可能让模型产生大面积的非法 token 模式。
• 训练时把 ptxas 编译是否通过做成一个额外的 loss 信号，虽然实现复杂，但效果远比单纯的语言建模损失好。没条件的话，至少要在后处理阶段用 ptxas 做过滤。
• 数据里专门保留一部分“错误修复示例”，即展示一段错误 PTX 和对应修改后的正确版本。这能显著提升模型利用编译器报错做修正的能力。
• 不要一上来就生成完整 kernel。先让模型学会生成一个空 kernel 的框架，再逐步填充寄存器、访存、同步逻辑，分阶段训练比端到端硬杠稳定很多。

6. 这类论文带来的一点思考

说回开头那个标题。AI 就是编译器，这句话在我实际验证过之后，给我的感觉更像是一个很好的研究宣言，而不是一个已经成立的事实。LLM 确实能学会像一个“初级编译器后端”那样生成 PTX 文本，它也确实在一些简单场景上表现出让人惊讶的正确性。但离替代 ptxas 里那几十个精心打磨的优化 pass，差距依然肉眼可见。

我个人在实际操作中体会最深的一点是：编译器后端的难点不在于某一条指令的选择，而在于“所有决策之间的相互制约”。寄存器分配影响调度，调度影响访存延迟隐藏，访存优化又反过来约束寄存器使用，这是一个环环相扣的组合优化问题。LLM 靠模式记忆很难统筹这一整条链路的全局一致性。

所以如果你问我会不会用手头这套方案替换现有 CUDA 构建流程，我的回答是：不会，至少现在不会。但这并不妨碍我认为这个方向价值极高——它把“编译器到底是怎么工作的”这个问题重新摆上了台面。过去几十年我们把这套工具链当基础设施用，很少有人质疑它的每一步决策是否合理。现在 AI 来了，逼着我们把指令选择、寄存器分配、调度约束这些概念重新拆开审视一遍，这本身就是很大的贡献。

最后再分享一个小技巧。如果你想快速获得“人类手工优化 vs 编译器生成 vs LLM 生成”三者的直觉差异，可以去 NVIDIA 的官方论坛和博客里找那些老鸟手写的 PTX 内嵌代码。很多人以为 PTX 只是编译器的一个中间产物，没人会手工写；实际上在极端性能优化的场景里，真有人手工调整 PTX 指令顺序、手动分配寄存器，就为了省几条指令。拿这些样本对比编译器默认输出，你会对“手工优化的价值天花板”有非常直观的认知。这个方向后续如果继续演进，我最期待的就是模型能把这些老手艺真正学会，而不是仅仅复制 nvcc 的常规行为。