Agent自动优化CUDA Kernel：冲上NVIDIA榜单第15名的实战拆解

我刚拿到这个题目的时候，第一反应是“这又是哪个榜单？”——毕竟叫 NVIDIA kernel 榜单一口气能冲到第 15 的事，圈子里很少有人公开拆过程。后来搞清楚了，是个公开的 GPU Kernel 性能评测排行，上面是各种算子的 CUDA 实现，按平均运行时的千分位比速度。以前我手动调过一阵子 CUTLASS，也踩过很多坑，所以对这种榜单一向是只敢看不敢玩。这次不一样，手头正好有一台带 A100 的机器，又一直在试各种 Agent 框架，于是咬咬牙决定做个实验：让 Agent 自己去写、编译、跑测试、改参数，24 小时不关机，看它到底能冲到什么位置。最后成绩是第 15 名，整个过程有惊无险，也踩了不少雷，所以特意整理一篇完整的拆解，给想用 Agent 做 CUDA Kernel 自动寻优的朋友少走弯路。

这个项目适合谁？如果你是做高性能计算的、写 CUDA Kernel 的，或者正在研究 AI Agent 如何落地到工程实操里，都可以从里面捞到一些思路。里面既有 Agent 的系统设计思路，也有 CUDA 优化里比较实用的指标和编译器参数，还有大量整夜调试时碰到的坑。我会把从环境搭建到最终提交的每一步都写清楚，不会有藏着掖着的地方。

1. 项目起步：目标设定与整体思路

1.1 榜单与任务界定

这个榜单的具体名字不方便多提，但可以描述成“一个面向 CUDA Kernel 优化挑战的公开排行榜”。评测方式很简单：给你一个预先写好的基础算子实现，比如一个大矩阵乘法的 GEMM 内核，然后你把跑得最快的优化版本提交上去，系统用统一的服务器、统一的 GPU 多次运行取中位数，最后按延迟排序。

我这次选的是一个偏计算密集的算子，官方基线代码是用最简单的 naive 方式写的，大概只用了单层循环，显存带宽利用率极低。基准得分大约在 12.5 ms。排行榜上第一名大概是 1.4 ms，第 15 名大概在 1.8 ms 左右。也就是说，要做的事很明确：在不改变计算数学结果的前提下，尽可能把运行时间往 1.8 ms 以下压。

用 Agent 来做这事，我当时定的目标不是“冲到前几名”，而是“让 Agent 在没有人工干预的情况下，独立完成几轮完整的、带有决策性的优化循环”。能跑到第 15 是意外之喜，但整个过程本身比排名更有价值，因为它验证了 Agent 在代码优化这件事上不是玩具。

1.2 为什么选择 Agent 而不是纯手工

我过去手动调过 GEMM，坦率讲，每次都是同样的流程：看 profile 结果，猜一个优化方向，改代码，重新编译，跑 benchmark，然后盯着数字叹气。一个 block size 参数我就得试 8 种组合，每次编译加运行要 40 秒，一晚上也就能试几百组。而且注意力很难维持，经常调着调着就忘了之前哪组参数效果最好。

Agent 的核心价值在于可以把“探索参数空间”这件无聊的事自动化，并且能利用大语言模型对“什么样的优化方向更合理”做出相对聪明的判断。它不像暴力搜索那样无脑试，而是每一轮都带着当前瓶颈数据去提出下一步假设。我搭的 Agent 在一个小时里就跑了超过 300 个变体，这个量是手动完全不现实的。

当然，Agent 不是万能的。你需要给它非常明确的任务边界和评价函数，否则它会像无头苍蝇一样乱编译。后面我会详细说明这套系统怎么落地。

1.3 整体技术选型

先说硬件。机器是单卡 A100 40G，内存 512G，Ubuntu 22.04，NVIDIA 驱动版本是 535.183.06，CUDA Toolkit 12.3。很多人会在装驱动这边卡住，比如“nvidia 控制面板找不到了”“nvidia-smi 正常但代码编译失败”，这些问题我这轮也碰到过，后面专门写一节解决经验。

软件层面我没有用特别重的 Agent 框架，而是选择了自己拼一个轻量级循环，配合 OpenAI 的 GPT-4 API 做代码生成和反思，然后用 Optuna 做参数搜索。之所以不直接上 LangChain 或 AutoGen，是因为 Kernel 优化这个场景里，主流程的循环是固定的：生成、编译、跑分、反馈。自己写循环能省掉很多框架层的抽象，也更容易排查问题。

CUDA 编译器用 nvcc，性能分析用 Nsight Compute，计时直接用 CUDA Event，不额外的计时器。最终提交的验证脚本是榜单平台提供的，保证结果可复现。

