我刚拿到这个题目的时候第一反应是“这又是哪个榜单”——毕竟叫 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内存 512GUbuntu 22.04NVIDIA 驱动版本是 535.183.06CUDA 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 -archsm_80 -O3 --use_fast_math -Xptxas -v -maxrregcount128 -dlto kernel.cu -o bench其中-Xptxas -v会把寄存器用量和 shared memory 使用量写到 stderrAgent 会把这些 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_size64、128、256、512、vector_width1、2、4、unroll_factor1、2、4、8、smem_size0、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 指定archsm_89。基线测试跑得很快官方给出的 naive Kernel 时间为 12.5 ms。我当时对这个数字有点无语因为同样是在 A100 上跑一个 4096 x 4096 的 FP32 GEMMCUBLAS 能做到 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 里时间段主要动作性能变化备注0-2h环境搭建 基线测试12.5 ms卡在驱动上但没影响整体2-4hAgent 首轮随机搜索 180 组8.1 ms发现 block_size512 效果最好4-8hLLM 模板生成 tiling 内核4.3 ms共享内存 向量化加载落地8-12hNsight 显示 bank conflict 偏高3.7 ms修改 swizzle 布局12-16hOptuna 调 unroll 和 maxrregcount2.5 ms组合找到后提升明显16-20h引入 double buffering1.9 ms与预期一致带宽打满20-24h最终精度校验 连续三次独立测时1.82 ms冲上第 15从时间线能看到Agent 并没有一直保持高速增长。它在 12 小时左右卡在 3.7 ms 附近连续两个小时没有任何变体能在 1% 的误差范围内超过当前最好结果。我当时差点想介入去手动调但还是忍住了结果 15 轮之后Optuna 找到一个组合block_size256, unroll4, maxrregcount96一下子把时间拉到 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_size256 unroll2的模板做微调本质上是在同一个山坡上反复爬根本没有跨到其他地貌。应对方案是我设计了一个“思想实验”触发器。当超 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-gpuutilization.gpu,memory.used --formatcsv的日志记录这个技巧强烈建议保留。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 或卷积并且用同样的系统打一遍榜看看能不能冲进前十。