本系列是论文和博客笔记 + 源码学习笔记，主要是关于 基于 LLM 的 GPU 内核代码生成。
随着AI辅助开发的模式的流行，大家开始尝试通过LLM来进行算子的生成和迁移，最近发现，PyTorch 发布了 KernelAgent，其尝试利用 Agent 来端到端实现torch模型优化及 Triton算子自动生成，关键做法是：多agent + 静态路由 + 子图分解 + 严格PASS验证 + 硬件profiler反馈。
因此，我找了一些论文和博客进行学习，具体可以参见文后“0xFF 参考”小节。
注：本文在几个月前就写好了，一直在草稿箱内，按照顺序，排到今天才放出来。
随着AI模型快速迭代与AI软硬件不断演进，AI产业对高质量算子的需求愈发强烈：
在深度学习模型的推理与训练过程中，绝大部分计算都依赖于底层计算内核（Kernel）来执行。计算内核是运行在硬件加速器（如 GPU、NPU、TPU）上的 “小型高性能程序”，它负责完成矩阵乘法、卷积、归一化等深度学习的核心算子运算。
高性能 GPU 内核主要依赖两种范式：手工编写和自动调优。
为什么 AI Infra 里还在大量招聘 GPU Kernel Engineer？ 论述了为何还需要Kernel Engineerr去手撸CUDA，具体如下：
虽然对 Kernel Engineer 的需求依旧强烈，但是人们也在思考：是否也能够借助LLM来模拟AI工程师的工作流程，凭借编译器反馈、硬件规格等丰富的信息，自动编写出准确且经过优化的内核代码呢？
实际上，LLM代码能力日益提升，为高性能 GPU 内核代码自动生成提供了一种极具潜力的方法。已有研究表明，现有 LLM 已具备一定的 GPU 内核生成能力。例如，英伟达工程师基于 DeepSeek-R1 设计了一套工作流程，在简单的 CUDA 内核生成任务中，该流程生成的内核在数值上全部正确，达到了 100% 的通过率。而且，在相关研究中，在完整的、复杂的机器学习架构上，平均加速达到了 50.8 倍！这验证了计算机体系结构的一个基本推论：系统的复杂性越高，自动优化的潜力就越大，因为人类工程师无法穷尽所有可能的优化组合。这里的性能提升是非线性的。
因此，基于LLM的算子生成技术逐渐成为业界研究热点。自2025年始，随着大模型的代码生成能力日益完善，用LLM来编写代码已在各类CodeAgent、AI IDE 中落地部署；在算子生成领域，各大知名高校、厂商纷纷开始投身到LLM生成算子的探索工作当中。
LLM生成算子的过程，就是给定一个PyTorch程序，让模型对其优化，然后生成一个包含自定义CUDA内核的PyTorch版本。在此期间中，模型可以自由决定优化哪些操作，以提高计算效率，其目标是自动生成正确且高性能的 GPU Kernel。
KernelBench是一个开源框架，旨在评估LLM在编写GPU内核方面的能力。KernelBench包含250个任务，涵盖了各种AI工作负载，并且易于扩展到新的工作负载。
下面给出 KernelBench 中的一个示例任务。每个任务都被封装在一个名为 Model 的类中。一个任务在 Model 类里主要包含两个核心函数：__init__和 forward；如有需要，还会提供辅助函数。AI算法通常在大型张量数据上进行操作。工作负载的最优内核取决于张量的大小和数据类型（如BF16、FP8），可以固定输入的 shape，仅通过随机生成的张量来改变数值。为此，我们提供了两个函数：get_inputs 和 get_init_inputs，分别用于生成初始化模型和运行前向传播时所需的随机参数。
下面给出某个模型针对上述任务规范进行优化后的示例输出。该模型不仅要生成 CUDA 内核代码，还要生成将其接入 PyTorch 框架所需的外围代码。评测框架会像调用普通 PyTorch 算子一样执行模型的 forward 函数，因此常见的做法是把 CUDA 代码直接内联到 Python 中。
算子开发领域主要存在四个挑战：
算子开发效率低，门槛高，依赖专业知识。手动方式来实现和优化算子时，开发者需同时掌握算法逻辑与硬件细节（内存架构、并行模型、指令集等），门槛极高，同时优化算子过程中，需通过反复实验调优线程块大小、tileing策略等参数，耗时且繁琐。
单一 LLM 生成算子存在正确性与性能缺陷。但是如果仅仅靠通用的LLM辅助开发的模式的话，不增加专业的软硬件知识和流程保证，易出现编译错误、数值不准确或性能不佳，难以满足生产需求。
跨平台与多 DSL 适配难度大。新的算子开发工具如triton/tilelang，目标是支持多种硬件平台，但是实际上，不同芯片类型的硬件架构差异还是很大的，很多算子实现还是会与硬件微架构耦合，迁移时需重新设计，可移植性差。
性能与开发效率/可移植性难以兼顾。追求开发效率与可移植性时，易牺牲算子性能，与人工优化后的算子性能差距大；如果要追求性能，现有优化方案多为单轮生成，无法系统性探索优化空间，性能提升有限。
我们从 KernelBench 入手来再深入。KernelBench 是让整个GPU编程自动化的起始催化剂。
下图展示了KernelBench评估语言模型（LM）生成高性能GPU内核的能力。KernelBench要求语言模型为给定的目标PyTorch模型架构生成优化的CUDA内核，并进行自动化评估。
这些任务根据包含的基本操作或PyTorch库函数的数量分为三个级别。
Level 1包含100个单个基本操作，如卷积、矩阵乘法等AI基础构建块。
Level 2包含100个操作序列，如卷积、ReLU和Bias的组合，这些操作可以融合成一个内核以提高性能。
Level 3包含50个完整的机器学习架构，如AlexNet和MiniGPT等，这些架构在运行训练和推理时对内核的性能要求极高。
在一次性基线评估中，LLM生成的内核平均在不到20%的任务中比PyTorch Eager更快。这表明，仅靠简单提示，LLM很难在性能上超越传统的PyTorch内核。另外，这些模型生成的内核存在大量执行错误、功能正确性问题，并且无法进行特定平台的优化。
具体来说，研究者发现：
我们再用论文 MultiKernelBench: A Multi-Platform Benchmark for Kernel Generation 作为印证，对于 GPU，论文得出了以下结论：
目前，基于 LLM 的 GPU 内核代码生成方法大致可分为两类：
目前猜测的几条趋势（可能有些已经过时）
从单agent走向多agent 编排
从“采样N次“走向“多轮RL+工具反馈”
从“求正确“走向“求实测加速”
严格沙箱与防reward hacking成为标配
领域专用语言（DSL）正在分化
“编译器 × LLM"协同设计兴起
基准开始“反通胀”
我们选择几个微调范式的示例来深入学习。
Kevin-32B（https://arxiv.org/pdf/2507.11948）由Stanford Cognition AI 实验室提出，生成语言为 CUDA，基座模型是 QwQ-32B，在 KernelBench L1 和 L2 级别展开实验。
论文作者利用 KernelBench L1 和 L2 上的 180 个题构建 RL 流程直接微调基座模型。使用的 RL 算法为 GRPO。训练流程分为了单轮训练（Single-Turn Training ）和多轮训练（Multi-Turn Training ）。
CUDA-L1: Improving CUDA Optimization via Contrastive Reinforcement Learning
CUDA-L1 的成功建立在一个精心设计的三阶段渐进式训练流程之上，旨在逐步提升 LLM 的 CUDA 编程与优化能力。三阶段像“模仿→比较→竞争”，每一步都把“能跑”→“较好”→“最好”做成独立里程碑；用对比学习做“护栏”，让强化学习始终在高速路而不是野地里飙车，最终拿到的模型既敢改代码，又不会改崩。
训练流水线总览如下：
基于微调（fine-tuning）的方法，在训练高性能 GPU 内核代码模型时面临多重瓶颈。
相比之下，Agent Pipeline 方法通过将生成过程拆解为多阶段模块，并结合编译器反馈、性能分析与工具链协作，实现了更灵活的优化和验证。但是，基于 Agent 的方法生成的代码很少是“纯血”Kernel（总是会调用高层 API，e.g. torch API）。
手工编写优化 GPU attention 内核既耗时又需要深厚经验。近期 LLM（如 DeepSeek-R1）在代码生成上表现亮眼，但仍难以一次性产出优化代码。因为把“写 GPU 内核”这件事直接扔给一次性的 LLM ，就像让一位刚毕业的学生闭眼一口气写完 500 行 CUDA——可能跑通，也可能带来灾难性后果。所以，需在推理阶段采用额外策略。
下面这段 prompt 是“relative positional embeddings attention kernel”任务里，用户扔给模型的示例输入。然而，LLM 偶尔会把不同语言或框架的语法“拌沙拉”，产出的代码要么编译不过，要么慢得离谱；而 GPU 线程怎么排布才算最优更是一门玄学，往往得“写-测-改”多轮才能拿到既对又快的内核。
Please write a GPU attention kernel to support relative position encodings. Implement the relative positional encoding on the fly within the kernel. The complete code should be returned, including the necessary modifications.
Use the following function to compute the relative positional encoding:
def relative_positional(score, b, h, q_idx, kv_idx):
return score + (q_idx - kv_idx)
When implementing the kernel, keep in mind that a constant scaling factor 1.44269504 should be applied to the relative positional encoding due to qk_scale = sm_scale * 1.44269504. The PyTorch reference does not need to scale the relative positional encoding, but in the GPU kernel, use:
qk = qk * qk_scale + rel_pos * 1.44269504
Please provide the complete updated kernel code that incorporates these changes, ensuring that the relative positional encoding is applied efficiently within the kernel operations.
而 Agent Pipeline 的做法是：把“写代码”拆成一条多角色、多步骤、可回滚的流水线，把“一次生成”变成“多轮对话 + 工具调用 + 验证反馈”，让 LLM 从“一次押宝”变成“持续迭代”。
另外，我们也能在 test-time scaling 上找到佐证。test-time scaling 是在推理阶段分配额外算力，让模型评估多条可能路径并选出最优解，从而模拟人类“分而治之”地拆解并解决复杂问题的过程。
具体而言，GPU 内核开发是「强硬件依赖、高专业门槛、多环节闭环」的复杂工程任务，单 LLM 智能体受限于「上下文理解上限、专业知识碎片化、无主动工具调用能力、无法闭环验证」等问题，只能生成「语法正确但性能拉胯、硬件不兼容、无工程可用性」的代码；而 Agent Pipeline 通过多智能体分工协作、流程化任务拆解、工具链闭环调用，让 GPU 内核生成从「单轮文本生成」升级为「端到端工程化开发」，最终输出硬件适配、性能达标、可直接部署的工业级 GPU 内核代码。
我们先来看看单 LLM 智能体生成 GPU 内核的核心痛点。
GPU 内核（CUDA/Triton）生成并非「简单代码编写」，而是需要硬件感知、性能调优、语法校验、真机验证、迭代优化的全流程工程能力，单 LLM 智能体在该场景下存在无法突破的先天缺陷：
GPU 内核开发需融合GPU 硬件架构（SM 版本、Tensor Core、共享内存）、CUDA/Triton 编程规范、算子优化原理（分块、访存、指令选择）、性能调优方法论四大类专业知识，且知识间高度耦合（如「分块大小需匹配 SM 共享内存」）。单 LLM 虽能记忆碎片化知识，但无法实现「知识的精准关联与场景化复用」，易生成「违反硬件约束」的代码。
GPU 内核的硬件适配性和性能有效性无法通过「文本推理」判断，必须依赖工具实时获取硬件信息、校验代码语法（CUDA Compiler、Triton 编译器）、真机运行验证性能（nvprof、torch.profiler）。单 LLM 是「被动生成型」智能体，无主动调用工具的能力，只能基于用户输入的模糊硬件信息生成代码，无法实现「实时硬件感知→针对性生成→工具验证→结果反馈」的闭环。
一个工业级 GPU 内核的生成任务，需拆解为「硬件信息采集→优化策略制定→内核代码编写→语法校验→真机性能测试→参数调优→最终定稿」多个子任务，且子任务间存在严格的依赖关系（如「先采集硬件信息，再制定优化策略」）。单 LLM 的上下文窗口有限，无法承载复杂任务的「流程记忆与逻辑衔接」，易出现子任务遗漏、步骤颠倒（如未校验语法直接生成代码），最终输出无工程价值。
GPU 内核优化是「试错 - 迭代」的过程 —— 即使初始代码语法正确，也可能因「分块不合理、访存效率低」导致性能不达标，需要基于真机测试结果调整参数（如修改 block size、调整精度）。单 LLM 是「单轮生成」模式，无法基于「性能测试反馈」进行针对性迭代，生成的代码往往是「理论可行，实际性能极差」。
GPU 内核生成通常存在多元优化目标（如「吞吐最大化 + 显存占用最小化」「延迟优先 + 硬件兼容性」），且目标间可能存在冲突。单 LLM 无「多目标权重权衡」的能力，易过度关注某一目标（如仅追求吞吐而忽略显存上限），导致代码无法满足实际场景需求。
Agent Pipeline 针对上述痛点，采用「任务拆解 - 专业分工 - 流程编排 - 工具闭环」的核心设计思路，将复杂的 GPU 内核生成任务拆解为多个「低复杂度、高专业性」的子任务，为每个子任务分配专属专业智能体，通过流水线式的协作与流程控制，实现端到端的工业级 GPU 内核生成。
比如， 下图是DeepSeek-R1 的迭代反馈Agent流程，这是NVIDIA，最早在 KernelBench 上做并 release 出结果的工作。
举例来说，Agent Pipeline的核心架构可以分为：5 大核心专业智能体 + 1 个流程调度智能体（GPU 内核生成专属）。
流水线的核心是「调度智能体统揽全局，专业智能体各司其职，工具链作为能力延伸」，各智能体通过统一的消息协议交互，子任务输出作为下一个子任务的输入，形成无人工干预的自动化流水线。以下是适配 GPU 内核（CUDA/Triton）生成的经典 Agent Pipeline 架构，各智能体的核心职责与工具调用能力高度贴合 GPU 开发场景：
| 智能体类型 | 核心职责 | 专属工具链（GPU 场景定制） | 与其他智能体的协作关系 |
|---|---|---|---|
| 流程调度智能体 | 全局任务编排，拆解子任务，分配给对应专业智能体，监控流水线进度，处理异常回滚 | 任务管理引擎、异常处理模块 | 调用所有专业智能体，接收各环节结果并决策下一步 |
| 硬件感知智能体 | 实时采集目标 GPU 硬件信息，生成标准化硬件报告，提炼核心优化约束 | torch.cuda API、nvidia-smi、NVIDIA 硬件特性映射库 |
向策略制定智能体输出「硬件信息 + 优化约束」 |
| 优化策略智能体 | 基于硬件报告制定专属优化策略（指令集、分块、精度、并行方式），明确性能目标 | GPU 架构 - 优化策略映射知识库、多目标权重权衡模块 | 向代码生成智能体输出「可执行优化策略」 |
| 内核编码智能体 | 基于优化策略编写 CUDA/Triton 内核代码，融入硬件适配逻辑和降级兜底策略 | CUDA/Triton 编程规范库、GPU 内核模板库（按算子 / 硬件分类） | 接收策略，向语法校验智能体输出「原始内核代码」 |
| 验证调优智能体 | 语法校验、真机编译运行、性能测试、参数自动调优，生成优化反馈 | CUDA Compiler、Triton 编译器、nvprof、torch.profiler、triton.autotune 调优引擎 | 向编码智能体输出「性能问题 + 调优建议」，最终输出达标代码 |
| 工程化封装智能体 | 对达标内核代码做工程化封装（类 / 函数封装、注释、测试代码、使用文档） | 代码规范检查工具、工程化模板库 | 接收最终代码，输出「可直接部署的工业级代码包」 |
GPU 内核生成的 Agent Pipeline 遵循「从硬件到软件、从策略到代码、从生成到验证、从调优到封装」的严格流程，子任务间无并行、无颠倒，确保每一步输出都为下一步提供可靠输入，具体流程如下：
Agent Pipeline 并非简单的「多智能体叠加」，而是通过专业分工、流程化协作、工具链闭环，解决了单智能体的先天缺陷，同时实现了「超越人工调优效率、媲美人工调优性能」的目标，其核心优势可概括为 6 点，且所有优势均深度贴合 GPU 内核生成的专业场景：
将 GPU 内核开发的四大类专业知识，精准分配给对应专业智能体（如硬件感知智能体掌握全量 GPU 架构知识，优化策略智能体掌握算子优化原理），每个智能体都是「领域专家」，能实现知识的精准记忆、场景化关联、专业级复用。相比单智能体的「碎片化知识堆砌」，Pipeline 能确保「硬件信息→优化策略→代码生成」的知识闭环，从根源上避免「违反硬件约束、性能优化无效」的问题。
所有专业智能体均具备主动工具调用能力，硬件信息无需用户手动输入（由硬件感知智能体实时采集），代码正确性无需人工校验（由验证调优智能体调用编译器），性能有效性无需人工测试（由验证调优智能体调用真机测试工具）。整个流程无人工干预，实现了「让代码生成过程自己 “感知硬件、验证自己、优化自己”」，这是单智能体和传统人工都无法实现的。
调度智能体将复杂的 GPU 内核生成任务，拆解为多个「低复杂度、小粒度」的子任务，每个子任务的处理都在 LLM 上下文窗口内，避免了逻辑断裂。同时，调度智能体通过严格的流程依赖控制，确保子任务按「硬件采集→策略制定→代码生成→验证调优」的顺序执行，无步骤遗漏、无逻辑颠倒，让复杂任务的执行更有序、更可靠。
Pipeline 引入「验证 - 反馈 - 迭代」的闭环机制，并非一次性生成代码，而是基于真机性能测试结果，持续调整优化策略和代码参数（如修改分块大小、调整精度、优化访存模式），直至性能达标。这种迭代能力让生成的代码不仅「语法正确、硬件兼容」，更能「性能达标、满足场景优化目标」，从「能用的玩具代码」升级为「好用的工业级代码」，解决了单智能体「一次性生成无优化」的痛点。
策略制定智能体内置多目标权重权衡模块，能根据用户输入的优化目标（如「吞吐优先」「显存优先」「延迟优先」），动态调整优化策略的权重（如吞吐优先则选择大分块，显存优先则选择低精度 + 层卸载）。相比单智能体的「顾此失彼」，Pipeline 能实现多目标的平衡优化，让生成的代码更贴合 LLM/RL/HPC 等实际工业场景的需求。
Agent Pipeline 采用 「模块化设计 + 标准化协议」，新增 GPU 架构（如 Blackwell）或新算子（如 UlyssesAttention）时，无需重构整个流水线，仅需更新「硬件感知智能体的硬件特性库」和「优化策略智能体的策略映射库」，即可快速适配新场景。同时，流水线能实时采集目标 GPU 硬件信息，生成针对性的代码，无需为不同 GPU 单独训练模型，实现了跨 GPU 架构、跨算子类型的泛化生成 。
KernelFalcon: Autonomous GPU Kernel Generation via Deep Agents
Automating GPU Kernel Generation with DeepSeek-R1 and Inference Time Scaling
DeepSeek-R1自写CUDA内核跑分屠榜！斯坦福学霸狂飙GPU编程自动化挑战人类
大模型能否为不同硬件平台生成高性能内核？南大、浙大提出跨平台内核生成评测框架MultiKernelBench
AKG kernel Agent：利用multi-agent进行kernel的生成和迁移
AKG KERNEL AGENT: A MULTI-AGENT FRAMEWORK FOR CROSS-PLATFORM KERNEL SYNTHESIS
RL 猛刷 CUDA 核：CUDA-L1: Improving CUDA Optimization via Contrastive Reinforcement Learning
MultiKernelBench: A Multi-Platform Benchmark for Kernel Generation
Ouyang A, Guo S, Arora S, et al. Kernelbench: Can llms write efficient gpu kernels?[J]. arXiv preprint arXiv:2502.10517, 2025.
Baronio, Carlo, et al. "Kevin: Multi-turn rl for generating cuda kernels."arXiv preprint arXiv:2507.11948(2025).
Li, Shangzhan, et al. "Autotriton: Automatic triton programming with reinforcement learning in llms."arXiv preprint arXiv:2507.05687(2025).
Li, Jianling, et al. "Tritonbench: Benchmarking large language model capabilities for generating triton operators."Findings of the Association for Computational Linguistics: ACL 2025. 2025.
Tjarko Lange, Robert, et al. "Towards Robust Agentic CUDA Kernel Benchmarking, Verification, and Optimization."arXiv e-prints(2025): arXiv-2509.
Chen, Wentao, et al. "CUDA-LLM: LLMs Can Write Efficient CUDA Kernels."arXiv preprint arXiv:2506.09092(2025).
此内容由惯性聚合(RSS阅读器)自动聚合整理，仅供阅读参考。 原文来自 — 版权归原作者所有。