2. Agent 寻优系统的核心设计

2.1 任务规划模块：拆解目标

Agent 的顶层是一个状态机，核心节点包括：分析基线、生成变体、编译执行、性能评估、瓶颈归因、策略调整。每个节点之间传递一个 JSON 格式的上下文数据包，里面包含当前 kernel 代码、编译参数、运行时间、Nsight Compute 的关键指标，以及前几步的历史记录。

为了让 LLM 能合理规划，我给它写了一套系统提示词，要求它把所有优化动作按风险等级分成几类：低风险的是调参数（block size、unroll factor），中风险是改循环结构和 shared memory 布局，高风险是彻底改写算法（比如从 tiling 改成寄存器阻塞）。Agent 每轮必须先从低风险开始，只有当收益曲线平坦时，才允许跳到高风险。

这个设计是因为我发现：让 LLM 一开始就自由发挥，它很容易生成一个“看起来很高端但完全跑不动”的 kernel。比如连续好几轮都在尝试 double buffer，却忽略了这个算子的瓶颈其实是在内存带宽，导致代码复杂度和实际性能完全不成比例。

2.2 代码生成与编译回路

代码生成环节，我采用的是“模板 + LLM 填空”的方法。先准备一个完整的 CUDA Kernel 模板，里面用宏定义控制关键参数，LLM 的任务是生成一组宏定义值，以及在注释<LLM_CODE>标记的位置插入循环级优化段。

这样的好处是编译出错的概率低。你让 LLM 直接输出一个完整 .cu 文件，经常遇到头文件缺失、函数签名不一致这问题，很容易把一晚上的时间全消耗在 Debug 上。用模板后，Agent 的生成内容被限制在一个安全范围内，编译失败率从 50% 降到了 15% 左右。

编译命令用的是这条：

nvcc -arch=sm_80 -O3 --use_fast_math -Xptxas -v -maxrregcount=128 -dlto kernel.cu -o bench

其中-Xptxas -v会把寄存器用量和 shared memory 使用量写到 stderr，Agent 会把这些 log 解析出来放进上下文。如果寄存器溢出，它会自动尝试调低maxrregcount，这是后期迭代很关键的一步。

2.3 性能评估与反馈机制

性能评估的工具是 CUDA Event 计时和 Nsight Compute。我写了一个简单的 runner，把一个 kernel 在固定输入下连续跑 50 次，取中位数作为最终得分。为什么不用平均值？因为 GPU 调度会有毛刺，平均值容易被极端值带偏，中位数更稳定。

Nsight Compute 也不是每轮都跑，因为它太耗时。我的策略是：每隔 5 轮跑一次完整的 profile，提取几个关键指标：achieved occupancy、smem bank conflict 次数、DRAM throughput utilization、L2 命中率。这些指标会被压缩成一段文本反馈给 LLM。

Agent 的反思节点会结合这些指标产生类似这样的判断：

• “当前 occupancy 只有 36%，说明并行度不足，应该减小 block size 或调整循环到 grid 维度。”
• “DRAM throughput 已经到 92%，说明访存瓶颈明显，下一步优先考虑向量化加载或启用 L2 persistence。”

为了让 LLM 能理解这些数据，我把所有指标都转成语义描述，而不是直接丢给它数字表格。实测下来，这类“翻译”能明显提高建议的可用性。

2.4 搜索策略与自动寻优

除了让 LLM 生成代码，参数搜索也同等重要。我定义了一个参数空间：block_size（64、128、256、512）、vector_width（1、2、4）、unroll_factor（1、2、4、8）、smem_size（0、16KB、32KB、48KB）。空间不算很大，但如果纯随机搜索，很容易错过最优组合。

我分两阶段：先用 Optuna 的 TPE 采样器跑 200 组随机参数，探出基本有潜力的区域；然后改用 CMA-ES 风格进化策略，对代码里的宏定义做小幅变异。这就像一个双保险，LLM 负责“语义级优化”，搜索器负责“数值级微调”。

在两个策略之上，Agent 还会维护一个“历史参数黑名单”。比如某一个 block_size 组合连续 3 次触发编译失败，或者性能比当前 best 差 5 倍以上，就直接加入黑名单，后续不再生成。这个机制看起来简单，但实际省掉了很多无意义的编译时间。

3. 实际操作过程：24 小时全记录

3.1 环境准备与基线测试

