1. 从“手写CUDA”到“Agent自动调优”这个方向到底在解决什么问题如果你最近一年在关注GPU高性能计算这个圈子应该能明显感觉到一个变化以前大家聊的是怎么手写shared memory、怎么避免bank conflict、怎么用warp shuffle做reduce而现在越来越多的人在讨论——能不能让一个智能体Agent自己去写kernel、自己去调参、自己去跑benchmark然后迭代。这个转变不是空穴来风它背后有一个非常现实的痛点。我做了几年CUDA相关的优化工作最深的体会就是写一个能跑通的kernel不难写一个跑得快的kernel非常难写一个在各种shape、各种dtype、各种硬件上都跑得快的kernel那基本是在折磨自己。一个典型的GEMM优化从naive版本到接近cuBLAS性能中间要经历tiling、register blocking、double buffering、swizzling、vectorized load等一大堆步骤每一步都有大量参数要调——tile size选多少、thread block怎么组织、每个thread算多少个输出、要不要用cp.async、要不要上Tensor Core这些决策空间组合起来轻松上百万种。人工调优一个kernel熟练工程师也要花好几天甚至几周。所以当LLM展现出代码生成能力之后一个很自然的想法就冒出来了能不能把“写kernel-编译-跑分-根据结果改代码”这个循环交给一个Agent去自动完成这就是CUDA Kernel Agent这个方向的核心命题。它要解决的不是“能不能生成CUDA代码”这种初级问题而是“能不能在给定硬件和问题规模下自动搜索出一个高性能实现”这个硬骨头。这篇文章我想系统梳理一下这个方向目前的研究格局。我会从问题定义开始拆解几个主流技术路线的设计思路分析它们各自的优势和局限然后落到实操层面——如果你想自己搭一个kernel agent有哪些坑是必须提前知道的。适合的读者是有CUDA基础、对LLM代码生成感兴趣、想了解自动kernel优化前沿的工程师和研究者。不需要你是CUDA专家但至少得知道什么是occupancy、什么是memory coalescing不然读起来会比较吃力。2. 核心问题拆解Kernel Agent到底难在哪2.1 为什么kernel优化不能简单套用“LLM写代码”那套很多人第一反应是现在LLM写Python、写C都挺溜的写CUDA应该也差不多吧实际试过就知道差距非常大。普通应用层代码你写错了顶多是逻辑bug跑起来慢一点但能出结果。CUDA kernel不一样它有三个要命的特点。第一正确性门槛高且反馈稀疏。一个kernel写出来编译通过不代表能跑能跑不代表结果对结果对不代表没有race condition。而且很多错误是静默的——比如shared memory没同步在小数据量下可能碰巧结果正确数据量一大就出错。Agent拿到的是一个二元信号对/错中间没有任何梯度信息告诉它“你哪里快接近对了”。第二性能反馈噪声大。同一个kernel跑两次耗时可能差5%甚至更多受GPU温度、频率boost、其他进程干扰影响。Agent如果拿单次测量结果做决策很容易被噪声带偏以为某个改动有效其实只是波动。这就要求测量协议必须做多次取中位数、warmup、固定clock等处理。第三搜索空间是组合爆炸的。前面说的tile size、thread组织、指令选择这些维度每个都是离散的、相互耦合的。你把tile从32改成64可能原来最优的thread block配置就废了。这种耦合性意味着不能各维度独立搜索必须联合优化而联合搜索的代价是指数级的。2.2 Agent需要具备的四种能力把上面这些难点翻译成对Agent的能力要求我认为至少要具备四样东西。代码生成能力是基础但这里的要求比一般代码生成高——生成的代码必须符合CUDA的编程模型约束不能出现非法内存访问、不能超出shared memory容量、不能违反同步语义。这需要模型对CUDA的硬件模型有相当深的理解而不是只会套模板。性能评估能力是核心。Agent必须能可靠地测量kernel耗时并且知道怎么排除噪声。这听起来简单实际上很多早期工作就栽在这里——测量不准后面所有决策都是空中楼阁。搜索策略能力决定效率。暴力枚举所有配置在预算内根本跑不完Agent需要知道哪些参数值得试、哪些方向大概率没戏、什么时候该放弃当前路线换一个思路。这本质上是一个带昂贵评估的黑盒优化问题。错误诊断与修复能力决定鲁棒性。编译失败、运行崩溃、结果错误是家常便饭Agent得能从错误信息里定位问题而不是一失败就重新生成一个全新的可能还是错的版本。2.3 评估基准的缺失与建立这个方向早期一个很大的障碍是没有公认的benchmark。你说你的Agent能自动优化kernel优化到什么程度算好跟谁比在什么硬件上比后来逐渐形成了一些共识性的评估维度。一是正确性验证必须有一套覆盖边界情况的测试用例不能只测一个shape。二是性能对标通常以cuBLAS、cuDNN、Triton这些成熟库作为baseline看Agent生成的kernel能达到baseline的百分之多少。三是搜索成本也就是达到目标性能花了多少次编译和测量这个指标直接决定了方法是否实用。四是泛化性在训练时没见过的shape和dtype上表现如何。我个人的经验是看一篇kernel agent的论文先看它的评估协议是否严谨。如果只报了一个shape的最优结果没有多次运行的方差没有和强baseline的对比那这个结果的可信度要打问号。3. 主流技术路线深度对比3.1 路线一基于搜索的迭代优化这是最直观的一条路线也是目前工作最多的方向。核心思路是把kernel优化建模成一个搜索问题Agent在每一轮生成一个候选实现编译测量后根据结果决定下一步往哪个方向走。具体实现上又分几种。一种是参数化模板参数搜索就是预先写好一个kernel骨架把tile size、unroll factor这些做成可配置参数Agent只负责搜参数。这种做法的好处是搜索空间可控、生成的代码一定合法坏处是天花板受限于模板本身——模板没想到的优化Agent永远搜不出来。另一种是自由代码生成迭代精化Agent直接生成完整kernel代码根据测量反馈修改。这种天花板高但搜索空间巨大容易在无效方向上浪费预算。实践中常见的是两者结合先用模板快速找到一个不错的起点再在这个基础上做自由修改。实操心得纯自由生成的Agent如果不加约束很容易陷入“每次生成一个完全不同的实现”的困境导致历史经验无法积累。我试过在prompt里强制要求“基于上一版修改只改一个变量”收敛速度明显提升。3.2 路线二基于编译反馈的闭环这条路线强调把编译器和profiler的信息充分利用起来。CUDA编译器nvcc会给出register usage、shared memory usage、spill情况等信息profiler如nsight compute能给出occupancy、memory throughput、compute throughput、stall reason等细粒度指标。一个成熟的kernel agent不应该只看“跑了多少毫秒”而应该看“为什么是这个耗时”。比如profiler告诉你memory throughput已经到90%了那说明是memory bound继续优化计算指令没用得想办法减少访存或者提高访存效率。如果告诉你occupancy很低是因为register pressure太大那就得想办法减少register使用。把profiler指标喂给Agent让它基于这些诊断信息做决策比单纯给一个耗时数字要有效得多。这也是我认为目前最有前景的方向之一——让Agent学会像人类专家一样读profiler。3.3 路线三多Agent协作与分工单个Agent能力有限那就上多个。常见分工方式有一个Agent负责生成代码一个负责review代码找bug一个负责分析性能瓶颈一个负责决定搜索方向。这种多Agent架构的好处是每个Agent可以专注于自己擅长的部分prompt可以写得更聚焦。但也有明显的代价通信开销大、容易互相甩锅、决策链条长导致迭代慢。我见过一些实现多个Agent讨论半天最后生成的代码还不如单Agent直接生成的。所以多Agent不是银弹关键看分工是否合理、通信协议是否高效。3.4 三条路线的横向对比维度搜索迭代路线编译反馈闭环多Agent协作实现复杂度中中高高搜索效率中高中性能天花板高高高对profiler依赖低高中调试难度低中高适合场景快速原型深度优化复杂任务分解这张表是我根据实际项目经验总结的不一定适用于所有情况但大方向应该没错。新手建议从搜索迭代路线入手跑通了再考虑引入profiler反馈。4. 实操从零搭建一个Kernel Agent的关键环节4.1 环境准备与测量协议先说环境。你需要一块支持CUDA的GPU驱动和CUDA toolkit版本要匹配。我建议固定一个版本组合不要频繁升级因为不同版本的编译器优化行为可能不同会导致你的实验结果不可复现。测量协议是重中之重。我的做法是每个kernel先跑10次warmup再跑50次取中位数。warmup是为了让GPU频率稳定取中位数是为了抗噪声。同时用cudaEvent做计时而不是CPU端计时避免launch overhead干扰。如果条件允许锁定GPU clock能进一步降低方差。# 测量kernel耗时的典型流程伪代码 import torch def measure_kernel(fn, warmup10, iters50): # warmup for _ in range(warmup): fn() torch.cuda.synchronize() times [] start torch.cuda.Event(enable_timingTrue) end torch.cuda.Event(enable_timingTrue) for _ in range(iters): start.record() fn() end.record() torch.cuda.synchronize() times.append(start.elapsed_time(end)) times.sort() return times[len(times)//2] # 中位数注意千万不要用time.time()在CPU端计时CUDA kernel是异步执行的你测到的可能只是launch时间。这个坑我踩过当时测出来kernel只要0.01ms还高兴了半天后来发现根本没等kernel执行完。4.2 Prompt设计与代码生成约束Prompt设计直接决定生成质量。我的经验是prompt里必须包含这几样东西目标问题的数学定义比如矩阵乘法的M、N、K、硬件约束shared memory大小、register上限、性能目标对标什么baseline、以及明确的代码规范用什么风格、要不要加注释。更重要的是约束。必须明确告诉Agent不能使用超过X KB的shared memory、每个thread block的thread数必须是32的倍数、必须处理边界情况。这些约束如果不写清楚生成的代码大概率编译不过或者跑出错误结果。还有一个技巧是给示例。在prompt里放一两个正确且高效的kernel作为参考Agent生成的质量会明显提升。这相当于few-shot learning让模型知道“好的kernel长什么样”。4.3 搜索策略与预算分配搜索策略决定了Agent在有限预算内能走多远。我常用的策略是分层搜索先粗粒度搜大方向比如用不用Tensor Core、用不用shared memory确定方向后再细粒度调参数。预算分配上我的经验是20%预算用于探索不同的大方向60%用于在最有希望的方向上精调20%留作备用——因为有时候一个看起来没戏的方向深入下去反而有惊喜。还有一个反直觉的经验不要过早收敛。有些Agent实现看到某个配置表现好就一直在附近微调结果陷入局部最优。适当保留一些随机探索反而更容易找到全局最优。4.4 错误处理与自动修复错误处理是区分玩具和实用系统的分水岭。一个实用的kernel agent必须能处理三类错误编译错误、运行时错误、结果错误。编译错误最好处理因为编译器会给出具体的行号和错误类型直接把这些信息喂回给Agent让它修复就行。运行时错误比如非法内存访问稍微麻烦一点需要Agent理解CUDA的错误码含义。结果错误最难因为Agent不知道错在哪只能靠对比参考实现来定位。我的做法是维护一个错误模式库把常见的错误和对应的修复策略存起来。当Agent遇到类似错误时先查库查不到再让LLM自由发挥。这样能大幅提升修复效率。5. 常见问题与排查技巧实录5.1 性能不达标怎么排查Agent生成的kernel性能不达标是最常见的问题。排查思路应该是从粗到细。先看occupancy。如果occupancy很低比如低于25%那基本可以确定是register或shared memory用太多了导致同时活跃的warp太少无法掩盖延迟。这时候要么减少register使用要么调整block size。再看memory access pattern。用profiler看global memory的load/store效率如果远低于理论带宽说明访存没有coalesce或者有严重的bank conflict。这是最常见的性能杀手。最后看instruction mix。如果计算指令占比很低大量时间花在地址计算、分支判断上那说明代码本身写得不够紧凑需要优化指令级并行。5.2 结果正确但性能波动大这个问题通常是测量协议的问题不是kernel本身的问题。检查几点有没有做warmup、有没有固定clock、有没有其他进程占用GPU、测量次数够不够。如果这些都排除了还是波动大那可能是kernel本身对数据敏感——比如某些数据分布导致分支预测失败率高。这种情况需要让Agent生成对数据不敏感的版本比如用无分支实现替代有分支实现。5.3 Agent陷入重复生成相似代码这是搜索策略的问题。Agent在某个局部区域反复打转说明探索机制不够。解决办法是引入“新颖性奖励”——如果生成的代码和之前的历史记录太相似就降低它的评分强制Agent探索新方向。另一个办法是定期“重启”——每隔N轮清空历史让Agent从一个全新的起点开始。这听起来很浪费但实际能有效跳出局部最优。5.4 常见问题速查表问题现象可能原因排查方向解决思路编译失败语法错误/资源超限看编译器报错行喂回错误信息让Agent修复运行崩溃非法访存/同步错误用compute-sanitizer检查边界和同步点结果错误race condition/逻辑bug对比参考实现小数据量逐步排查性能差occupancy低/访存差看profiler指标针对性优化波动大测量噪声/数据敏感检查测量协议固定clock/多次取中位数重复生成搜索陷入局部最优看历史记录引入新颖性奖励/重启6. 这个方向后续可以怎么扩展聊完现状说说我对这个方向未来走向的一些判断。不一定对但可以作为参考。从单kernel到全模型。目前大部分工作聚焦在单个kernel的优化但实际应用中一个模型有几十上百个kernel它们之间有依赖关系。怎么在全局层面做优化——比如算子融合、内存复用——是更大的课题。这需要Agent理解计算图而不只是单个kernel。从离线搜索到在线自适应。现在的做法是离线搜好一个kernel然后部署。但实际负载是动态变化的输入shape可能随时变。能不能让Agent在运行时根据实际负载动态选择和调整kernel这需要极低的搜索开销目前还很难做到。从手工特征到学习到的搜索策略。现在搜索策略很多还是手工设计的启发式规则。用强化学习或者元学习来训练搜索策略让它从大量优化任务中学习“什么样的kernel该怎么优化”是一个很有潜力的方向。跨硬件泛化。在A100上搜出来的最优配置搬到H100上可能就不是最优了。怎么让Agent具备跨硬件泛化能力或者能快速适应新硬件是一个实际部署中必须解决的问题。我个人在实际操作中的体会是这个方向目前还处于“能做出demo但离产品化有距离”的阶段。最大的瓶颈不是LLM的代码生成能力而是搜索效率和评估可靠性。谁能把这两个问题解决好谁就能让kernel agent真正落地。如果你也想入这个坑建议先从一个小而具体的kernel入手——比如一个简单的element-wise操作或者reduction——把整个闭环跑通再逐步增加复杂度。别一上来就搞GEMM那个坑太深容易劝退。