1. 这个标题到底在说什么事先把标题拆开看。“AI到底能不能自己造AI”这句话背后其实藏着一个被反复争论了好几年的问题我们能不能让一个模型去设计、训练、优化另一个模型甚至让它去改自己的训练流程、改自己的推理代码、改自己跑在GPU上的算子。注意这里说的不是“AI帮人写几行代码”这种浅层辅助而是让AI真正参与到“造AI”这个链条里——从数据构造、训练策略、奖励设计一直到最底层的GPU Kernel实现。“别吵了有人做出来了”这句话指向的是近一两年一批很实在的工作用Agent框架把“模型研发”这件事拆成可执行的任务流让一个或多个Agent在受控环境里反复试错最终产出能跑通的训练脚本、能编译的Kernel、能提升指标的SFT/RLVR数据。热搜词里出现的Agent、RLVR、SFT、GPU Kernel基本就是这条链路上的四个关键节点。我先把这四个词用大白话过一遍不然后面没法聊。Agent在这里不是聊天机器人而是一个能调用工具、能读写文件、能执行命令、能根据反馈调整下一步动作的执行体。它和普通“对话式AI”最大的区别是它有循环有状态有目标会失败会重试。SFTSupervised Fine-Tuning监督微调。简单说就是拿一批“输入-输出”配对的数据让模型学会按某种格式、某种风格、某种推理路径来回答问题。它是让基座模型变成“能用”的第一步也是最耗数据质量的一步。RLVRReinforcement Learning with Verifiable Rewards可验证奖励的强化学习。它的核心是奖励不是人打的而是由程序自动判定的。比如代码题跑一遍测试用例过了就给奖励数学题对答案对了就给奖励。这样一来模型可以自己生成大量尝试靠“可验证”的信号来进化而不需要人类逐条标注。GPU Kernel就是跑在GPU上的计算核心函数。深度学习框架里那些矩阵乘、卷积、归一化、注意力最终都要落到Kernel上。Kernel写得好不好直接决定训练和推理快不快、显存省不省。而“在GPU上执行的全流程”这个问题问的就是从算子定义、编译、调度到实际执行这一整条链路。把这四个词串起来标题的意思就清楚了有人用Agent框架结合SFT和RLVR让AI去自动生成、优化GPU Kernel甚至进一步去参与模型训练流程的设计。这不是概念演示而是有可运行产物、有指标对比、有失败记录的工程实践。我写这篇东西不是要复述某篇论文而是想从一个实际动手的人的角度把这条链路拆开讲清楚它到底怎么跑起来的哪些环节最容易崩哪些参数不能乱设哪些“看起来很美”的做法其实根本落不了地。适合谁看适合已经写过Agent、跑过SFT、调过Kernel或者至少对其中一两块有实操经验的人。纯小白也能看但需要你对“模型训练不是点一下按钮”这件事有基本认知。2. 整体设计思路为什么是Agent加RLVR加Kernel这条线2.1 为什么不让AI直接写完整训练脚本很多人第一反应是既然要让AI造AI那直接让一个大模型输出一整套训练代码不就行了我试过这条路在真实场景里基本走不通。原因不是模型不够聪明而是训练脚本的正确性依赖太多隐式上下文数据路径、tokenizer版本、分布式配置、显存预算、混合精度策略、梯度累积步数、学习率调度任何一个环节错一点跑起来就是loss不降或者直接OOM。更关键的是一次性生成的代码没有反馈闭环。模型不知道它写的batch size会不会爆显存不知道它选的learning rate会不会导致梯度爆炸不知道它调用的Kernel在目标GPU架构上能不能编译通过。没有反馈就没有修正最后只能靠人肉逐行review那还不如自己写。所以真正可行的思路是把“造AI”拆成可验证的小任务让Agent在循环里试错。每个小任务都有明确的成功判据比如“这段Kernel能编译通过”、“这个训练脚本跑100步loss下降”、“这批SFT数据能让模型在验证集上提升2个点”。Agent不需要一次做对它只需要在多次尝试中逼近正确。2.2 Agent在这里的角色不是“写手”而是“实验员”我见过太多Agent项目把Agent当成一个更花哨的代码生成器。但在“AI造AI”这个场景里Agent的核心价值不是生成而是执行、观察、调整。具体来说一个合格的研发Agent需要具备这几个能力文件系统操作读写脚本、配置文件、日志。命令执行跑训练、跑编译、跑测试。结果解析从stdout、stderr、日志文件里提取关键指标和报错信息。状态记忆记住上一次尝试改了什么、结果如何避免重复踩同一个坑。预算控制知道自己的时间、显存、API调用次数有限不能无限试错。这五条里最容易被低估的是结果解析。很多Agent框架把大量精力花在prompt编排上却对“怎么从一堆日志里判断这次实验是成功还是失败”处理得很粗糙。结果就是Agent看到“loss: nan”还以为训练在正常进行继续往下跑浪费大量算力。2.3 RLVR为什么比人工奖励更适合这个场景在“AI造AI”的链路里奖励信号必须自动化否则整个循环跑不起来。RLVR的价值就在这里它把“好不好”变成一个可以程序判定的问题。以Kernel生成为例奖励可以这样设计判定项判定方式奖励权重编译通过调用编译器返回码为0必须满足否则直接0分数值正确与参考实现对比误差在阈值内高权重执行速度与基线Kernel对比耗时中权重显存占用统计峰值显存低权重代码可读性静态检查如行数、注释比例极低权重或不计这种奖励设计的好处是完全客观不需要人类标注也不依赖另一个模型的主观判断。坏处是容易过拟合到奖励函数本身Agent可能生成一个只对当前测试用例正确、换个shape就崩的Kernel。所以实际使用中测试用例集必须足够多样而且要定期轮换。2.4 SFT在这个链路里到底干什么SFT不是用来“教模型写Kernel”的那是RLVR的活。SFT在这里的作用是给Agent一个像样的起点。具体来说如果直接让一个通用模型去写CUDA Kernel它大概率会输出语法错误、API误用、线程块配置离谱的代码。但如果你先用一批高质量的“Kernel描述-正确实现”配对数据做SFT模型至少能学会Kernel函数的基本结构。线程索引、共享内存、同步点的常见写法。不同算子如softmax、layernorm、matmul的典型优化模式。有了这个起点RLVR阶段的搜索空间会小很多Agent不用从零开始摸索语法而是把精力放在性能优化和边界处理上。我个人的经验是SFT数据不在多而在“干净”。一千条经过人工验证的Kernel实现比十万条从网上爬的代码片段有用得多。因为SFT阶段一旦引入错误模式后面RLVR要花很大代价才能纠正。3. 核心细节解析从任务定义到Kernel落地3.1 任务定义怎么把“造AI”拆成Agent能接的活这是整个项目里最容易被忽视、但最决定成败的一步。任务拆得太粗Agent无从下手拆得太细Agent变成只会执行固定步骤的脚本失去搜索能力。我的做法是按“可验证产物”来拆。每一个子任务都必须有一个明确的、程序可判定的产出物。比如子任务A生成一个能编译通过的CUDA Kernel实现指定算子。子任务B让该Kernel在给定输入下数值正确。子任务C让该Kernel的执行时间低于基线。子任务D生成一批SFT数据使模型在验证集上指标提升。子任务E设计一个训练配置使loss在1000步内稳定下降。每个子任务都对应一个Agent可以独立完成的循环。Agent不需要理解整个“造AI”的宏大目标它只需要知道当前子任务的输入、输出和成功判据。注意任务拆解时一定要避免“隐式依赖”。比如子任务C依赖子任务B的产物那就要把B的输出路径明确告诉Agent而不是让它自己去猜。3.2 Agent框架选型为什么我最终选了轻量编排而不是重型框架市面上Agent框架很多从重型的多Agent协作平台到轻量的单文件循环都有。我在这个项目里试过三种方案最后落在一个轻量编排加自定义工具层的方案上。重型框架的问题在于抽象层太厚。当你需要Agent去执行一个具体的编译命令、解析一个具体的报错时框架的抽象反而成了障碍。你花在“怎么让框架支持这个操作”上的时间可能比直接写一个循环还多。轻量方案的核心就是一个while循环while not done and step max_steps: observation execute_action(action) action agent.plan(observation, memory) memory.append((action, observation)) done check_success(observation)这个循环里真正需要精心设计的是三件事action空间Agent能做什么。我通常限制为文件读写、命令执行、结果查询三类。observation压缩日志太长必须截取关键部分再喂给Agent。失败处理连续失败多少次后切换策略或者直接终止。实操心得observation压缩做得好不好直接决定Agent的“智商”。我一般会把stderr的前50行和后50行保留中间用省略号代替对于训练日志只保留loss、lr、grad_norm这几个关键指标的最后若干条。3.3 GPU Kernel生成的关键约束让Agent生成Kernel最容易翻车的地方不是算法而是硬件约束。以下是我在实际项目中总结的必须显式告诉Agent的约束目标GPU架构不同架构的warp大小、共享内存容量、寄存器数量都不同。不告诉Agent它可能写出在A100上能跑、在消费级卡上直接编译失败的代码。线程块大小必须是warp大小的整数倍且不能超过硬件上限。常见取值128、256、512。共享内存预算每个block能用的共享内存有限超了要么编译失败要么occupancy暴跌。寄存器压力寄存器用太多会导致occupancy下降反而变慢。数值精度float32、float16、bfloat16的累加行为不同必须明确。这些约束不能只写在prompt里最好做成可执行的检查工具。比如写一个脚本在Kernel编译后自动检查线程块大小、共享内存用量、寄存器数量把结果反馈给Agent。这样Agent不用“记住”所有约束它只需要根据反馈调整。3.4 RLVR的奖励设计别让Agent学会“作弊”奖励设计是RLVR里最危险的部分。我踩过的一个坑是早期奖励只看“编译通过数值正确”结果Agent学会了一个取巧策略——把所有计算都放到一个线程里串行执行。这样数值肯定正确编译也能过但性能惨不忍睹。后来我加了性能奖励但又出现新问题Agent开始生成只对当前测试shape优化的Kernel换个shape就崩。解决办法是测试集随机化每次评估时输入的shape、数据类型、内存布局都从一个大池子里随机采样Agent无法针对固定输入做特化。还有一个隐蔽的坑是奖励黑客。比如奖励里有一项是“代码行数少”Agent就会把代码压缩成一行可读性极差但奖励更高。所以任何非核心的奖励项权重都必须压得很低或者干脆去掉。奖励项早期设计修正后编译通过必须必须数值正确必须必须执行速度无与基线对比分档给分代码行数有权重高去掉测试shape固定随机采样显存占用无低权重3.5 SFT数据构造质量比数量重要一个数量级SFT数据构造的核心原则是每一条数据都必须是“可执行且正确”的。我见过太多项目为了堆数据量把网上爬的代码片段直接拿来用结果模型学了一堆错误模式后面怎么调都调不回来。我的做法是从已有正确实现出发比如PyTorch的官方Kernel、经过验证的开源实现。自动生成描述用规则或模型生成“这个Kernel做什么、输入输出是什么、关键优化点在哪”的描述。人工抽检随机抽10%的数据实际编译运行确认正确。去重和清洗去掉重复的、过于简单的、格式混乱的样本。数据量方面我的经验是500到2000条高质量样本就足够让模型学会基本模式。再多边际收益递减而且清洗成本急剧上升。4. 实操过程从零跑通一个Kernel生成Agent4.1 环境准备与依赖安装先说明这里的环境是Linux加NVIDIA GPUCUDA工具链齐全。如果你没有GPU后面的Kernel编译和执行部分跑不了但Agent框架和SFT部分可以先用CPU模拟。核心依赖# 基础Python环境 python -m venv agent_env source agent_env/bin/activate # Agent框架相关 pip install openai anthropic # 按你用的模型API来 pip install pyyaml jinja2 # 配置和模板 # Kernel相关 pip install torch numpy # CUDA工具链需要单独安装确保nvcc可用 nvcc --version注意不要在一个环境里混装多个版本的CUDA。我见过有人系统里同时有CUDA 11和12编译时链接到错误的版本报错信息完全看不懂。4.2 任务配置文件设计Agent需要一个清晰的任务描述。我用YAML来定义每个子任务task: name: generate_softmax_kernel description: 实现一个CUDA Kernel完成softmax计算 input_spec: shape: [batch, seq_len, hidden] dtype: float16 success_criteria: - 编译通过 - 与PyTorch参考实现误差小于1e-3 - 执行时间不超过基线的1.5倍 constraints: max_threads_per_block: 1024 max_shared_memory_kb: 48 target_arch: sm_80 max_attempts: 20 timeout_seconds: 300这个配置里success_criteria必须可程序判定。不要写“代码质量好”这种模糊条件Agent无法理解你也无法自动检查。4.3 Agent主循环实现主循环是整个项目的骨架。我把它简化到最核心的逻辑import subprocess import json from pathlib import Path class KernelAgent: def __init__(self, config, model_client): self.config config self.client model_client self.memory [] self.attempt 0 def run(self): while self.attempt self.config[max_attempts]: self.attempt 1 # 1. 让模型生成或修改Kernel代码 code self.generate_code() # 2. 写入文件 kernel_path Path(workspace/kernel.cu) kernel_path.write_text(code) # 3. 编译 compile_result self.compile_kernel(kernel_path) if not compile_result[success]: self.memory.append({ attempt: self.attempt, error: compile_result[stderr][:2000], action: compile_failed }) continue # 4. 运行正确性测试 correctness self.test_correctness() if not correctness[success]: self.memory.append({ attempt: self.attempt, error: correctness[message], action: correctness_failed }) continue # 5. 性能测试 perf self.test_performance() self.memory.append({ attempt: self.attempt, perf: perf, action: success }) if self.check_success(perf): return {status: success, code: code, perf: perf} return {status: max_attempts_reached, memory: self.memory}这个循环里memory的构造方式非常关键。不要把完整日志塞进去只保留“尝试编号、动作、关键错误或指标”。我一般限制memory总长度在4000 token以内超了就丢掉最早的记录。4.4 编译与执行的具体命令编译Kernel用nvcc但要注意几个参数nvcc -archsm_80 -O3 -shared -Xcompiler -fPIC \ -o kernel.so kernel.cu \ 2 compile_error.log-archsm_80指定目标架构必须和实际GPU匹配。-O3开启优化但有时会导致编译时间过长可以先用-O2调试。-shared -fPIC生成动态库方便Python调用。2 compile_error.log把错误单独存文件方便截取。执行测试时我通常写一个Python脚本加载编译好的.so用ctypes或torch的C扩展机制调用然后和PyTorch参考实现对比。import torch import ctypes # 加载编译好的Kernel lib ctypes.CDLL(./kernel.so) # 构造输入 x torch.randn(8, 128, 512, dtypetorch.float16, devicecuda) ref torch.softmax(x, dim-1) # 调用自定义Kernel具体调用方式取决于你的封装 # ... # 对比 diff (out - ref).abs().max().item() print(fmax diff: {diff})实操心得数值对比时float16的误差阈值不要设得太死。我一般用1e-3作为相对误差上限绝对误差用1e-2。设太严会导致Agent永远无法通过设太松又会放过错误实现。4.5 性能测试与基线对比性能测试最忌讳“只跑一次”。GPU上有太多噪声时钟频率波动、其他进程干扰、缓存状态不同。我的做法是预热10次不计时。正式跑100次取中位数。同时记录显存峰值。和PyTorch基线用同样的方式测保证公平。import time def benchmark(fn, warmup10, iters100): for _ in range(warmup): fn() torch.cuda.synchronize() times [] for _ in range(iters): start time.perf_counter() fn() torch.cuda.synchronize() times.append(time.perf_counter() - start) times.sort() return times[len(times) // 2] # 中位数基线就用PyTorch的对应算子。如果自定义Kernel比基线慢超过1.5倍就算失败Agent需要继续优化。4.6 SFT数据构造的实际流程SFT数据构造我走的是“半自动加人工抽检”的路线收集正确实现从PyTorch源码、经过验证的开源项目中提取Kernel实现。生成描述用规则模板生成“功能描述、输入输出规格、关键优化点”。构造训练样本把描述作为输入Kernel代码作为输出。自动验证每条样本都实际编译运行确认正确。人工抽检随机抽10%人工看代码和描述是否匹配。格式化统一成模型训练需要的格式比如JSONL。{ instruction: 实现一个CUDA Kernel完成softmax计算。输入shape为[batch, seq_len, hidden]dtype为float16。要求使用共享内存优化。, output: __global__ void softmax_kernel(...) { ... } }数据量控制在1000条左右训练1到2个epoch。太多容易过拟合太少学不到模式。4.7 RLVR训练的关键参数RLVR训练我用的是常见的策略梯度方法核心参数如下参数取值说明batch size32每次采样32个Kernel尝试learning rate1e-6比SFT小一个数量级KL系数0.01防止偏离SFT模型太远奖励缩放归一化到0-1避免奖励尺度差异过大采样温度0.8保持一定探索性最大生成长度2048 token覆盖大多数Kernel实现KL系数特别重要。设太小模型会为了奖励走偏生成一堆无法编译的怪代码设太大模型几乎不更新RLVR白跑。0.01到0.05之间是比较稳的范围。5. 常见问题与排查技巧实录5.1 Agent反复生成同样的错误代码这是最常见的问题。原因通常是memory没有正确反馈给模型或者反馈信息太模糊。比如只告诉模型“编译失败”不告诉它具体哪一行、什么错误模型只能瞎猜。解决办法把编译器的完整错误信息截取关键部分放进memory。在prompt里明确要求“不要重复上一次的错误”。如果连续3次生成相同错误强制切换策略比如让模型先输出伪代码再输出实现。5.2 Kernel编译通过但运行时报非法内存访问这种问题最折磨人。编译通过说明语法没问题但运行时崩通常是线程索引越界或共享内存越界。排查步骤用compute-sanitizer跑一遍它会告诉你具体哪一行越界。检查线程块大小和网格大小是否匹配输入shape。检查共享内存数组的索引是否超出声明大小。检查是否有未同步的读写。compute-sanitizer --tool memcheck ./test_kernel注意compute-sanitizer会显著拖慢执行速度只用于调试不要用于性能测试。5.3 数值正确但性能远低于基线这说明Kernel能跑但优化没做好。常见原因线程块太小比如只用32个线程GPU利用率极低。共享内存没用或用法不对该用共享内存的地方直接读全局内存。寄存器压力过大导致occupancy下降。内存访问不连续导致缓存命中率低。排查时先用nvprof或nsight compute看occupancy和内存吞吐再针对性调整。5.4 RLVR训练中奖励不升反降这通常意味着奖励函数有漏洞模型找到了某种“作弊”路径但这条路径在评估时被惩罚了。比如模型发现生成超长代码能提高编译通过率但评估时因为性能差被扣分。解决办法检查奖励函数是否有可以被“钻空子”的项。增加测试集的多样性防止过拟合。降低非核心奖励项的权重。如果问题持续回退到SFT模型重新设计奖励。5.5 SFT后模型只会模仿不会创新这是SFT的固有局限。SFT教的是“像训练数据那样写”不是“写出更好的”。如果SFT数据里全是基础实现模型就只会基础实现。解决办法SFT数据里混入一些优化技巧明显的样本。SFT之后必须跑RLVR让模型在奖励驱动下探索更优解。不要指望SFT一步到位它只是起点。5.6 常见问题速查表现象可能原因排查方向Agent重复错误memory反馈不足检查错误信息是否完整传入编译通过但崩溃索引越界用compute-sanitizer定位数值正确但慢优化不足看occupancy和内存吞吐奖励不升奖励函数有漏洞检查是否有作弊路径模型只会模仿SFT数据单一混入优化样本跑RLVR训练OOMbatch size或序列太长减小batch开梯度累积loss变nan学习率太大或数据有问题降lr检查数据清洗5.7 几个我踩过的坑第一个坑是过早引入多Agent协作。一开始我觉得多个Agent分工合作会更高效结果发现通信开销巨大而且一个Agent的错误会传染给其他Agent。后来改成单Agent加工具层反而更稳。第二个坑是奖励函数设计得太复杂。我一开始把编译、正确性、性能、代码风格、显存占用全塞进奖励里结果模型完全懵了不知道往哪个方向优化。后来砍到只保留编译、正确性、性能三项训练立刻稳定。第三个坑是测试集太小。早期只用固定shape测试模型生成的Kernel换个shape就崩。后来改成每次从池子里随机采样shape泛化能力明显提升。第四个坑是忽略编译时间。有些Kernel实现用了大量模板元编程编译要几分钟Agent的循环根本跑不起来。后来加了编译超时限制超过30秒直接判失败。6. 这条链路还能怎么扩展跑通Kernel生成之后我试过把同样的思路往上游推。比如让Agent去设计SFT数据配比或者去调整RLVR的奖励权重。效果有但不如Kernel生成那么稳定因为上游任务的“可验证性”更弱奖励信号更模糊。另一个方向是跨算子泛化。目前Agent在一个算子上调好的策略换一个算子往往要重新调。如果能让Agent学会“迁移优化经验”比如把softmax的共享内存优化思路迁移到layernorm上那价值会大很多。我试过用few-shot示例来引导迁移有一定效果但还不稳定。还有一个很实际的方向是把Agent接入现有的训练框架。现在很多团队都有自己的训练pipeline如果Agent能直接读取pipeline配置、生成优化建议、甚至自动提交实验那就能真正嵌入研发流程而不是停留在demo阶段。这块我还在摸索核心难点是权限控制和实验隔离不能让Agent的试错影响到主分支。最后说一个我个人的判断“AI造AI”目前最成熟的落点就是Kernel生成和训练配置调优这两个环节因为它们有明确的可验证信号。再往上的模型架构设计、数据配方设计短期内还是得靠人。Agent可以辅助搜索但还做不到自主决策。所以别被标题里的“有人做出来了”冲昏头做出来的是一部分不是全部。但这一部分已经足够让日常研发效率提升一个档次了。