环境这部分我花了大概两小时，因为踩了几个软件层面的坑。首先是在 Ubuntu 上装驱动。我用的显卡是 A100，所以驱动版本不能太老。最开始手动装一个 535 的 driver，结果 nvidia-smi 正常，但一用 PyTorch 就报 CUDA error。后来又折腾了“nvidia 驱动安装”过程中常见的黑屏问题，最后是直接把驱动用apt install nvidia-driver-535-server重装一遍才稳。

这里有一个很容易被忽略的细节：如果装了新的 CUDA Toolkit，一定检查一下nvcc --version和nvidia-smi里的驱动版本是否匹配。A100 的 sm_80 架构需要 CUDA 11.1 以上，我直接用 12.3 就没问题。如果用的 4090 这种 Ada 架构，就得让 nvcc 指定arch=sm_89。

基线测试跑得很快，官方给出的 naive Kernel 时间为 12.5 ms。我当时对这个数字有点无语，因为同样是在 A100 上跑一个 4096 x 4096 的 FP32 GEMM，CUBLAS 能做到 1.2 ms 左右。这说明官方结题代码留下了巨大的优化空间，Agent 有得玩。

3.2 关键参数与算子分析

正式跑 Agent 之前，我先手动分析了一遍算子。计算类型是 FP32，矩阵边长 4096，存在较明显的访存密集特征。naive 实现的问题在于每个线程重复读取多次全局内存，并且循环结构完全没利用 GPU 的并行层次。

我给 Agent 提供的第一份分析报告里写清了三个优先方向：

• 使用 32x32 的 block tile 降低全局内存流量
• 对 A 和 B 矩阵分别使用__ldg()和向量化 float4 加载
• 在共享内存里做一次数据布局转换，解决后续 bank conflict

这个其实是我手动调 GEMM 的老套路。把这份报告放在 Agent 上下文中，是为了避免它一开始就去搞什么 warp specialiation 或者 split-K 之类的进阶技巧，先把基础的 tiling 做好再说。

3.3 优化迭代记录（小时级）

下面是这次迭代的真实时间线，我记录在项目 log 里：

从时间线能看到，Agent 并没有一直保持高速增长。它在 12 小时左右卡在 3.7 ms 附近，连续两个小时没有任何变体能在 1% 的误差范围内超过当前最好结果。我当时差点想介入去手动调，但还是忍住了，结果 15 轮之后，Optuna 找到一个组合block_size=256, unroll=4, maxrregcount=96，一下子把时间拉到 2.8 ms。

这给了一个启发：Agent 自动寻优的曲线不是平滑下降的，经常是长时间不动，然后突然下跌。你必须有足够的耐心，并且设置合理的“无进展中断”阈值。我是允许连续 60 轮无提升才自动暂停，否则会提前错过那波爆发。

3.4 最终提交与排名冲刺

到了第 22 小时，Agent 已经稳定跑出 1.82 ms 的中位数成绩。但提交之前还有一道坎，就是正确性验证。平台要求输出结果和原版 Kernel 对比的 max relative error 小于 1e-3。Agent 生成的某些代码为了用-use_fast_math，把精度搞坏了，误差达到 2e-2，导致不通过。

我设置了一个自动化检查：每个变体跑完后，除了算性能，还要跑一个小的随机输入测试，对比 CPU 参考结果。如果误差超标，哪怕性能再高也直接丢弃。这样一来，最终提交的变体是经过 1200 多个候选里筛出来的，既快又满足精度约束。

提交时我手动看了下排名，发现自己已经跳到了第 15。后面五个小时其实没有再动，因为排名已经稳定。最后我关掉 Agent 时，它一共尝试了 1378 个变体，其中有 214 个编译成功，47 个正确性验证通过，12 个比当前最好成绩更快。整个过程只消耗了 3 次人工干预，而且都是因为 GPU 温度过热导致测试环境异常，跟代码逻辑无关。

4. 常见问题与排查技巧实录

4.1 编译器与运行时问题

这一轮里最常碰到的编译问题是寄存器溢出。当maxrregcount设为 255 时，-Xptxas -v会提示 “REGISTER SPILL”，意思是在高并行度场景下寄存器用太多，导致部分变量被挤到本地内存。本地内存访问会让性能和全局内存差不多，非常坑。

解决办法有两类：一是降低每个线程的寄存器预算，比如设成 128 或 96；二是调整 block size 让每个线程处理的元素少一点。很多新手看到maxrregcount就无脑调大，实际上这个值要和 occupancy 一起权衡。Agent 的做法很直接：如果 spill 出现，就把maxrregcount减 16 然后重新编译，直到 spill 消失，再小范围浮动找出性能峰值。

