摘要：Hazy Research 提出 CUDA DSL 可能逐步退休，因为 Agent 已能在测试与 profiler 反馈下直接生成目标优化代码。本文分析 megakernel、抽象复用、验证 oracle 和适用边界。
斯坦福 Hazy Research 在 2026 年 8 月发布的《Retire the Abstractions》很像一封写给自己作品的提前退休通知。两年前，他们还在介绍 ThunderKittens：一个嵌入 CUDA 的轻量 DSL，用 tile、共享内存布局和 warp 级操作，把 Hopper 上令人头疼的 WGMMA、TMA、寄存器排布包装成较短、较容易复用的代码。去年，他们又为了把完整的 Llama-1B 前向过程塞进一个 CUDA megakernel，构造了 GPU 端解释器、指令模板、共享内存分页和跨 SM 同步框架。如今，同一批研究者却说，CUDA DSL 可能是下一批“退休”的抽象，连 ThunderKittens 也不例外。
这项判断既没有宣告 CUDA 过时，也没有主张让大模型随手生成 kernel 后直接上线。文章提出的变化更微妙：过去，人类依靠代码抽象来卸载认知负担；现在，coding agent 可以在不完整的自然语言说明、参考实现、测试与 profiler 反馈之间反复工作，直接生成面向某张卡、某组 shape、某种模型结构的底层实现。于是，抽象作为认知卸载器的价值可能下降，规格、约束和验证系统的重要性随之上升。
这一区分决定了我们应当如何理解“退休”二字。可能退场的是固定实现和统一编程表面；知识本身与工程纪律仍会保留。
megakernel 为什么值得冒险
【事实】Hazy Research 所说的 megakernel 远超相邻两个 pointwise op 的常见 kernel fusion：它把一次模型前向中的几十乃至约一百个操作组织进单次 GPU kernel 启动。其 Llama-3.2-1B 实验针对 batch size 为 1 的低延迟解码：此时矩阵规模偏小，工作负载主要受模型权重读取带宽限制，传统推理引擎却要连续启动大量 kernel。前一个 kernel 的尾部 thread block、启动与撤销成本、下一阶段的数据加载延迟，会在时间线上形成一段段“气泡”。
他们在 H100 上测得，普通 CUDA stream 启动一个简单 kernel 的成本约为 2.1 微秒，CUDA Graph 可降到约 1.3 微秒，但在上百次边界面前仍然可观。团队把整个 Llama-1B 前向合并后，报告 H100 的显存带宽利用率达到 78%，相对 vLLM 超过 2.5 倍、相对 SGLang 超过 1.5 倍；B200 版本的单次前向低于 680 微秒。这些数字有明确前提：特定模型、单序列、低延迟、BF16、指定 GPU 与其测试栈。它们不能直接外推成“所有推理都快 2.5 倍”。
megakernel 的性能来源也不神秘。原来由 kernel 边界隐式保证的全局顺序，被改写成 kernel 内细粒度调度：当上一条指令还在存储结果时，下一条指令可以预取权重；MLP 中间结果可以切成多个 chunk，消费者不必等待整个张量完成；不同操作还能共享有限的 shared memory。Hazy 的实现为每个 SM 预先生成指令序列，在 GPU 端由解释器执行，并把 H100 的部分共享内存切成 16 KiB 页面，由指令显式申请、释放和转交。依赖关系则用全局内存中的计数器表达，生产者完成后递增计数，消费者等待目标值。
【判断】这类设计把“框架调度问题”变成了“单个并行程序的正确性问题”。它消除了 kernel 边界的税，却同时拆掉了 CUDA 原先免费提供的隔离栏。过去，错误常表现为某个算子输出不对；现在，错误可能来自页面过早复用、原子操作可见性、跨 SM 死锁、某种输入长度才出现的竞争，甚至只是某次编译后寄存器压力改变，导致原本漂亮的流水线崩掉。megakernel 的代价很明确：团队用难度更高的验证换取更紧凑的执行时间表。
Triton 与 ThunderKittens 如何分担工程复杂度
Triton 是理解这场争论的好参照。CUDA 的基本视角接近“标量程序、成组线程”：程序员为线程写逻辑，再亲自处理线程块、共享内存、同步与访问合并。Triton 把视角提升到 blocked program，程序员操作一块数据，编译器再完成线程映射、合并访问、预取、向量化、共享内存分配与部分 tensor core 指令选择。它远超简单的 Python 语法糖，把大量 GPU 调度知识固化进语言语义和编译器分析。
ThunderKittens 选择了另一处落点。它仍嵌在 CUDA 中，保留硬件可见性，同时把寄存器 tile、shared-memory tile、layout、warp/warp-group 操作包装起来。缺少能力时，开发者可以退回原生 CUDA。Hazy 团队曾用约百行 ThunderKittens 写出 H100 上的 FlashAttention-2 风格前向 kernel，并报告在其配置中超过当时的 FA2 基线。它的设计信念是：AI 算子与现代 GPU 都偏爱 tile，那么应当让 tile 成为程序员直接操作的基本对象。
这些 DSL 的价值至少有四层。它们压缩代码量；把 swizzle、bank conflict、barrier 等易错细节封装起来；为团队提供共同词汇；还把硬件知识变成可复查、可测试、可版本化的资产。最后两层尤其容易被“Agent 会写代码”的兴奋掩盖。抽象不只是帮助一个高手少敲几行字，它也让十支团队共享同一套 tile 语义、同一批回归测试和同一种 code review 尺度。
《Retire the Abstractions》的尖锐之处，在于它把 Agent 放到了编译器原先的位置之上。传统编译器接收精确语言，补全机器级细节；Agent 可以接收“不够精确但包含意图”的说明，阅读参考代码，调用编译器和 profiler，遇到错误再修改。Hazy 团队称，对于他们本来就能不用抽象完成的任务，例如优化 GEMM，Agent 在被告知 PTX 指令和 warp specialization 设计后，已能很快逼近先进实现。对于此前必须依赖框架才能管理的 megakernel，Agent 还做不到 one-shot，但可以把原来写进复杂 C++ 模板的部分规则改放在 prompt 中，陪着专家逐步解开竞争和死锁。
【判断】这里自动化的核心可以概括为搜索，代码生成只是其中一步。Agent 生成候选实现，编译器过滤语法与类型错误，测试过滤明显的语义错误，profiler 和计时器给出优化方向，专家再提供搜索空间中最关键的先验。GPU kernel 恰好拥有较密的反馈：能否编译、结果是否接近 reference、耗时多少、带宽与 occupancy 如何，许多指标都可机器读取。因此它比需求模糊、正确性难观测的业务系统更适合 agentic optimization。
Agent 实际上要怎样造出一个 megakernel
把过程展开，它更接近自动化实验室，和普通聊天窗口相去甚远。输入端需要一份可执行的 reference，以及清楚的部署分布：模型权重格式、batch 和序列长度范围、KV cache 布局、允许的 dtype、目标 GPU、延迟或吞吐目标。Agent 先建立朴素 CUDA 版本，确认调用边界和数值基线，再通过 profile 判断瓶颈究竟是 HBM、shared memory、指令发射、同步还是 launch gap。只有这一步成立，融合才有意义；若算子本来已接近计算上限，盲目做成 megakernel 可能只会增加寄存器压力。
随后进入结构搜索。Agent 可以尝试不同的 tile、warp 数、流水级数、producer-consumer 分工和 shared-memory 布局；对端到端融合，还要同时搜索 SM 的任务分配、指令顺序、buffer 生命周期与依赖计数。每轮都要经历编译、功能测试、sanitizer、基准和 profiler 分析。较慢方案并非简单删除，它们提供反例：例如更深预取因寄存器溢出变慢，更多融合因 occupancy 降低而失败，某种原子协议在长序列触发争用。这些轨迹应沉淀为下一轮生成的上下文，避免随一次对话结束而消失。
专家的角色也没有消失，其工作重心会由逐行实现转向制定搜索边界。专家会告诉 Agent：某代 GPU 上应优先验证 TMA，哪类矩阵形状值得用 tensor core，哪些同步不能依赖未文档化行为，什么 profile 形状才符合物理直觉。Hazy 所说“仍需告诉 Agent 使用什么 PTX、怎样做 warp specialization”，正说明当前系统主要降低了执行与迭代成本，尚未独立发明完整性能模型。
【判断】短期内，职责分离的流水线比一个全能 Agent 更可靠：生成者追求新方案，验证者主动寻找反例，基准器在隔离环境测量，发布器只接受满足策略的制品。即使这些角色由同一种模型承担，也应隔离上下文和权限，避免生成者为了保住自己的加速结果而放宽 tolerance、删掉失败 case 或改变 reference。人类 code review 中“作者不兼任唯一审批者”的原则，在机器生成时代仍然成立。
复用不会消失，只会换一种载体
如果每个模型、每张 GPU 都生成一份专用 CUDA，最直接的质疑是：还复用什么？Hazy 自己也承认，十个团队使用 ThunderKittens，可以共同积累 tile 语义的测试；十个团队分别生成 bespoke megakernel，则会得到十组互不分层的问题。
“代码不复用”不等于“没有复用”。未来可复用资产可能从模板库迁移为五类东西：
- 意图规格：算子语义、允许的数值误差、shape 范围、动态维度和别名规则；
- 不变量：哪些 barrier 前必须发生什么，哪些 buffer 的生命周期不得重叠，累加精度与确定性要求是什么；
- oracle：可信 reference、差分测试、性能模型、profile 图应该呈现的结构；
- 硬件知识：特定架构可用的 PTX 指令、对齐要求、TMA/WGMMA 约束、已知编译器缺陷；
- 生成与评测轨迹：哪些方案失败、在哪些 shape 退化、为何选定最终版本。
这仍然是一套“库”，其中的主角会由可链接的 C++ 模板转为可供 Agent 编译出实现的知识包。最终 CUDA 可以像目标文件或 autotuning cache 一样被保存，以便复现和部署，却不再被视为项目唯一的源代码。
【判断】这种复用模式对高度专用的内核很有吸引力，但它不会全面替代传统库。cuBLAS、cuDNN、CUTLASS、Triton kernel 或 ThunderKittens primitive 能把巨额优化成本摊到海量调用上，还承担 ABI 稳定、平台覆盖和长期兼容责任。若一个算子调用广泛、接口稳定，经过多年打磨的共享实现依旧有强大经济性。Agent 生成最有优势的地方，是“共享库覆盖不到、但运行规模足以支付优化成本”的长尾：新模型结构、固定 shape、罕见融合、端到端低延迟路径，以及新硬件刚发布而生态尚未跟上的窗口期。
稀缺资源转向 oracle
Hazy 的原文有一句比“DSL 要退休”更重要的话：只有当一个 oracle 能比抽象活得更久时，才有资格让那层抽象退休。
oracle 不能只是一段 PyTorch reference。KernelBench 将 LLM 生成 kernel 的评估拆成编译、正确性和性能三关，这是很好的最小框架；但它早期的典型正确性配置，是在预设 shape 上生成 5 组随机输入，以绝对、相对误差各 1e-2 进行比较。这样的设置适合低成本基准，却不足以证明生产正确性。候选 kernel 可能只对被测 shape 正确，可能越界但测试没触发，可能在 NaN、Inf、极值、非连续张量或特定对齐下失败，也可能存在低概率数据竞争。浮点并行归约还会让“相等”的定义依赖 dtype、累加顺序和业务容忍度。
面向生产的 oracle 至少应包含四条证据链。
其一是语义差分：覆盖边界 shape、随机 shape、异常数值、不同 stride 和真实分布样本，并明确 FP32/TF32/BF16/FP8 下的误差预算。其二是内存与并发安全：运行 compute-sanitizer 类工具，加入竞争检测、超时与死锁检测；对关键同步协议做状态机检查，必要时引入形式化验证。已有研究开始为 AI 生成 CUDA 证明内存安全、无数据竞争乃至语义等价，这恰好说明“跑过几组测试”离生产验收还有很远。其三是性能 oracle：不仅记录平均延迟，还看分位数、冷启动、功耗、显存占用、寄存器与 shared memory 使用，以及多 stream、并发请求下是否仍有收益。其四是部署一致性：锁定 GPU SKU、driver、CUDA、编译器 flags 和时钟策略，保留可复现的构建与基准环境。
验证还应防止 Agent 钻空子。计时前漏同步、复用陈旧输出、降低精度却沿用宽松阈值、只针对公开测试 shape 特化，都可能制造虚假加速。性能基准要隔离编译时间与运行时间，做 warm-up，随机化候选顺序，报告置信区间，并用独立隐藏测试集验收。对 megakernel 还要检查端到端时间线：局部 kernel 很快，不代表整体服务更快；占满 SM 也可能压制并发、破坏调度弹性。
【判断】Agent 让实现变便宜后，oracle 会成为新的成本中心。团队若没有能力写出比生成代码更可靠的验收系统，就不应删除原有抽象。自然语言 prompt 无法自动成为规格，reference 也未必给出正确答案；二者都可能携带历史 bug。信任必须“上移”，这意味着审查对象要由几千行 diff 转向规格完备性、测试强度和证据覆盖，审查力度不能下降。
性能会更快，也会更碎
传统 DSL 倾向于在可移植性、开发效率和峰值性能之间找公共点。Agent 可以为 H100 与 B200、短序列与长序列、固定 batch 与动态 batch 分别生成代码，因此更容易接受激进特化：常量折叠 shape，固定 tile，融合跨层操作，选择架构专属 PTX，甚至根据一次 profile 重写 pipeline。实现成本下降后，过去“不值得手调”的长尾配置也可能拥有专用 kernel。
代价是性能资产碎片化。一个在 H100 SXM、CUDA 某版本上胜出的 megakernel，换到 PCIe 型号、MIG、不同 driver 或下一代编译器后可能退化；固定 batch=1 的极低延迟实现，也可能牺牲吞吐与并发公平。生成系统必须把“适用域”当成一等公民：每个 artifact 都附带硬件指纹、shape contract、数值 contract、基准结果与回退路径。线上监控发现漂移时，应自动退回已知安全实现，禁止 Agent 在生产流量上即兴修改底层代码。
因此，较可信的形态应采用离线生成、沙箱编译、分层验证、基准竞赛、人工或策略门禁、签名发布，避免“一段 prompt 每次现编”。生成出的二进制仍应进入版本管理与制品仓库。所谓 disposable implementation，指的是允许重新生成，同时仍须保持可追溯。
这里还有一个容易被忽略的成本账。生成 kernel 消耗模型推理、GPU 编译和大量基准时间；验证复杂 megakernel 还会占用稀缺的新卡。若目标路径每天只运行几百次，节省的微秒可能永远抵不过搜索成本。相反，大规模在线推理、强化学习 rollout、实时语音和人机交互场景会把低延迟收益重复数十亿次，专用实现才有合理回报。团队应计算“优化投资回收期”，而不只是展示最快的一次 benchmark。
哪些团队现在适合尝试
最适合率先采用这条路线的，是拥有 GPU 性能专家、稳定工作负载和完善评测基础设施的模型系统团队。它们知道该提示 Agent 使用哪类 PTX、如何设计 warp specialization，也能看懂 Nsight 时间线是否合理；固定模型与固定集群让特化收益能被长期摊销；大规模推理的每几个百分点又足以覆盖生成和验证成本。研究新架构的团队也适合：Agent 可以较快补齐新算子的高性能实现，减少“理论更优、实现太慢”造成的误判。
中型团队可以采取混合策略：常规算子继续使用 PyTorch/Triton/成熟库，只把 profiler 证明为瓶颈、且输入域清晰的少数路径交给 Agent；要求每个生成 kernel 都有 reference、隐藏测试、基准报告和保守回退。这样能获取长尾特化收益，又不必接管整个 CUDA 工具链。
不适合立刻“退休抽象”的，是缺少 CUDA 专家、硬件型号繁杂、输入高度动态、强合规或安全关键的团队。对它们而言，DSL 不只是认知负担，也是知识传递、人员培训、review 边界和供应商责任的载体。若无法判断一个 barrier 为什么安全，Agent 给出的解释也不能替代证据。此时更合理的选择是让 Agent 在 Triton、CUTLASS 或受约束模板内生成代码，把搜索空间控制在已有语义与测试框架之中。
抽象退场后仍有契约
《Retire the Abstractions》最值得保留的是它对“项目是什么”的重新定义；对 CUDA DSL 时间表的大胆预测反倒居于其次。代码库过去同时承担意图、实现、知识、测试和协作界面。Agent 把实现的边际成本压低后，这些角色可能拆开：规格描述想要什么，不变量描述绝不能发生什么，oracle 决定何时可信，生成器负责为当前目标找到一份实现。
【判断】DSL 不会整齐地集体退场。它们会在不同位置变形：有些继续作为人类教学与协作语言；有些退到 Agent 的内部表示；有些成为限制生成空间的安全边界；还有些被拆成硬件知识库、测试集和性能模型。最终留下的“抽象地板”，正是纯粹的 contract——再聪明的执行者，也不能替需求方决定数值误差是否可接受、延迟和吞吐谁更重要、什么输入必须支持。
所以，Agent 生成 CUDA megakernel 的未来不会让软件工程退回手工作坊。恰恰相反：当代码可以批量生成，规格、验证、制品治理和回退机制必须比今天更严密。我们或许会少维护一层漂亮的 C++ 模板，却要维护一套更诚实的知识系统。
抽象可以退休。责任没有退休。
来源
- Hazy Research, Retire the Abstractions, 2026-08-05.
- Hazy Research, Look Ma, No Bubbles! Designing a Low-Latency Megakernel for Llama-1B, 2025-05-27.
- Hazy Research, GPUs Go Brrr, 2024-05-12.
- Triton Documentation, Programming Guide: Introduction.
- Stanford Scaling Intelligence Lab, KernelBench: Can LLMs Write GPU Kernels?, 2025.