最近在折腾Agent自动写GPU内核这件事时我几乎把“生成代码、丢给编译器、看报错、再改”这套循环跑吐了。LLM明明能写出看起来像模像样的CUDA可一旦涉及性能、数据布局、调度策略它就像在没有仪表盘的车里飙车——看不到当前IR长什么样不知道哪个pass生效了也不清楚为什么带宽利用率只有三分之一。后来我把思路切到TIRx Harness这类Open Compiler Harness上才真正把“Agentic GPU Programming”从玄学变成了可调、可复现、可审计的工程流程。TIRx Harness简单说就是架在LLM Agent与GPU编译器、运行时之间的一层开放且结构化的控制平面。TIRx里的TIR指TVM体系里的Tensor Intermediate Representationx在这里取exchange的意思可以理解为“交换/接驳层”。它把TIR的读取、变换、代码生成、编译启动、profile反馈整体封装成Agent可调用的窄接口同时保留版本追踪、审计日志和回滚能力。这篇文章我会从背景痛点、模块设计、实操流程、关键取舍、问题排查五个角度把这类Harness拆开讲清楚。适合正在做AI驱动算子优化、编译器工具链、GPU内核自动生成或者纯粹被LLM写CUDA气到想换思路的读者。1. 为什么Agent写GPU代码急需一个Harness1.1 粗糙循环的三个致命问题先看现在最常见的Agent写CUDA的循环模型读任务描述生成一份CUDA代码丢给nvcc或tvmc编译把编译错误文本塞回上下文让模型再改。这套流程在小玩具上能跑但到了真实GPU编程场景问题非常明显。第一个问题是反馈的信噪比极低。编译器报错是给人眼看的里面充满了文件路径、行号、类型名称碎片、宏展开痕迹。Agent拿到这一大坨文本等于让它在噪声里捞信号。就算编译通过性能反馈也往往只是“运行耗时xx毫秒”这样一个粗糙标量它无法告诉Agent瓶颈到底在访存、共享内存bank conflict还是线程束发散。没有结构化的中间信息模型只能瞎猜猜错就再来一轮每一轮都是token和时间成本。第二个问题是IR状态完全不可见。Agent生成的是源代码文本但GPU编译器真正处理的是IR是已经做过类型推断、循环结构分析、内存布局标定的中间表达。源代码层看起来没问题到了IR层可能循环轴被错误绑定、vectorize宽度无效、shared memory都有溢出风险。Agent看不到IR快照就等同于一个医生只看化验单上的“异常”二字却看不到具体指标没法对症下药。第三个问题是优化动作缺乏约束和审计。让Agent直接改代码它可能同时动三个无关的东西改坏了根本不知道是哪一步引起的。真正的编译优化应该是受控的变换序列先向量化再绑定线程块再调整展开因子。每一步都需要可回滚、可对比。没有Harness层的约束Agent就是在没有安全气囊的车上做漂移。1.2 Harness是“仪表盘”而不是“自动驾驶”很多人会问现在AutoTVM、Ansor这类自动调优器不是已经能搜索最优schedule了吗还需要Harness做什么这里要分清楚定位。AutoTVM负责的是在固定搜索空间里用代价模型和进化策略找最优配置它是“自动驾驶”而TIRx Harness要做的不是取代搜索而是给Agent一个“可编程的仪表盘和控制通道”。你可以这样理解Agent是驾驶员Harness是带仪表盘、限速器和行车记录仪的底盘。Agent可以看到当前的IR结构、循环层次、内存足迹、线程绑定状态它可以请求执行一个变换但每个变换都会产生新版本IR并且记录审计日志它也可以调用profile模块拿回结构化的运行反馈。Harness不会替Agent做最终决策但会让Agent的每个决策都建立在可验证的信息之上。这也是Harness与普通编译器wrapper的本质区别。普通的“编译-运行”脚本只是个黑盒调用Harness则是一层契约它明确规定Agent能查什么、能改什么、不能改什么并且把这些能力暴露成稳定的、可编程的接口。这样Agent的推理能力才能真正发挥而不是在字符串和临时文件之间反复挣扎。2. TIRx Harness的整体架构与核心模块拆解2.1 四个层次视图、变换、执行、代理从架构上看TIRx Harness这类开放编译器接驳层通常会拆成四个模块IR视图层、变换注册层、执行反馈层、Agent接入层。这四个层次的目标很明确把编译器内部能力“API化”但又不让Agent直接触碰编译器内部的可变状态。IR视图层负责给Agent提供一个不可变的IR快照snapshot。它只允许查询不允许原地修改。比如输入一个prim_func视图层可以返回它的循环结构、线程绑定、buffer shape、数据类型、内存跨度。变换注册层管理所有可用的编译变换和pass pipelineAgent只能通过名称和参数请求变换由Harness验证合法性后在新的IR版本上执行。执行反馈层负责把IR编译成可运行的内核执行、计时、收集硬件计数器或profile数据然后打包成结构化的性能报告。最后的Agent接入层把前面三层包装成LLM友好的调用接口通常是一组Python函数或JSON协议。这种分层最大的好处是职责单一。视图层只管“读”变换层只管“改”执行层只管“跑”接入层只管“对话”。每一层都不依赖LLM的具体实现今天你接OpenAI的模型也好接本地微调的CodeLlama也好甚至接一个随机搜索策略也好接口都是一样的。这种设计让Harness本身与具体的Agent解耦未来换模型几乎零成本。2.2 IR视图层让Agent看到编译器的真实状态IR视图层是TIRx Harness最值得花功夫的地方。GPU代码优化里“看不见”是最大的敌人一个内核最后跑得慢原因可能藏在循环层级、内存布局、线程映射任何一个地方。视图层要做的就是把IR的这些关键属性显式地呈现给Agent。我自己的习惯是实现这样一组查询接口snap ctx.read_ir(version12) snap.loop_structure() # 返回嵌套循环的层级与轴名 snap.thread_binding() # 返回线程块/线程与循环轴的绑定关系 snap.memory_footprint() # 返回全局内存、共享内存、寄存器估计占用 snap.access_pattern() # 返回访存的连续性、步长、对齐信息 snap.call_graph() # 返回内联函数、外部调用、库函数依赖以TVM的TIR为例一个简单的saxpy内核在IR里长这样T.prim_func def saxpy( A: T.Buffer((8192,), float32), B: T.Buffer((8192,), float32), C: T.Buffer((8192,), float32), alpha: T.float32, ): for i in T.grid(8192): with T.block(compute): vi T.axis.S(8192, i) C[vi] alpha * A[vi] B[vi]IR视图层拿到这个prim_func之后可以回答很多源代码层面根本看不出来的问题。比如循环i是否已经被split成外层块循环和内层线程循环A的访问是否是连续的float32还是可以合并成float4buffer的memory scope是global还是shared这些信息对优化决策至关重要但你在纯CUDA源码层面很难自动获得。Harness的价值就在于把这些编译器内部信息变成可查询的一等公民。2.3 变换层受控的可组合优化原语IR视图解决“读”变换层解决“改”。但这里的“改”不是让Agent直接输出一堆新IR节点而是让它从一组经过验证的优化变换中选择并组合。这就是为什么open compiler harness比让Agent从头生成TIR要可靠得多。变换层通常维护一个注册表每个变换都有名称、参数模式、前置条件、副作用描述。比如ctx.register_transform( vectorize_width, predicatefragment_valid_on_contiguous_axis, apply_fnauto_vectorize, params_schema{width: int, target_axis: str}, ) ctx.register_pass_pipeline( fast_speedup, stages[simplify, vectorize_width, auto_bind, unroll_tail], )Agent通过ctx.apply(ir_version12, transformvectorize_width, params{width: 4})来请求变换Harness先检查前置条件比如目标轴是否连续、宽度是否能整除通过后产生一个IR的新版本并且保留原来的版本。每个变换都会记录ir_version_before和ir_version_afterAgent可以随时对比两个版本之间的差异。这种设计有意把Agent的“动作空间”限制在预设的变换集合里。原因是直接生成合法IR对LLM来说是要同时保证语法、类型、边界、绑定关系全部正确难度极大而选择变换只需要理解“此时应该向量化”这个语义实际变换的正确性由编译器保证。这就像给Agent一堆乐高积木而不是一堆乐高零件的图纸拼装难度完全不同。2.4 执行反馈层反馈给Agent的不该是文本而是结构化数据执行反馈层是闭环的落点。IR变换做完最终还是要编译、运行、测性能。这一层最核心的设计决策是反馈必须是结构化的而不是自然语言。我见过很多Agent系统把profile输出拼成一大段文本塞给LLM比如“kernel avg time 123.5 us, achieved occupancy 0.63, memory throughput 45.2%”这种文本浪费token且难以稳定解析。更合理的做法是让Harness返回一个带明确字段的字典Agent接入层再做一次摘要与序列化。举个例子一次典型执行反馈可以组织成{ ir_version: 12, compile_status: ok, launch_status: ok, duration_us: 231.4, achieved_occupancy: 0.63, bandwidth_utilization: 0.48, bottleneck_location: loop_axis_i, vector_width1, resource_estimate: {regs_per_thread: 32, shared_mem_bytes: 0} }这份数据如果转成prompt里的文字可以写成“IR版本12运行正常耗时231微秒占用率0.63带宽利用率只有0.48瓶颈在循环轴i的向量宽度为1”。模型看到这种信息做决策和看到“编译失败请检查代码”是完全两个量级的效率。执行反馈层还需要做的一件事是重复测量。GPU运行有噪声单次测量不可信。我一般每个配置跑5次去掉最高最低后取中位数。这个细节看似简单但能显著减少Agent被偶然抖动误导的情况。Harness应该在反馈里附带sample_count字段让Agent知道这个数据点的可信度。3. 实操用TIRx Harness跑通一个Agent化算子调优闭环3.1 环境准备与基础配置下面这部分我基于社区里做Harness类项目最常见的工程实践来展开具体环境会因人而异但大致链路是通用的。环境方面建议Python 3.10以上CUDA toolkit 12.x一个能调用的LLM接口OpenAI兼容格式、本地vLLM均可以及带有TVM或TIRx运行时依赖的Python环境。安装完依赖后核心是初始化Harness上下文from tirx import TIRxContext, HarnessAgent ctx TIRxContext( backendcuda, targetsm_89, # 按实际显卡填写比如RTX 4090对应sm_89 workspace_dir./workspace, enable_auditTrue, ) # 注册一批预置变换 ctx.register_default_transforms() ctx.register_pass_pipeline(speedup, [simplify, auto_bind, vectorize_width, unroll_tail])这个阶段最容易踩的坑是target设置错误。很多人图省事用--targetnvidia这种宽泛写法结果代码生成时生成的是通用NVCC路径没有针对具体架构做指令调度和寄存器分配优化。应该显式指定sm版本像是sm_89或sm_90a否则后面看到的性能数据会明显偏低。3.2 初始化种子内核并完成首轮反馈Harness不要求Agent凭空造内核你可以给它一个朴素但可运行的种子实现让优化工作从“改进”开始而不是从“创造”开始。这非常符合编译器从业者的习惯先有correctness再谈performance。seed_ir ctx.load_or_build_prim_func(saxpy, srcexamples/saxpy.py) version ctx.commit_ir(seed_ir, messageinitial naive implementation) feedback ctx.run(version_idversion, profile_rounds5) print(feedback.summary())首次反馈通常不会好看向量元素就是逐个循环带宽利用率大致在40%到60%之间这很正常。它的意义在于给Agent一个基线锚点。没有基线后面所有优化都是自说自话。3.3 Agent会话让模型基于结构反馈做决策真正进入Agent化流程后循环结构大概是这样一个模式agent HarnessAgent(llm_clientllm, ctxctx) session agent.new_session(taskoptimize saxpy bandwidth) session.observe(initial_feedback) for step in range(8): plan session.ask_for_plan( # 传给模型当前IR摘要 最近一次反馈 动作空间 ir_summaryctx.summarize(version), feedbackfeedback.get_prompt_payload(), allowed_actionsctx.action_schema(), ) if plan.finished: break # Harness层负责合法性和回滚 result ctx.apply( version_idversion, transformplan.transform, paramsplan.params, ) if not result.ok: session.observe_rejection(result.reason) continue version result.new_version feedback ctx.run(version_idversion, profile_rounds5) session.observe(feedback)这个循环有几个工程细节值得强调。一是每个plan必须显式包含transform和params不能让Agent直接输出一整段IR代码否则你会把编译器安全边界直接暴露给模型。二是rejection信息也要回传给Agent比如“vectorize_width在axis上不合法因为该轴not contiguous”这能帮模型快速修正策略。三是建议每步都调用ctx.diff(version_before, version_after)把IR差异摘要付给Agent让模型知道自己的操作产生了什么实际影响。3.4 从naive到向量化的实例复盘我用一个实际案例来说明整条链路如何工作。假设saxpy的初始IR是逐个元素循环首轮反馈显示瓶颈位置在loop_axis_i, vector_width1带宽利用率0.48。Agent读取这些信息后给出的第一个优化动作通常是vectorize_width4目标axis为i。Harness执行变换后IR的循环体从标量运算变成了连续4个float32的向量运算# 变换前 C[i] alpha * A[i] B[i] # 变换后示意 C[i:i4] alpha * A[i:i4] B[i:i4] # float4 宽度这看起来只是简单替换但底层涉及内存对齐检查、循环尾处理、bank conflict规避一堆问题。Harness的变换层自动处理这些约束Agent只需要知道“把宽度调到4”这个宏观决策即可。第二轮的反馈数据通常会有明显改善带宽利用率可能从0.48提升到0.7左右。此时Agent看到bottleneck字段开始变化不再指向vector width而是转向occupancy或thread数配置。它就可以继续调整block sizeplan { transform: set_block_size, params: {block_size: 256, grid_size: 1024} }这里有个常见误区block size不是越大越好。block太大可能寄存器溢出导致per thread可用寄存器变少block太小又会造成调度开销占比过高。较好的做法是让Harness在反馈里直接带上resource_estimate字段Agent看到寄存器占用和共享内存占用后自然会避开明显不合理的配置。整个过程中每一步都有IR版本号和profile数据对应最后收敛时的IR与初始IR的差异一目了然。对比基线最终效果可能是带宽利用率从0.48上升到0.72耗时降低约三成。这个提升幅度对单个算子来说已经很可观而且整个优化过程Agent只做了六到八次决策token消耗远低于反复手写CUDA的玩法。3.5 会话状态保存与结果审计Agentic优化跑完不是终点把过程保存下来往往更有价值。我会建议在Harness语境里把一次完整的会话记录成session.json里面包含{ task: saxpy bandwidth optimize, initial_ir_version: 1, final_ir_version: 7, steps: [ {from: 1, to: 2, action: vectorize_width, params: {width: 4}, feedback_delta: 0.20 bandwidth}, {from: 2, to: 3, action: set_block_size, params: {block_size: 256}, feedback_delta: 0.05 occupancy} ], total_cost_usd: 0.04, final_speedup: 1.38 }这个文件至少有三种用途一是有问题可以精确回溯是哪一步引入的二是用来评估LLM的决策质量比如它有没有反复试同一动作三是后续做Agent策略微调的时候这些历史是最有价值的训练样本。Harness的audit能力越强这套系统在团队里的可信度就越高。4. 设计背后的关键决策为什么Harness必须是这个样子4.1 为什么IR快照必须是不可变的在TIRx Harness里IR版本一旦提交就不会被修改所有变换产生新版本。这不是为了炫技而是有非常实际的原因。第一个原因是并发与隔离。Agent会话可能有多个如果它们共享同一个可变IR对象一个会话的改动会污染另一个会话的观察结果。不可变快照让每个会话都拥有自己独立的视角互不干扰。第二个原因是回滚能力。优化的过程本质上是试错Agent会把IR从v1一路变换到v15中途难免改出一个跑得更慢的版本。没有不可变快照你想回到v9就变得非常困难除非你在每一步都手动保存一份拷贝。版本化快照天然支持任意回退。第三个原因是审计与复现。团队协作时你需要能够回答“这个最终版本是怎么从初始版本一步步变过来的”。不可变版本链就是最自然的答案载体。每条边都有明确的变换动作和参数每一步都能复盘。4.2 为什么限制Agent直接修改IR而是操作变换这是我在设计Harness时思考最多的一个点。最初的想法很简单既然LLM连C和CUDA都能生成那直接让它生成TIR脚本应该也行。但实际上LLM直接生成IR的问题非常明显。IR是强约束的结构。TIR的合法程序必须满足SSA形式、类型一致、buffer边界可解析、thread binding合理、block迭代域与buffer形状匹配甚至还有TVM家特有的T.axis语义约束。让LLM一次性生成完整且合法的IR本质上是在要求模型做一个非常高难度的结构化生成任务。它的失败模式多到让人绝望语法对了类型不对类型对了边界不对边界对了layout又不合理。错误信息还特别晦涩模型的修正效率很低。相反把Agent的动作空间收敛到“选择变换”之后模型的任务从“从头构造一个合法IR”变成了“从当前IR出发选择一个能改善性能的合法步骤”。这是一个决策问题不是生成问题。决策空间内每个动作的正确性由编译器保证模型只需要承担策略选择的智力负担。实测下来这个模式收敛轮次和成功率都远超直接生成。当然这不意味着Agent完全不生成IR。在种子初始化阶段Agent生成一份“近似合法”的prim_func脚本再由Harness做类型补齐和layout归一化这仍然是一个重要场景。但进入迭代优化阶段以后主导逻辑应该是变换选择而不是代码重写。4.3 为什么反馈必须结构化不能直接在prompt里堆文本LLM的上下文窗口再大也是有限资源而GPU内核的profile输出、编译日志、IR dump都是非常吃token的东西。如果让Agent直接面对原始输出它会很快在噪声中迷失。结构化反馈的核心价值有二。第一是信息密度高一个包含duration、occupancy、bandwidth utilization、bottleneck_location的字典用100个token就能表达过去几千行日志才能传递的信息。第二是方差小文本解析有随机性模型可能漏看某个关键指标而字段明确的JSON则保证每次都能被完整接收。你在设计get_prompt_payload()时不应该把整个反馈字典原封不动丢给模型而是先做一层摘要。比如只保留top-3瓶颈指标、最近两次IR版本的差异摘要、以及当前资源估计。把“全量数据”留给Harness日志把“决策所需最小信息”交给模型。这个最小集的设计通常需要试验几次才能稳定但值得投入。4.4 为什么这种Harness形态能跨后端扩展TIRx Harness选择以TIR为核心表示不是偶然因为TIR是后端无关的。同一个IR视图层可以接CUDA后端也可以接ROCm、Metal或OneAPI。执行反馈层的接口抽象成Runner每种后端只是一个插件实现而已。这意味着你为CUDA写的变换、Agent策略、反馈摘要逻辑未来切换到AMD或Intel GPU时依然沿用。Harness的“开放性”不仅指代码开源更指它的后端边界清晰。这样做让Agentic GPU Programming不再绑定到某一家硬件厂商而是变成一套可以在不同算力平台上复用的方法论。我在实际使用中还体会到这种跨后端设计对策略演进的好处。如果今天只针对CUDA优化Agent学的很多经验是CUDA特定的但如果你把变换抽象成“向量化”“平铺”“线程绑定”这些通用动作模型学习到的策略就能迁移到更多硬件平台。这也是Harness与直接在CUDA源码文本上做Agent优化相比长期价值大得多的原因。5. 常见问题与排查技巧实录5.1 Agent反复尝试同一个无效变换怎么办很多跑过Agent优化的人第一个头疼的问题就是模型陷入死循环重复提交vectorize_width8之类结果被拒的动作。根因通常不是模型笨而是反馈里缺乏“为什么被拒”的解释信息。Harness的run()返回rejection_reason之后一定要把这个reason原样传给Agent而不是只给一个布尔值。如果加了reason还是重复尝试就需要在接入层做去重。我在Harness里维护一个tried_configs集合对每次plan的(transform, params)做签名重复的动作直接拦截并提示“该配置已在第3步尝试过当时带宽利用率0.52”。这种硬性去重能有效终结循环。还有一个隐藏原因搜索空间太窄模型不知道除了调vectorize还能干什么。这时候需要扩充允许动作列表或者把cost model的中间估算结果也暴露给Agent让它知道还有别的地方可以优化。5.2 IR行号与源码行号对不上这个问题几乎必然出现。TIR在生成CUDA源码时经过了多层改写IR节点和最终源码的位置映射变得很弱。你在feedback里写bottleneck_location: loop_axis_i, line 9这个line 9指的是IR层的行号Agent完全无法对应到实际代码。更好的做法是在IR节点上挂origin元数据。Harness在做代码生成时尽量为每个IR节点保留它来源的源码标记反馈时也使用IR路径而不是行号比如blockcompute, loopi, iter0这种结构化的位置描述。这样Agent虽然拿不到行号但能理解当前瓶颈在哪个逻辑层。5.3 pass顺序冲突导致性能回退变换不是越激进越好。一个很典型的坑是先做了unroll_tail再做auto_bind导致循环轴绑定关系乱掉反而出现性能回退。Harness要做的就是记录pipeline的顺序并在Agent生成transform组合时附带一个“顺序合法性校验器”。例如unroll不能在任何依赖loop structure分析的pass之前vectorize的要求是访存连续等。这些约束可以在register_pass_pipeline阶段配置成静态规则。Agent请求pipeline时Harness先整体校验一遍不合法就拒绝并返回哪个阶段的条件未满足。这样既保留灵活组合能力又避免无意义尝试。5.4 显存越界没有立刻报错后面才随机崩溃GPU编程里最阴间的bug就是不报错。Agent生成的IR可能存在某个分支没有做边界保护实际触发的数据量刚好没越界换个输入规模就崩。Harness层必须主动做静态边界检查而不是等launch报错。我在变换层实现了safe_bounds_analysis它检查IR中每个buffer访问的索引范围是否落在声明区间内检查步长和尾处理是否覆盖完整。如果某个变换会引入潜在越界直接拦下。这个检查同样可以作为Agent决策的一个可见属性让模型理解“当前候选方案虽然性能高但有越界风险”。5.5 Harness会话挂起导致整个流程卡死Agent化流程一旦跑起来如果某个内核launch永远不返回整个会话就卡死了。我建议把执行反馈层的launch调用统一包进超时机制比如内核执行超过3秒直接kill并标记launch_statustimeout同时把栈信息和占用资源记录进日志。这个超时时间最好可配置因为有的极热点确实需要较长时间。更稳妥的方式是让Agent的运行环境与主进程隔离跑在子进程或独立容器里。这样即使某个内核彻底悬挂主流程也只是收到一个session_crashed信号而不是整个训练或调优任务全部终止。一些实操中的真心话我自己的工具链最早只是给LLM套了个编译脚本跑通后才发现真正的瓶颈从来都不是“模型写不出代码”而是“模型看不到自己写的东西在编译器和硬件之间发生了什么”。TIRx Harness这类开放编译器接驳层给我的最大启发是Agent的能力边界由反馈质量决定而不是由模型参数量决定。你给它一个结构清晰、可验证、可回滚的动作空间它就能在我们的领域里发挥出远超文本生成的价值。最后分享一个特别有用的小技巧在每次反馈里附上“本次未被采纳的候选变换及其原因”比如“你已经考虑过vectorize_width8但因为axis长度不能被8整除而作废”。这听起来是在给Agent负面信息但实际效果非常显著。我加了这个字段之后同一个任务的收敛轮次从18轮降到了9轮因为它大幅减少了模型对空空间的探索成本。如果你也在搭类似的Agentic GPU Programming链路值得试一试。