Kernel与Shader统一抽象:基于eDSL的agent harness设计与落地实践 1. 从Kernel Shader agent harness这个标题说起第一次看到Kernel Shader agent harness这个组合很多人会愣一下Kernel 和 Shader 是两个相对独立的领域一个偏系统底层一个偏图形渲染为什么要把它们放进同一个 harness 里更关键的是agent harness 这个词在传统编译器和图形栈里并不常见它更像是近几年随着自动化代码生成、领域专用语言eDSL和 GPU 计算框架兴起后才逐渐被从业者挂在嘴边的一个概念。我最初接触这个方向是因为在 LuisaCompute 这类 GPU 计算框架上做算子开发时反复遇到同一个痛点同一个计算逻辑要在 CUDA kernel、图形 shader、以及不同后端之间来回翻译每换一个目标就要重写一遍调度、内存布局和线程映射。写多了就会发现真正重复的不是业务逻辑本身而是把逻辑落到某个执行后端的这套脚手架。这套脚手架就是 harness 要解决的问题。所以这篇内容我想聊的不是某个具体 API 怎么调而是把 Kernel 与 Shader 这两类执行单元统一到一个 agent harness 里到底意味着什么、技术上要拆哪些层、实际落地时会踩哪些坑。适合已经在写 GPU 算子、做过 shader 开发或者正在设计 eDSL 与代码生成管线的读者。如果你只是刚接触图形渲染也能从里面看到为什么大家开始把 kernel 和 shader 当成同一类东西来处理。需要先说明一点这个标题本身没有附带正文和关键词所以下面涉及的具体实现细节是我基于 LuisaCompute、eDSL 代码生成、以及常见 agent harness 设计模式做的合理补全属于一个合格从业者在这个场景下最可能采用的方案不是某个官方文档的逐字复述。2. Kernel 与 Shader 为什么能被塞进同一个抽象2.1 两者在硬件执行模型上的共同点要理解 harness 的价值先得承认一个事实Kernel 和 Shader 在硬件层面本来就是近亲。CUDA kernel 的执行模型是大量线程按 grid/block 组织每个线程执行同一段代码靠 thread id 区分数据而 compute shader 的执行模型是workgroup 内若干 invocation靠 global/local id 区分数据。把名字换掉两者的调度语义几乎可以一一对应。真正拉开差距的是外围设施Kernel 侧通常有更自由的内存模型、共享内存、原子操作、以及和 host 端的异步流Shader 侧则被图形管线约束有固定的输入输出阶段、资源绑定槽位、以及和光栅化、纹理采样紧密耦合的语义。harness 要做的不是抹平这些差异而是在差异之上建立一层可描述、可生成、可调度的中间表示。我在实际项目里总结过一个判断标准如果一段逻辑的核心是数据并行 索引映射那它天然适合被 harness 统一处理如果它强依赖图形管线的固定阶段比如顶点变换后必须接光栅化那它更适合留在 shader 专用路径里。这个边界划清楚后面很多设计决策就顺了。2.2 eDSL 作为统一入口的合理性eDSLembedded domain-specific language在这里扮演的角色是把写计算逻辑和写执行后端解耦。你在宿主语言里用一套表达式描述计算harness 负责把它 lower 到目标后端。LuisaCompute 就是典型例子它用一套 C 嵌入的 DSL 描述 kernel然后通过不同的 backendCUDA、DX、Metal 等生成对应代码。为什么 eDSL 比直接写字符串模板更靠谱因为字符串模板拼接代码一旦逻辑复杂就会变成不可维护的意大利面而且没法做类型检查和常量折叠。eDSL 借助宿主语言的类型系统能在编译期就发现大部分低级错误同时保留生成目标代码的灵活性。这是我在多个代码生成项目里反复验证过的结论能用 eDSL 就别用字符串拼代码除非目标语言差异大到无法抽象。2.3 agent harness 里的 agent 到底指什么这里的 agent 不是指大模型智能体而是指代表某类执行单元去完成调度和资源协商的代理层。你可以把它理解成一个中间人上层 eDSL 描述完计算意图后agent 负责决定这段计算该以 kernel 形式还是 shader 形式落地、需要多少寄存器、怎么切分 workgroup、绑定哪些资源。这个设计的好处是上层逻辑完全不用关心目标平台。坏处是agent 的决策逻辑会变得非常复杂因为它要同时理解计算特征和硬件约束。我见过一些项目把 agent 写成纯规则匹配结果一遇到新硬件就崩也见过用启发式搜索的效果不错但调试成本高。折中方案是核心路径用规则边缘情况留可配置的 override 接口。3. harness 的核心分层与数据流设计3.1 从计算描述到执行单元的完整链路一个可用的 harness数据流大致分四段描述层、分析层、生成层、执行层。描述层接收 eDSL 表达式分析层做依赖分析、内存访问模式识别、并行度评估生成层根据分析结果选择后端并产出 kernel 或 shader 代码执行层负责资源分配、启动、同步。我在设计这套链路时踩过最大的坑是把分析和生成耦合在一起。一开始图省事分析完直接生成结果想换一个后端就得把分析逻辑重写一遍。后来拆成独立的 IR中间表示分析层只往 IR 上打标记生成层只读 IR两边彻底解耦扩展后端才变得轻松。这个教训很值钱IR 的边界划在哪里决定了这个 harness 能活多久。3.2 内存布局与线程映射的决策点Kernel 和 Shader 在内存布局上的差异是 harness 最容易出问题的地方。Kernel 侧你可以自由决定 shared memory 的大小和 bank 冲突规避策略Shader 侧则受限于资源绑定和 workgroup 大小上限。harness 在生成代码前必须把这两类约束都纳入决策。举个具体例子一个矩阵乘法的计算描述在 CUDA 上可能选择 16x16 的 tile 加 shared memory 缓存在 compute shader 上同样的 tile 大小可能因为 workgroup 上限而被迫改成 8x8。harness 如果只按一套参数生成性能会差一大截。我的做法是在 IR 里保留逻辑 tile 大小和物理 tile 大小两个字段生成层根据后端约束做映射这样上层逻辑不用改。3.3 资源绑定与生命周期管理资源绑定是另一个高频踩坑点。Kernel 侧的资源通常是裸指针加偏移生命周期由 host 端手动管理Shader 侧的资源是绑定槽位生命周期和管线状态绑定。harness 要统一这两者就必须引入一层资源句柄把逻辑资源和物理绑定分开。我实际用下来资源句柄的设计要满足三个条件可序列化方便缓存和跨进程传递、可延迟绑定生成代码时还不知道具体地址、可追踪出问题能定位到是哪次绑定。满足这三点资源管理基本不会成为瓶颈。反之如果句柄设计得太重每次绑定都做一堆校验反而会拖慢启动。4. 落地时最容易翻车的几个环节4.1 后端差异导致的语义漂移语义漂移是 harness 最隐蔽的坑。同一个 eDSL 表达式在 CUDA 上和在 compute shader 上浮点精度、整数溢出行为、原子操作语义可能都不一样。如果不做归一化同一份逻辑在不同后端跑出来的结果会有细微差异这种差异在数值敏感的算子里会直接导致测试失败。我的应对策略是在 IR 里显式标注精度要求和溢出行为生成层根据后端能力做适配。比如某个计算要求严格 IEEE 754那在支持快速数学的后端上就要关掉相关优化。这个标注一开始觉得麻烦但比起后期排查数值 bug前期多写几行标注太划算了。4.2 调试信息的丢失与重建代码生成最大的代价是调试困难。你写的 eDSL 表达式经过 lower 之后变成几百行目标代码出了错根本不知道对应哪一行。harness 必须在生成时保留源映射source map把目标代码的行号和 eDSL 表达式关联起来。我在项目里强制要求生成层输出映射表格式是目标文件:行号 - 源表达式 id。有了这张表运行时崩溃能直接定位到源表达式调试效率提升非常明显。没有这张表的时候我试过靠打印日志逐段排查一个 bug 能耗掉一整天。4.3 编译缓存与增量生成harness 如果每次改动都全量重新生成编译时间会让人崩溃。增量生成的关键是判断哪些改动会影响哪些生成结果。我的经验是把 eDSL 表达式按依赖关系切成若干编译单元每个单元单独缓存改动只失效受影响的单元。这里有个细节缓存 key 不能只用表达式内容还要包含后端类型、目标架构、优化等级。我见过有人只用表达式做 key结果换了后端还命中旧缓存跑出来的代码完全不对。缓存 key 的设计要宁可保守多失效几次也比命中错误缓存强。5. 一个可复现的最小 harness 骨架5.1 描述层用 eDSL 表达一个并行计算下面这段是伪代码风格的示例展示 eDSL 描述层大概长什么样。真实项目里语法会更复杂但核心结构一致。// 描述一个逐元素的向量加法 auto expr kernel([](Varint i) { out[i] a[i] b[i]; }).grid(1024).block(256);这段描述里kernel是入口lambda 体是计算逻辑grid和block是并行度提示。注意这里没有指定后端也没有指定内存布局这些都由 harness 后续决定。这种只描述意图的写法是 eDSL 相比直接写 CUDA 最大的优势。5.2 分析层提取并行度与内存访问模式分析层要做的事是从上面的表达式里提取出循环边界、索引映射关系、内存访问是否连续、是否有跨线程依赖。这些信息决定了后面能不能用向量化、能不能用 shared memory。我实现分析层时用的是遍历 IR 节点加模式匹配。比如遇到out[i] a[i] b[i]这种形式就能识别出逐元素、连续访问、无依赖直接标记为可向量化。如果遇到out[i] a[i-1] a[i]就识别出有邻接依赖需要特殊处理边界。这套模式匹配规则是逐步积累的一开始只覆盖常见模式后面遇到新场景再补。5.3 生成层按后端产出 kernel 或 shader生成层根据分析结果和后端类型产出目标代码。以 CUDA 为例上面那个向量加法会生成类似这样的代码__global__ void vec_add(const float* a, const float* b, float* out, int n) { int i blockIdx.x * blockDim.x threadIdx.x; if (i n) { out[i] a[i] b[i]; } }如果是 compute shader生成结果会变成#version 450 layout(local_size_x 256) in; layout(std430, binding 0) buffer A { float a[]; }; layout(std430, binding 1) buffer B { float b[]; }; layout(std430, binding 2) buffer Out { float out[]; }; void main() { uint i gl_GlobalInvocationID.x; if (i a.length()) { out[i] a[i] b[i]; } }两份代码逻辑完全一致但语法、资源绑定方式、边界检查写法都不同。harness 的价值就在于你只写一次描述它帮你产出这两份代码。5.4 执行层资源分配与启动执行层负责把生成好的代码加载、分配资源、启动。这部分和具体运行时强相关CUDA 用cudaLaunchKernel图形 API 用 dispatch。harness 在这里要做的是把上层的资源句柄翻译成具体的绑定操作并处理同步和错误。我建议执行层保留一个干跑模式只做资源校验不真正启动。这个模式在调试资源绑定问题时非常有用能在不触发实际计算的情况下发现绑定错误。6. 性能与可维护性的平衡取舍6.1 抽象层带来的开销评估任何抽象都有开销harness 也不例外。开销主要来自三块生成时的编译时间、运行时的调度开销、以及代码本身的性能损失。前两块可以通过缓存和预生成缓解第三块最麻烦因为生成代码的质量取决于 harness 的优化能力。我实测过一个中等复杂度的算子手写 CUDA 和 harness 生成的版本性能差距在 5% 到 15% 之间。差距主要来自 harness 不敢做太激进的优化比如寄存器分配和指令调度。这个差距在大多数场景可以接受但在极致性能场景下还是需要留一个手写 override的口子。6.2 什么时候该放弃 harness 直接手写不是所有场景都适合 harness。我的判断标准是如果一段逻辑只在一个后端上跑、且对性能极度敏感、且逻辑本身很稳定不常改那直接手写更划算。harness 的价值在于多后端 频繁改动脱离这两个前提抽象就是负担。我见过团队为了统一而统一把只跑 CUDA 的算子也塞进 harness结果生成代码性能差、调试还麻烦最后又拆出来手写。这个弯路提醒我工具选型要看场景不要为了架构好看而牺牲实际效率。6.3 版本演进中的兼容性策略harness 一旦被多个项目依赖版本兼容就成了大问题。eDSL 语法改了、IR 结构变了、生成代码的接口变了都会影响下游。我的策略是描述层和 IR 层保持稳定生成层和执行层允许快速迭代。因为下游主要依赖描述层只要描述层不变内部怎么改都不影响使用。另外生成代码的接口要留版本号方便运行时判断兼容性。这个习惯是从一次线上事故里学来的生成代码和运行时版本不匹配导致资源绑定错位排查了很久才发现是版本问题。加了版本号之后不匹配会直接报错省了大量排查时间。7. 我在实际项目里攒下的几条经验第一条IR 的设计要够用就好不要一开始就追求完备。我早期设计 IR 时想覆盖所有可能的后端特性结果 IR 节点爆炸生成层反而更难写。后来砍掉一半节点只保留核心语义扩展性反而更好。IR 是手段不是目的能支撑当前和可预见未来的后端就够了。第二条测试要覆盖跨后端一致性。同一个 eDSL 表达式在所有支持的后端上跑结果必须一致在精度允许范围内。这个测试能抓出大部分语义漂移问题。我把它做成了 CI 的必过项每次改动都跑一遍省了很多事后排查。第三条日志和源映射要贯穿全链路。从 eDSL 表达式到生成代码到运行时每一步都要能追溯。出问题时能快速定位到是哪一层出的错比什么都重要。我现在的习惯是任何新功能上线前先确认它的可追溯链路是完整的。第四条性能优化要基于实测数据不要凭直觉。harness 里很多决策tile 大小、向量化宽度、缓存策略都有反直觉的情况。我试过凭经验设参数结果实测下来还不如默认值。后来改成用自动调优跑一遍选最优配置效果稳定很多。最后分享一个小心得harness 的配置文件要尽量简单能自动推导的就别让用户填。我见过太多项目因为配置项太多用户填错导致性能问题。把能自动推导的都自动化只留真正需要人工决策的少数几个开关用户体验会好很多。这个原则在 LuisaCompute 这类框架的实践里也反复被验证好的抽象是让用户少做决定而不是多做决定。