我不太确定自己当时是不是真的能冲上那张榜单只是隐约觉得如果能用一个 Agent 在 24 小时内自动把 CUDA kernel 调优搞定这件事本身就是值得记录的。上周我试了一把一个专门为 NVIDIA kernel 性能优化设计的 Agent在我的一张 A100 卡上连续跑了 24 小时在公开的 kernel 性能基准榜上从最初的 57 名一路爬到了 15 名。整个过程没有人工改过一行 kernel 代码所有优化决策都是 Agent 自己做出来的。这篇文章就是把这次自动寻优的完整过程、技术选型和踩坑经验一次性拆开写给同样在搞 GPU kernel 优化、或者正在做 Agent 开发的人。1. 项目目标Agent 到底在解决什么问题先说榜单背景。这个榜单是一个半开放社区的 kernel 性能评测排行榜基于一组固定的 benchmark 输入收集不同人提交的 CUDA kernel 实现按平均运行时间排序。类似 NBody、GEMM、卷积这些经典算子都覆盖每天自动更新排名。我盯了很久想尝试一下“不手动调参只靠 Agent 自主寻优”能不能在这个榜单上打出成绩。为什么不手动调因为 CUDA kernel 的优化空间实在太大。gridDim、blockDim、寄存器分配、向量化宽度、循环展开因子、共享内存大小、访存模式每一项排列组合起来都是天文数字。人类专家能靠直觉快速排除一部分但面对一个陌生 kernel 时直觉也会失效。尤其是现代 GPU 架构A100 和 H100 的行为差异很大一个在旧卡上最优的配置换新卡可能直接性能腰斩。Agent 在这里解决的是“在线决策”问题它不是一次性跑几百个配置取最优而是边跑边学根据每次实验的反馈决定下一步生成什么样的变体。这个思路和传统 autotuner如 TVM Ansor、Halide不同后者更多是在一个已经定义好的模板里搜索参数而 Agent 还能修改 kernel 的结构比如改变循环顺序、改 tiling 策略、合并核函数。这才是标题里“自动寻优”比“自动调参”更准确的原因。榜单排名的波动性也很有意思。每天都有新提交前 20 名的时延差距常常在 1% 以内一个优秀的优化可能只让你上升几名。所以我最初的目标不是第 15而是“能稳定进入前 20 就算成功”。最后跑到第 15属于意料之外但也说明这套 Agent 方法确实有效。2. 系统设计Agent 框架选型与寻优策略2.1 先分清两个人人容易混淆的概念harness 和 agent很多刚接触 Agent 开发的人会把 harness 和 Agent 混在一起。简单说harness 是“执行容器”负责编译、加载、运行、计时、收集指标Agent 是“决策大脑”负责分析指标、生成下一个 kernel 候选。两者通过一组标准接口通信。在这次项目里harness 是一个 Python 服务封装了 nvcc 编译、CUDA event 计时、Nsight Compute 调用以及 submission 脚本。Agent 则是一套基于 LLM 的决策循环它不能直接执行代码只能通过工具调用 harness 提供的接口。这样做的最大好处是安全Agent 生成的代码不会直接接触宿主系统只在隔离的编译沙箱里跑。这个分离也方便并发。我一开始让 Agent 和 harness 跑在同一进程结果 Agent 等待编译时整个进程被阻塞并发完全做不起来。后来改成 harness 以独立子进程运行Agent 在等待实验结果时可以继续生成下一批候选实验吞吐量提升了接近 3 倍。所以如果你在搭类似系统第一件事就是把 harness 进程化、接口化。2.2 Agent 的整体架构四层循环我采用的是四层循环结构而不是一个单层的“LLM 生成-评估”循环调度层维护一个实验队列控制并发 GPU 任务数量避免多个实验同时占用同一个 GPU 导致测量互相干扰。生成层由 LLM 根据当前最优 kernel、历史实验结果、瓶颈分析报告生成多样化候选。评估层调用 harness 对候选进行编译、运行、计时返回结构化指标时延、occupancy、寄存器数、带宽利用率。记忆层存储所有实验记录包括 kernel 源码、参数、性能指标、失败原因供生成层做上下文参考。每一轮迭代Agent 会从记忆层取出最近 N 条记录连同当前最优 kernel 和一段由工具生成的瓶颈摘要一起交给生成层。LLM 输出的是完整的 kernel 源码或一组参数修改建议。生成层的 prompt 里明确要求“只输出 CUDA C 代码不要解释”这样能减少 token 消耗和解析失败率。2.3 为什么选择“LLM 贝叶斯优化”的组合纯 LLM 生成的问题在于生成的随机性太高可能连续十次都在做无意义的改动。纯贝叶斯优化又难以处理结构层面的变化只能搜索连续参数。所以我实际用的是混合策略连续参数blockDim、gridDim、unroll factor、向量化宽度由贝叶斯优化负责采用期望改进准则选点。结构变化tiling 方式、循环顺序、共享内存使用策略由 LLM 负责但会给它提供当前 kernel 的关键性能指标比如访存带宽是否饱和、occupancy 是否过低、是否存在 bank conflict。这个组合刚好互补。贝叶斯优化擅长在局部空间做精细调整LLM 擅长跳出局部陷阱。整个 24 小时里性能的提升有一半来自结构变化Agent 从最简单的 naive 实现先是改成了 vectorized load之后又引入 shared memory tiling最后把 grid-stride loop 改成固定次数循环。这些改动都是 LLM 依据瓶颈摘要自主做出的。2.4 并发与任务调度Agent 怎么扛并发并发是这次实验能 24 小时跑完的关键。单个 GPU 同时只能跑一个耗时测量任务但编译和代码分析不占 GPU。我在一张 A100 上跑又把编译放到 CPU 上GPU 只负责计时。调度层维护一个简单的 token bucket最多 2 个 GPU 任务并行其余候选排队编译任务最多 4 个并发防止 CPU 过载导致编译时间成为瓶颈。Agent 的“扛并发”能力不是靠异步框架而是靠任务队列解耦。我用了最朴素的queue.Queue加多个 worker 线程Agent 生成候选后丢进队列harness worker 消费队列执行编译和评测。实测下来整个 pipeline 的吞吐大约是平均每 40 秒完成一个候选24 小时大约跑了 2100 个实验。没有这个并发设计单线程串行至少需要 4 倍时间也就没法在 24 小时内冲榜。3. 核心实现细节搜索空间、奖励函数与工具链3.1 搜索空间定义哪些参数值得 Agent 去动我最初把搜索空间定义得特别大包括 gridDim、blockDim、动态共享内存、寄存器上限、指令级优化等十多个维度。结果 Agent 花了大量时间在低维空间里瞎试。后来我重新整理了搜索空间分成“必须搜索”和“由启发式决定”两类让 Agent 的注意力聚焦在真正影响性能的参数上。下面是我最终采用的搜索空间表格参数类型范围/取值影响blockDim.x连续整数32~1024步长 32决定线程束数量与占用率gridDim.x连续整数由输入大小与 blockDim 推导影响负载均衡向量化宽度离散枚举1/2/4float1/2/4显著影响访存带宽利用率循环展开因子连续整数1~8减少循环开销但增加寄存器压力共享内存 tiling 尺寸枚举8/16/32/64提升数据复用避免 bank conflict是否使用 padding布尔true/false影响 bank conflict 概率是否使用__restrict__布尔true/false影响编译器自动向量化编译器优化标志枚举O2/O3 快运算影响数值精度与速度单看这些参数人工搜索很容易陷入局部最优。比如 blockDim 设为 256 可能是多数情况下的好选择但配合向量化宽度 4 和循环展开 2 之后最优 blockDim 可能变成 192。Agent 的优势就在这里它能通过历史实验数据发现参数之间的交互效应而不是孤立地调每个参数。3.2 奖励函数设计不能只看运行时间如果奖励函数只是“运行时间越小越好”Agent 很快就会作弊它会倾向于生成一个只对特定输入 size 有效的 kernel而榜单用的是多个输入 size 的平均值。所以我把奖励函数设计为score -mean(runtime_over_all_inputs) - 0.1 * std(runtime_over_all_inputs) - 0.05 * compilation_fail_count其中 compilation fail count 是一个惩罚项如果 Agent 生成的 kernel 编译失败它会记录到一个会话级计数器超过 3 次后暂停生成该方向的变体。std 惩罚项是为了防止 Agent 牺牲某几个输入 size 来换平均值。测试集包含 8 组不同规模的数据从 128x128 到 4096x4096 都有。评估流程也要严谨。运行时测量不是只跑一次而是连续跑 10 次丢弃前 2 次作为 warmup再用剩余 8 次的中位数作为该候选的最终成绩。中位数比平均值抗噪因为 CUDA kernel 偶尔会受到时钟降频或后台任务影响。3.3 工具链Nsight Compute 与 nvcc 的正确用法Agent 需要知道当前 kernel 的瓶颈在哪才能做出有效的结构修改。纯靠运行时间是做不到的所以我在 harness 里集成了 Nsight Computencu进行采样分析。每次评估时不会频繁调用 ncu因为它本身有较大开销。我的策略是每 20 个实验对当前最优 kernel 做一次完整 profiling提取几个关键指标achieved occupancy实测占用率memory throughput存储吞吐compute throughput计算吞吐是否存在 uncoalesced memory access合并访存情况shared memory bank conflict 比例这些指标以文本形式存入记忆层。LLM 看到“memory throughput 达到 90% 但 compute throughput 只有 40%”这种信息后通常就能意识到应该减少冗余访存、提升数据复用而不是傻傻地调 blockDim。实测下来瓶颈摘要的质量直接影响生成质量这也是为什么工具链集成在 Agent 系统里如此重要。3.4 Agent 与工具交互的代码骨架下面是我在生成层和 harness 之间定义的接口骨架供参考# harness.py class KernelHarness: def compile(self, source: str, flags: str) - CompileResult: 调用 nvcc 编译返回是否成功和错误信息 def benchmark(self, kernel_path: str) - BenchmarkResult: 加载 kernel对每个输入 size 计时返回平均和中位时延 def profile(self, kernel_path: str) - ProfileResult: 调用 ncu 返回瓶颈指标 def submit(self, kernel_path: str) - SubmitResult: 将 kernel 提交到榜单队列 # agent.py def agent_loop(): memory load_memory() current_best load_best() while not timeout: context build_context(memory, current_best) suggestion llm_generate(context) if suggestion.is_parameter_change(): next_params bayesian_optimizer.suggest(memory) candidate apply_params(current_best, next_params) else: candidate suggestion.source_code harness.compile(candidate) result harness.benchmark(candidate) memory.add(candidate, result) update_best(current_best, result)这段代码虽然简化了但核心思路不变Agent 总是在“记忆里发生了什么”和“当前最优是什么”的基础上做决策。没有记忆的 Agent 只会随机乱试24 小时可能只在原地打转。4. 24 小时实战过程记录从 57 名到 15 名的寻优轨迹4.1 前 1 小时准备 baseline 和榜单规则刚开始我没有急着让 Agent 开跑而是先用一个手动编写的朴素 kernel 作为 baseline在测试集上测出平均时延。当时这个 naive kernel 的排名是 57 名离前 50 还有一段距离。同时我把榜单规则逐条读了一遍允许用什么编译器版本、是否需要提交完整项目、对cudaMalloc和cudaMemcpy的计时是否有豁免。这些规则直接影响 Agent 的优化方向比如榜单明确说明计时只包含 kernel 执行时间不包含 host 侧内存拷贝那 Agent 就不需要浪费时间优化拷贝部分。前一个小时的另一项重要工作是确认 harness 的测量结果和榜单官方结果没有系统性偏差。我提交了一个 baseline 版本等官方更新后对比差异在 2% 以内可接受。如果这一步不做后面所有实验结果都可能是自嗨。4.2 第一阶段1-6 小时盲目搜索期前五个小时 Agent 的输出比较混乱它尝试了各种参数组合但平均性能提升很慢。从记忆数据看这一阶段主要做了三件事调 blockDim、调向量化宽度、开关__restrict__。这些改动只能带来 5%-10% 的提升但排名仅仅从 57 到 49。此时的问题在于 Agent 没有对瓶颈做深入分析ncu 的 profiling 还没有触发记忆里全是参数和结果缺少结构信息。我在第 4 小时检查了一次中间结果发现 Agent 已经开始重复尝试相似的配置说明贝叶斯优化器过度信任了早期的局部最优。于是我给优化器加了一个“exploration bonus”在采集函数里提升不确定区域的权重。之后 Agent 开始尝试更大的 tiling 尺寸这个改变直接让时延下降了约 20%。4.3 第二阶段6-18 小时瓶颈分析和结构性优化第 8 小时我第一次让 Agent 对当时的 best kernel 跑了一次完整的 Nsight Compute profile。结果非常清晰memory throughput 只有 35%而 compute throughput 是 65%。这意味着 kernel 被访存瓶颈卡住了但并不是因为带宽不够而是因为访存模式混乱。Agent 在后续的 prompt 中收到了这个信息开始尝试 shared memory tiling。这个阶段大约生成了 800 个候选其中最有效的一次改动是把二维 tile 从 16x16 改成 32x32并加上 padding 避免 bank conflict。这个优化让 occupancy 从 40% 提升到 65%时延又降了 25%。排名进入前 30。之后 Agent 又根据 profiling 发现寄存器溢出主动把循环展开因子从 8 降到 4性能反而提升。这说明 Agent 确实理解了“寄存器压力”和“占用率”之间的权衡。4.4 第三阶段18-24 小时冲刺前的精细调整接近第 20 小时最佳成绩已经能排在第 18 位左右。此刻榜单前面的选手差距极小想再往前需要非常精细的优化。Agent 的贝叶斯优化器开始集中搜索 blockDim 和向量化宽度的交互区域连续跑了几百个只差几个线程数的配置终于找到一个组合让某些输入 size 的时延再降 3%。加上用了--use_fast_math标志允许以极小的精度损失换取更高性能最终平均成绩排到了第 15。这一段经历让我印象很深最后 4 小时的提升量其实很小但名次上升了 3 位因为榜单头部密集。如果只按“性能绝对提升”来看Agent 早期做得更成功但按“冲榜目标”来看这个阶段才是决定胜负的关键。给 Agent 设定排名作为动态目标而不是单纯优化时延是这次设计里比较明智的决定。下表是整个 24 小时各个关键里程碑的汇总时间点策略重点最优时延相对 baseline榜单排名0hbaseline1.00x575h参数搜索 向量化0.88x499hshared memory tiling0.74x2815h展开因子 padding 调整0.68x1920h寄存器压力优化0.66x1824h微调和 nvcc 标志0.64x15这个轨迹说明Agent 寻优的收益不是线性的而是呈现出两到三次大的跳跃每次跳跃都对应一次结构级修改。参数级优化只能带来平滑的小幅提升。5. 常见问题与避坑实录5.1 编译环境问题kernel header files not in any这个坑几乎每个跑过 CUDA 项目的人都会遇到。Agent 生成的第一批候选里有两个编译失败错误信息是make: common.mk:82: *** kernel header files not in any。原因是环境变量CUDA_PATH没有正确指向 CUDA 安装根目录或者 Makefile 里包含了多余的头文件路径。解决办法是在 harness 启动时显式设置export CUDA_PATH/usr/local/cuda-12.3 export PATH$CUDA_PATH/bin:$PATH export LD_LIBRARY_PATH$CUDA_PATH/lib64:$LD_LIBRARY_PATH还有几个候选是因为 nvcc 版本和驱动不匹配导致编译时找不到libcudart.so。如果没有一个统一的编译环境Agent 很容易在这种环境问题上浪费大量时间。我的经验是harness 必须用容器化环境固定 CUDA 工具链比如用 Docker 封装 nvcc 和驱动而不是寄希望于宿主机环境永远一致。5.2 测量噪声问题CUDA kernel 的运行时间测量比想象中更容易受噪声干扰。最开始 Agent 报告的同一个候选在不同时间跑出来的成绩波动能达到 15%这会让优化器完全失效。解决办法是增加重复次数和采用中位数同时确保没有其他进程占用 GPU。有一个细节很容易被忽略cudaEventRecord只测 GPU 时间不包含 kernel launch 的 CPU 开销如果你的 kernel 执行时间很短低于 5 微秒launch 开销可能成为主导。榜单的输入 size 较大kernel 执行时间在毫秒级所以这个问题不严重。但如果你优化短 kernel一定要考虑把多个 kernel 调用包进同一个 CUDA graph 来减小启动延迟。5.3 Agent 生成无效代码LLM 生成代码的失败率不低。在我的实验里约 12% 的生成结果存在语法错误、数组越界声明、或者调用不存在的 API。一个关键技巧是让 Agent 在生成代码时附带一段“我预期这个改动会带来什么影响”的说明虽然我明确告诉它“不要输出解释”但在单独的 metadata 字段里写理由能显著提高代码有效性。因为这些理由让模型更聚焦于逻辑一致性。另外harness 编译报错信息也会回传给 Agent。Agent 看到报错后可以自行修复比如“错误argument of type float4* is incompatible with parameter of type float*”它会意识到需要修改指针类型。如果连续多次修复失败则丢弃候选避免无限循环。5.4 榜单规则里的隐性限制冲榜之前一定要细读规则。有些榜单会禁止使用__launch_bounds__或者限制最大寄存器数有些则会要求 kernel 对所有输入 size 都保持正确性不允许使用任何“作弊”行为比如根据输入 size 分支到不同的优化路径。Agent 是黑盒它不知道这些规则所以需要人工在 prompt 里写清楚约束还要在 harness 里做规则校验比如检查 kernel 是否包含switch或将输入硬编码。我就在这里吃过亏Agent 生成了一个针对 4096x4096 矩阵高度优化的分支其他 size 走的路径性能一般整体平均反而下降了。后来我在奖励函数里增加了“每个输入的时延偏离最优 ratio 不能超过 30%”的约束才避免了这种投机行为。5.5 避坑清单速查坑点症状对策CUDA 环境变量缺失编译报 kernel header files not in any容器化固定环境显式设置 CUDA_PATH测量噪声大同一 kernel 成绩波动 15%重复 10 次取中位数GPU 被其他任务抢占冲刺阶段成绩突降调度层独占 GPU禁止并发占卡Agent 生成越界代码编译报 out-of-bounds反馈编译错误让 Agent 自修复对特定 size 过拟合平均分不高但个别 size 超好增加每个 size 的偏差惩罚寄存器溢出时延不降反升降低 unroll factor检查 ncu 寄存器指标这些坑看起来都很基础但在 Agent 自动化场景下任何一个都会成倍消耗时间。人工调优时遇到环境问题可以顺手解决但 Agent 不会像人一样灵活“绕过”环境问题它只会不断编译失败再尝试这非常浪费时间。所以为 Agent 打造一个稳定、干净的 harness 是投入产出比最高的一步。6. 事后复盘这套方法论能迁移到哪里这次实验结束之后我一直在想Agent 自动寻优的本质到底改变了什么。它并没有替代人的分析能力而是把“实验-反馈-调整”这个循环压缩到了分钟级。人看一份 Nsight Compute 报告需要几分钟Agent 读取文本后可以在几十秒内生成下一版 kernel。24 小时几千次实验这已经超过了绝大多数人类工程师能承受的调优工作量。和传统 autotuner 比起来这套方案的优势在于结构搜索能力。TVM Ansor 也做 kernel 搜索但它的搜索空间由一系列预定义的模板组成模板之间的切换逻辑是固定的。而 LLM Agent 可以跨模板操作甚至可以提出我们没想到过的 tiling 变体。代价是它不够稳定需要记忆层和贝叶斯优化器兜底。但 Agent 也有明显的边界。它在单个 kernel 的调优上很有效但如果任务变成“端到端模型性能优化”搜索空间会爆炸Agent 的上下文也会被大量无关信息淹没。我之前尝试让它优化一个包含多个 kernel 的流水线效果远不如单 kernel因为 LLM 很难同时追踪多个 kernel 之间的依赖关系。针对这种场景可能需要更复杂的规划模块比如把一个 pipeline 级优化问题拆成多个子问题再让多个 Agent 分头解决。就我自己来说这次经历最大的收获不是第 15 名而是验证了一套“LLM 负责结构探索、优化器负责参数精调、harness 负责稳定执行”的通用框架。之后我又用同样的架构去调一个自定义的反向传播算子虽然没有榜单可冲但在一个内部测试集上拿到了 1.7 倍的加速。所以这个方法并不仅限于冲榜它对任何有明确指标反馈的优化问题都适用。如果你也想尝试我建议从一个小目标开始先挑一个你手头已知最优配置的 kernel让 Agent 自动跑 8 小时看看它能不能找到你已知的最优解。如果能说明 harness 和 prompt 设计基本合格再去挑战未知问题。顺便说一句给 Agent 提供 profiling 反馈时记得把原始输出的关键数值抽取成结构化文本不要直接把 Nsight Compute 的完整终端输出丢给模型。上下文一长模型很容易忽略关键数字实验效果会大打折扣。最后一个小技巧在评价 Agent 表现时别只看最终排名还要看它能否复现自己的成功路径。如果 Agent 只是偶然冲进前 15下一次就崩溃那这个系统就没有任何实用价值。我在设计时专门加了一个“回放”步骤24 小时结束后让 Agent 基于记忆重新生成当时的最佳 kernel看能否得到相同的时延。只有这一步通过我才会相信那 24 小时不只是运气。
