CUDA Kernel和AI Agent这两个词放在一起2026年之前很多人会觉得它们是两条平行线一个在最底层做GPU并行计算一个在最上层做智能决策。但2026年的技术版图里这两条线开始加速交汇——Kernel不再只是被手写和被调优的对象Agent也不再只满足于调用现成的推理API。这篇研究综述想把“CUDA Kernel Agent”这个交叉领域完整拆开它到底在解决什么问题、技术栈长什么样、最小实验怎么搭、未来会往哪走。适合两类人读一类是做算子库、推理引擎优化的工程师想搞清楚Agent能怎么帮自己干活另一类是搞AI Agent框架的开发者想弄明白为什么底层Kernel的选择会直接影响Agent的体验和成本。我先说个总体感受。我在查相关资料的时候发现大家关心的东西特别杂有人在问WSL2里怎么装CUDA、多版本CUDA怎么共存有人在问llama.cpp为什么报kernel不兼容还有人在追Agent开发框架、Agent记忆、Agent安全甚至“命令行coding agent”这类新产品也在讨论范围内。这些话题表面上散实际上都指向同一个问题Agent要真正在高性能GPU环境里跑起来、跑得快、跑得稳Kernel层是绕不过去的。2026年的“CUDA Kernel Agent研究”本质上是在研究“底层并行计算与上层智能系统如何双向奔赴”。这篇综述我就按这个逻辑往下拆。1. 为什么2026年大家都在看CUDA Kernel Agent1.1 Kernel开发的范式迁移从手写优化到系统自治传统上CUDA Kernel开发是一条很“硬”的路。你得懂线程块怎么切、共享内存怎么排布、bank conflict怎么避开、occupancy怎么拉满还得知道每一代GPU架构的脾气——Turing、Ampere、Hopper、Ada、Blackwell的寄存器文件大小、Tensor Core指令差异、L2分片方式全都不一样。我见过很多性能团队一个GEMM Kernel能调一个月最后靠的是一摞实验记录和老师傅手感。但2024到2026年这波大模型浪潮把算子规模推到了一个手写跟不上的量级。FlashAttention、PagedAttention、MoE分组GEMM、结构化稀疏、长序列分块新型算子的出现速度远快于手工优化速度。工业界已经大面积转向cutlass、Triton、自动调优工具和“代码生成编译反馈”的流水线。这个阶段Kernel开发的瓶颈不再是“谁会写CUDA”而是“谁能高效地搜索Kernel配置空间、快速试错、自动迭代”。这正好是Agent擅长干的事。1.2 Agent并不只是跑在Kernel上的“应用层”很多人一听到AI Agent第一反应是“这跟GPU底层的Kernel有什么关系”。其实关系大得离谱。Agent的推理底座就是LLM服务而LLM服务的前向计算全是一堆CUDA Kernel在撑着QKV投影是GEMM注意力分数是FlashAttention KernelMoE路由是分组GEMMKV Cache管理是PagedAttention的Kernel。实测下来Agent首token延迟和推理吞吐几乎由这些Kernel的优化水平决定。再往工具层看Agent要调用工具比如向量检索、代码执行、图像处理、文档解析绝大多数计算密集步骤最终还是会落到预编译好的CUDA Kernel或加速库上。所以我说一句不太好听但很实在的话2026年谈Agent工程化如果不碰Kernel层你其实是在隔靴搔痒。Kernel决定Agent的效率和成本Agent反过来也应该有能力和意识去选择、生成、调优Kernel这就形成了双向依赖。1.3 这个交叉领域到底要回答哪三个问题我在整理这个领域时慢慢发现所有研究和工程探索其实都在回答三个具体问题。把这三个问题想清楚综述的骨架就立住了。谁来定调度哪个Kernel现在大多是手工指定或者靠算子库的启发式规则。未来Agent能不能根据上下文、输入形状、硬件状态动态决策在normal GEMM、Triton融合版本、Tensor Core特化版本之间做选择。谁来生成并验证新Kernel传统编译器能做简单生成但面对新型算子Agent能不能读懂语义、生成候选实现、编译反馈、循环优化形成一个自动迭代闭环。谁来解释Kernel执行中的异常Profiling数据、错误日志、SASS反汇编这些信号复杂又冗长Agent能不能抓住关键信息做自我诊断和自我修复。这三个问题分别对应了Kernel侧的“决策维度”、“生成维度”和“可观测维度”。后面的章节基本都围绕它们展开这也是我认为2026年CUDA Kernel Agent研究最核心的贡献点。2. Kernel侧的Agent化自动调优、生成式编程与自我诊断2.1 自动调优的演进从人肉搜索到Agent搜索先看自动调优这条路。早期阶段Kernel调优就是“人肉搜索”改一个blockDim跑一次benchmark记录耗时再改。后来有了参数化模板加搜索算法比如TVM的AutoTVM、Triton的autotune它们能在小规模参数空间里跑随机搜索或贝叶斯优化。这个思路有效但搜索空间一旦变大就露怯——为什么因为Kernel调优的参数不只是block大小还有vectorization宽度、swizzle样式、流水线阶段数、double buffer开关、指令调度方式甚至可以细分到“要不要用cp.async”。Agent进来之后整个搜索范式变了。它不再机械地枚举参数而是像经验丰富的性能工程师一样读历史实验结果然后做出推理判断。我在实验里试过一个很朴素的场景让模型看最近10组tile size和对应耗时再决定下一组候选参数。有意思的是模型会自己总结规律比如“在M4096时用大tile更好因为L2命中率上来了”这种抽象能力是传统搜索算法不具备的。2026年这个方向已经成熟到可以产品化了Agent作为调优循环的决策中枢底层保持模板自动编译和benchmark验证不变但搜索策略从“数值优化”变成了“经验推理”。2.2 生成式Kernel编程把编译反馈变成Agent的“强化学习信号”比自动调参更激进的做法是让Agent直接生成Kernel代码。这里的核心难点不是“模型会不会写CUDA”而是“如何让模型在真实编译器的反馈中迭代”。我把这套机制叫“编译反馈闭环”它的工作方式是这样的Agent生成一段Triton或CUDA C代码交给编译器编译如果编译失败就把错误信息返回给AgentAgent根据错误修改如果编译成功就运行基准测试把性能数据返回Agent继续优化下一轮。实际操作中有一个很关键的工程问题nvcc的报错信息又长又拗口直接塞给Agent会浪费大量上下文窗口。我的做法是先做一层“错误摘要器”把冗长日志压缩成“第42行类型不匹配第43行数组越界”这种结构化描述。这个前置处理步骤特别重要我见过好多项目在这个环节翻车最后模型完全不知道错在哪。另一个经验是做生成式Kernel编程优先从Triton入手而不是直接挑战CUDA C。Triton的编程模型更接近“算子语义描述”索引、布局、同步这些脏活全丢给编译器Agent要控制的参数更少、反馈更快。这就像让一个新司机先开自动挡先把路跑熟再回来学手动挡。等Agent对Triton的生成稳定了再试着让它产出CUDA C和PTX级别的代码那个阶段才是真正的硬核优化。2.3 把Profiling数据变成Agent能读懂的“信号”Agent要调优Kernel光看一个耗时是不行的你得让它知道慢在哪里。ncu和nsys能给出海量指标但原始输出对Agent不友好甚至对人也不友好。我在自己的最小实验里做了一层极简的Profile Summary每次跑完Kernel用ncu --metrics抓10个关键指标比如achieved occupancy、memory throughput、L2 hit rate、SM busy rate然后整理成一段JSON摘要喂给Agent。这套做法效果很明显。有一次Kernel延迟偏高Agent读了摘要发现memory throughput在90%以上但SM busy rate只有30%立刻判断是访存瓶颈下一轮就开始改数据排布而不是死磕计算指令。这就是可观测性映射到Agent决策的过程把Profiler指标和性能根因之间建立关联Agent才能做有方向的搜索而不是盲猜。很多团队忽略了这个环节觉得直接把ncu原始文本丢给模型就行实践下来上下文消耗大、效果差纯属浪费Token。说到环境这块很多新手一上来就栽在工具链上。有用户问“CUDA Samples找不到”其实装了CUDA Toolkit之后Samples往往在/usr/local/cuda/samples部分发行版需要单独git clone cuda-samples再用make编译还有用户问“怎么确认CUDA是否真的装好了”最直观的就是nvidia-smi看驱动nvcc --version看编译器。这些看似琐碎的排查恰恰是Kernel Agent落地最常被卡住的环节后面第4章我会专门把环境方案摊开讲。3. Agent侧的Kernel化工具调用、推理底座与异构执行3.1 工具即内核Agent的每一次“动手”都在触达Kernel换到Agent这一侧来看。当我们说“Agent会调用工具”底层到底发生了什么我拆过一个文档总结类的Agent它的工作链路是先调用Embedding模型给文档做向量化——这一步是GPU上的GEMM Kernel再调用向量检索服务找相关片段——这一步涉及索引扫描和相似度计算又落在Kernel上最后把检索结果拼给LLM生成摘要——还是Kernel。可以说Agent的很多工具调用本质是一系列预编译CUDA Kernel的包装。所以2026年有一个很明显的产品化方向把常用计算能力封装成“工具包”或“技能包”底层挂载预编译好的CUDA实现Agent层只做选择与编排。比如社区里在做的Agent Skills、第三方Agent工作台你会发现它们提供的“技能”越偏底层效果越好。原因很简单预编译的加速库经过了充分优化比Agent即时生成的代码稳定得多。我的判断是短期内Agent不至于替代cuBLAS、cuDNN这些库而是会把它们变成自己工具箱里的“标准零件”。3.2 推理引擎里的Kernel实际上锁死了Agent的“手感”Agent的真实使用体验很大程度上由推理引擎的Kernel优化决定。vLLM、TensorRT-LLM这些框架里PagedAttention的Kernel决定了长上下文下KV Cache的吞吐MoE推理的Grouped GEMM Kernel决定了路由计算的延迟连续批处理的调度策略又影响着整个服务的排队时间。Agent要“想得久”因为现在很多Agent会先吐出大量思维链再给结论这意味着生成Token数暴涨对推理Kernel的吞吐压力更大。这里我特别想提一个大家经常踩的坑很多人用llama.cpp跑本地模型时遇到“kernel不兼容”的报错一脸懵。其实这背后就是Kernel与硬件算力不匹配的问题。llama.cpp对CUDA架构版本非常敏感编译时必须指定目标架构的compute capability比如你的显卡是Ada架构sm_89如果编译产物只包含sm_70的SASS运行时就会直接报错。检查方法也简单nvidia-smi看驱动版本deviceQuery看算力然后编译时明确指定archcompute_89,codesm_89。这个例子特别能说明问题哪怕你写的是Python调用llama.cpp底层Kernel的匹配问题照样会一票否决。Agent工具链的稳定性从来不是只看API层还得看Kernel层。3.3 从单卡到集群多工具组合与宏内核调度Agent的复杂任务往往会触发一组并发计算这时候单个Kernel的调优就不够了得上升到“宏内核调度”的层面。CUDA Graph就是一个典型工具它可以把多个Kernel捕获成一个图一次性提交给GPU执行极大降低Kernel Launch的开销。大模型推理框架已经默认这么干了把prefill阶段的多个算子捕获成一个Graph阶段内所有Kernel在GPU上连续执行省掉几千次CPU-GPU同步。对Agent系统来说CUDA Graph带来的启发是工具的编排不能只看逻辑层还得看物理执行层。如果Agent在一轮思考里需要调用检索、重排、摘要三个工具理想情况是这三个工具对应的Kernel能在一个CUDA Graph里批量提交而不是来回同步等结果。我在实践里试过把“向量检索前缀重排”两个工具合并到一个Graph里执行整体延迟下降了接近30%效果非常直观。再往上走还有集群层面的调度GPU时间切片、MPS多进程共享、多卡并行这些都会成为Agent服务化平台必须考虑的资源底座。4. 实操与上手搭建一个最小CUDA Kernel Agent实验4.1 先把环境铺好WSL2、Docker与CUDA多版本共存聊了这么多得说点能直接用的东西。我建议想研究这个方向的人环境直接用“WSL2 Docker NVIDIA Container Toolkit”这套组合或者原生Linux配conda。最省心的路线是Windows上装好WSL2在WSL2里装好最新NVIDIA驱动对应的CUDA Toolkit然后所有Agent实验跑在Docker容器里隔离做得干净想换CUDA版本不用折腾宿主机。关于多版本CUDA我要多说两句。很多人纠结“装哪个CUDA版本”其实驱动、CUDA Runtime、CUDA Toolkit是三层东西驱动是底层的向下兼容Runtime是程序实际链接的库Toolkit是你用来编译的那套开发工具。一台机器上完全可以共存多个CUDA Toolkit版本只要用软链接/usr/local/cuda切换或者在编译时通过-I和-L指定不同路径就行。第一步检查驱动和算力nvidia-smi nvcc --version /usr/local/cuda/bin/deviceQuery第二步确认Docker能调用GPUdocker info | grep -i runtime docker run --gpus all --shm-size8g -it nvcr.io/nvidia/cuda:12.4.1-runtime-ubuntu22.04 bash第三步在容器内安装Python依赖pip install torch triton容器跑起来后deviceQuery能正常打印你的显卡信息就说明环境通了。常见报错是nvidia-container-cli找不到驱动多半是宿主机没装好nvidia-container-toolkit或者--gpus all语法没被识别按这个顺序排查基本能解决。4.2 用什么模型当Agent大脑最顺手做Kernel Agent实验选模型是个性价比问题。我的建议是一开始别纠结本地部署直接用能稳定输出结构化JSON的商用API来写Agent逻辑把精力集中在“Kernel生成与优化循环”上而不是先被推理框架的部署细节淹没。等流程跑通、效果验证了再考虑用开源模型加本地推理来做生产力部署。这一步很关键因为Kernel调优本身就有很强的“试探—反馈—修正”特性模型输出格式不稳定的话下游解析和编译环节会连环出错。Agent框架方面你用LangChain或AutoGen都行但我个人其实更推荐自己写一个几十行的LLM循环给模型一个System Prompt告诉它“你现在是CUDA Kernel调优工程师”给它一个工具函数函数签名是submit_kernel(config) - result它每次输出配置参数你解析JSON、编译运行、返回结果如此循环。别急着上重型框架先跑通一个最小闭环后面再谈记忆、规划、评估这些高级能力。这也是我复盘下来觉得新手最容易走偏的地方——一上来就搭Agent框架结果半天没碰Kernel。4.3 最小闭环让Agent生成并优化一个GEMM Kernel下面我把最小闭环的代码骨架贴出来这个实验我实际跑过非常说明问题。目标很简单让Agent生成一个矩阵乘法Kernel先验证正确性再根据性能反馈循环优化。我建议用Triton来写Kernel模板原因前面说过反馈快、参数少、Agent容易学。核心模板长这样import triton import triton.language as tl import torch triton.jit def matmul_kernel( a_ptr, b_ptr, c_ptr, M, N, K, stride_am, stride_ak, stride_bk, stride_bn, stride_cm, stride_cn, BLOCK_M: tl.constexpr, BLOCK_N: tl.constexpr, BLOCK_K: tl.constexpr, ): pid_m tl.program_id(0) pid_n tl.program_id(1) offs_m pid_m * BLOCK_M tl.arange(0, BLOCK_M) offs_n pid_n * BLOCK_N tl.arange(0, BLOCK_N) offs_k tl.arange(0, BLOCK_K) a_ptrs a_ptr offs_m[:, None] * stride_am offs_k[None, :] * stride_ak b_ptrs b_ptr offs_k[:, None] * stride_bk offs_n[None, :] * stride_bn acc tl.zeros((BLOCK_M, BLOCK_N), dtypetl.float32) for k in range(0, K, BLOCK_K): a tl.load(a_ptrs) b tl.load(b_ptrs) acc tl.dot(a, b) a_ptrs BLOCK_K * stride_ak b_ptrs BLOCK_K * stride_bk c_ptrs c_ptr offs_m[:, None] * stride_cm offs_n[None, :] * stride_cn tl.store(c_ptrs, acc)然后写一个循环让Agent来改参数configs [ {BLOCK_M: 64, BLOCK_N: 64, BLOCK_K: 32}, {BLOCK_M: 128, BLOCK_N: 128, BLOCK_K: 32}, ] history [] for cfg in configs: # 编译并运行kernel latency_ms, correctness run_benchmark(M4096, N4096, K4096, **cfg) history.append({**cfg, latency_ms: latency_ms, correctness: correctness}) # 把history序列化为文本让Agent输出下一组参数 next_cfg agent_suggest(history)这个循环是核心每一次Agent拿到的不只是“上一轮耗时”而是“一系列参数组合与效果的轨迹”它可以基于轨迹做出类比推理——比如发现BLOCK_M增大导致寄存器溢出就下调觉得BLOCK_K太小导致访存重复就上调。整个过程中Agent其实是在做一种“黑盒优化加经验启发”的混合策略这就是跟传统贝叶斯优化最大的区别。我实际跑下来的参数规律是这样的对于M、N、K都是4096的方形矩阵BLOCK_M128, BLOCK_N128, BLOCK_K32通常比小Tile要好因为更大的Tile能提升数据复用率但如果矩阵形状变成“瘦高型”比如M256、K16384大BLOCK_K反而先把K维切得更碎配合流水线才能压满带宽。这种“形状-参数匹配”的规律正是Agent最该学的东西。你可以在提示词里给它一两组“黄金示例”比如直接告诉它“当M远大于N时优先保留BLOCK_M比较大的配置”它会少走很多弯路。4.4 实战中碰到的典型坑与排查顺序这个最小实验看着简单真做起来坑不少。我把一年来踩过的坑整理成一张排查表按出现频率排序方便你照着查。现象根因解决办法程序跑起来报“no kernel image”编译产物不匹配GPU算力检查显卡Arch编译时指定archcompute_xx,codesm_xxAgent生成的代码反复编译失败nvcc报错信息太长模型抓不住重点加一层错误摘要器把日志压缩成结构化信息性能数据忽高忽低Agent被噪声忽悠Profile与Benchmark混在一起跑干扰严重单独用ncu做Profile正式Benchmark时关闭一切追踪Agent越改越离谱开始乱调参数提示词里缺少约束模型随机发散给Agent一个“参数边界”并在输出JSON里做schema校验容器内无法使用GPUNVIDIA Container Toolkit没装好确认docker info里有nvidia runtime重装toolkit还有一个安全相关的提醒让Agent直接生成并执行CUDA代码是有风险的。Kernel编译过程中会展开模板、执行宏、做异步操作这些都可能被滥用。我的做法是所有实验跑在Docker容器里用非root用户运行容器不挂载宿主机敏感目录网络默认关闭。Agent的生成物永远只是候选最终要经过人工和CI回归测试才能进生产。这个习惯一定要从一开始就养成等出现问题再补代价就大了。5. 挑战、边界与2026年趋势判断5.1 正确性验证生成Kernel最容易忽略的“最后一公里”做一个Kernel生成Agent最难的其实不是“生成”而是“验证”。GEMM这种简单算子正不正确可以拿CPU参考结果对齐但一旦涉及融合算子、原子操作、混合精度想靠reference验证就非常困难了。我自己做的时候会专门设计一套“运行时正确性网”对每个生成的Kernel自动注入断言把输出和参考结果比对按rtol和atol容差判断是否通过。这里有一个特别容易误导人的细节浮点运算是不可结合的同一个算法换个block大小累加顺序变了结果就会有细微差异。Agent必须理解“误差在容忍范围内不算错”否则它会陷入无穷无尽的“修正代码”循环把本来就正确的实现改得一团糟。所以我在提示词里会明确说性能优化时只要正确性指标在容差内就不允许为了安全性牺牲可感知的性能。5.2 安全边界动态生成Kernel是一把双刃剑谈到Agent化安全是绕不开的话题。我支持Agent参与Kernel生成与研究但必须给它划清边界。动态生成的代码直接上生产这在当前的验证体系下是不负责任的。更现实的风险是Agent不只是生成了代码它还能改编译参数、改环境变量、改临时目录这些操作轻则把机器搞乱重则影响同一台机器上的其他GPU任务。我建议的标准流程是Agent负责在“隔离沙箱”里自由探索生成候选Kernel和调优报告人或者传统编译工具链负责终审最终产物进入一个带版本控制的Kernel仓库再通过CI跑完整的正确性、性能回归。把Agent定位成“高效的研究员”而不是“自动上线的生产者”——至少在2026年的工程成熟度下这是最稳妥的边界。5.3 2026年最值得关注的关键趋势最后聊几个我认为接下来会加速的方向算是一个基于现状的展望。第一个是“Kernel即服务”。未来Agent平台会内置大量预编译的加速原语Agent做的事情更多是“选择与编排”而不是“从零生成”。这就像今天你不会自己写排序算法而是调标准库一样。第二个是“编译前端与后端分离”。Agent停留在高层算子语义描述层面具体到线程排布和指令选择由Triton这类编译器去处理。第三个是“可观测性作为第一公民”。Profiling数据会越来越结构化成为Agent学习和决策的核心信号源这也是我在第2章强调的东西。第四个趋势我特别有感触异构兼容会从“口号”变成“工程问题”。现在大家讨论AMD显卡能不能跑CUDA代码、其他加速硬件怎么迁移本质上都是在问“Kernel的抽象层能不能再往上抬一层”。一旦算子的语义描述和硬件解耦Agent的价值会更大——因为它面对的就不再是一套封闭指令集而是一个开放的能力描述空间。这个方向能走多远我还在观察但趋势已经摆在那里了。我个人的实际操作体会是把Profiling数据喂给Agent之后它确实会去调整Block大小但前几次改的方向几乎全是大Tile策略感觉它默认了“越大越好”。只有当你把寄存器溢出、共享内存占用这些反馈也送给它它才会慢慢学会平衡。这让我意识到Kernel Agent的瓶颈之一不是模型能力而是“我们能不能把底层反馈翻译成模型听得懂的语言”。如果你要做这方面的实验我建议从最小的GEMM闭环开始先把反馈链做扎实再谈扩展算子类型和Agent策略。这个领域最迷人的地方在于它逼着你同时理解硬件和智能系统的语言而这两套语言正在2026年第一次真正开始对话。