另一个坑是 CUDA 12.3 和 Ubuntu 22.04 组合下，偶尔会报 “Cannot find an eligible device” 的错误。这个问题一般不是驱动坏了，而是运行环境中LD_LIBRARY_PATH指向了 PyTorch 自带的 CUDA 运行时库，导致版本冲突。我的规避方式是在 runner 脚本里显式export LD_LIBRARY_PATH=/usr/local/cuda/lib64:$LD_LIBRARY_PATH，并且不要全局 source 一个乱写的 PATH。

4.2 Agent 陷入局部最优的应对

局部最优是自动寻优系统最烦的问题。在我这个项目里，Agent 在 3.7 ms 附近卡了两个多小时，后来分析日志发现，它一直在围绕同一个block_size=256 + unroll=2的模板做微调，本质上是在同一个山坡上反复爬，根本没有跨到其他地貌。

应对方案是我设计了一个“思想实验”触发器。当超 40 轮无提升时，Agent 会强制丢弃当前 best 相关代码，直接从基线代码开始，但带着一个新的优化风格提示词，比如“尝试用浮点飞马算法减少乘除操作”或者“尝试用 8-stage pipeline”。这个思路后来帮我找到了 double buffering 这条路，因为之前的代码结构完全是单缓冲的。

不要害怕丢掉当前最好成绩，局部最优丢掉并不可惜，怕的是永远没有机会跳到更好的区域。当然，自动丢弃 best 是有代价的，所以我在代码里做了一个备份：当前 best 会永久保存到一个hall_of_fame/目录，即使 Agent 后来越调越差，最差也只是回退到之前的成绩，不会掉出排行榜。

4.3 性能抖动与数据可信度

评测连续跑 50 次取中位数，看起来挺稳妥，但前几次跑分经常偏高。后来查了原因，是 GPU 要经历一个“预热”阶段，SM 频率、内存控制器状态都要稳定下来。我的 runner 会在正式计时前先跑 5 次空跑，让 GPU 进入稳定状态，然后再计时。

还有一次比较困惑的问题是在同一台机器上，Agent 某个变体某次测出 1.7 ms，之后连续几次都是 2.1 ms。我一开始怀疑 Agent 作弊或者缓存，后来发现是我的 A100 的 PCIe 带宽被别的任务抢占了。排查方法很简单：用nvidia-smi看有没有其他进程占用了 GPU 显存或 sm。从那时起，我就在 runner 里加了一句nvidia-smi --query-gpu=utilization.gpu,memory.used --format=csv的日志记录，这个技巧强烈建议保留。

4.4 实用避坑经验

这里写几条我这趟下来最值得记住的避坑经验。

第一，不要在一开始就引入复杂的 Agent 框架。先自己手写一个 100 行左右的循环，验证你的 kernel 模板和环境没问题之后，再加LLM推理和搜索策略。框架的抽象会把 bug 藏得很深，你反而不知道是 Agent 的问题还是框架的问题。

第二，所有中间产物必须落盘。每个候选 kernel 的代码、编译命令、stderr 输出、跑分表格都要写到独立文件目录。这样即使 Agent 崩溃，你还能从崩溃点附近恢复。我这次就是有个晚上进程意外退出，靠着 log 回放才恢复迭代。

第三，给你的 Agent 设定一个时间预算。我用了一个全局计时器，每轮迭代之后判断剩余时间和当前性能，如果剩余时间不足 2 小时，就自动切到“微调模式”，只改参数不再动代码结构。这保证了最后提交前不会因为一个高风险实验把排名丢掉。

第四，千万不要忽略正确性校验。性能再高的 Kernel，如果输出误差超过阈值，提交上去也会被判无效，白白浪费几个小时。我把正确性校验放在性能测试之前，这样能节省大量时间为无效变体跑分。

5. 最后再分享一个小技巧

前面这些写完后，我其实还想补充一个经验：如果你真的想复制这个项目，最关键的其实不是 Agent 框架或 LLM 选择，而是怎么把 GPU 性能数据转换成 LLM 能“读得懂”的语言。我最初直接把 Nsight 的一堆原始指标丢给 GPT-4，它给出的建议非常空泛。后来我把指标转成类似“smem bank conflict 占比超过 40%”“warp stall 原因中 long scoreboard 占比最高”这样带因果关系的描述后，Agent 的优化成功率直接翻倍。

这件事让我重新认识了 Agent 自动寻优的本质：它不是一个“全自动黑盒”，而是你用工程化的方式把行业经验固化成一种可交互的上下文。真正要动脑的地方，还是在你定义任务、设计反馈回路、判断边界条件这些环节。以后如果有机会，我还想试试更复杂一点的算子，比如 attention 或卷积，并且用同样的系统打一遍榜，看看能不能冲进前十。