我不太确定自己当时是不是真的能冲上那张榜单，只是隐约觉得，如果能用一个 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/4（float1/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 Compute（ncu）进行采样分析。每次评估时，不会频繁调用 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） | 榜单排名 |
|---|---|---|---|
| 0h | baseline | 1.00x | 57 |
| 5h | 参数搜索 + 向量化 | 0.88x | 49 |
| 9h | shared memory tiling | 0.74x | 28 |
| 15h | 展开因子 + padding 调整 | 0.68x | 19 |
| 20h | 寄存器压力优化 | 0.66x | 18 |
| 24h | 微调和 nvcc 标志 | 0.64x | 15 |
这个轨迹说明，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 小时不只是运气。
