最近大半年，我让 LLM 帮我写 CUDA 算子的频率越来越高，省是真省，坑也是真多。你会发现它生成的 kernel 基本逃不出四种状态：语法能编但逻辑错、逻辑对但慢到不能用、表面完整但内存访问越界、以及直接对着错误 API 一顿幻觉输出。最可怕的是，很多问题不是一眼能看出来的，等我把这些代码塞进工程再跑回归，大概率就是一个通宵。被折腾几次之后，我实在受不了了，就在 LLM 生成脚本和真实编译之间加了一道自动拦截层，用一套双轨脚本先做快速过滤。这篇文章就把这套 CMP-D V2.0-Lite 的设计思路和落地过程完整吐出来，包括它的静态检查规则、动态校验逻辑、以及为什么这套东西能拦掉大约七成垃圾 CUDA 算子。
这套方案没有用到多高深的技术，就是“静态特征扫描 + 动态编译验证”两个轨道并行，再统一汇聚成 PASS / REVIEW / REJECT 三个结论。但它恰好戳中 LLM 生成代码的核心痛点：模型擅长“造句”，不擅长“验证”，所以我们要做的事情，就是把验证环节自动化、前置化、可解释化。
1. 先用一个场景说清楚：LLM 生成的 CUDA 算子到底哪些环节会出问题
1.1 从一条指令看 LLM 生成的失效模式
我平时的主要工作场景是用大模型辅助写各类 AI 框架底层算子，包括常见的 Elementwise、Reduce、MatMul、卷积以及一些偏门算子。听起来很理想，但真正实操过的朋友都知道，LLM 生成的 CUDA 代码有一个典型特点：上下文越长，越容易在关键细节上“一本正经地胡诌”。比如让它写一个把 tanh 融合进自定义算子的 kernel，它可能先给出正确的数学表达式，但接着就在 blockIdx 和 blockDim 的维度映射上犯浑，或者干脆把threadIdx.x * blockDim.x + threadIdx.x这种低级错误写出来。最让人头疼的还不是语法错误，而是那种“看起来完全正常、放到 CPU 参考实现一对比就翻车”的代码。我把这些现象统称为“无效生成”，它们占了生成总量的大头。
后来我把一批历史生成的坏算子拉出来统计，大概分成三类：
- 第一类，编译层问题。占比最高，大约四成。包括拼错 API 名、类型不匹配、头文件缺失、
__global__写成了__global、动态共享内存没加 extern 声明等。 - 第二类，语义层问题。占比约两成。代码能编过，但数值或索引逻辑不对，跑出来的结果和预期差得离谱。
- 第三类，性能层问题。占比一成半。算法和语义都对，但访存局部性极差、并发粒度不合理、存在大量分支发散，实际跑起来比手写版本慢三到五倍。
这三类加起来，已经超过七成。而且这三类问题非常有规律，完全可以用脚本自动识别。这就是我下定决心做拦截系统的直接原因。
1.2 为什么人工 review 兜不住这个场景
有人可能会问：让团队里的资深工程师 review 不就行了？问题是，LLM 生成算子的频率远高于人工 review 的吞吐能力。我自己一天能让大模型生成二三十版 kernel，如果每一版都要人肉看一遍，基本不用干别的活了。更麻烦的是，人工 review 对“编译能过”的代码往往放松警惕，而 LLM 翻车恰恰就翻在这些隐蔽地方。也就是说，高频生成带来的筛选压力，已经超出了一个资深工程师人工 review 的合理负载上限。
所以，我的结论是：必须在“生成”和“进工程”之间放一个自动化的闸口。它的目标不是替代人类专家，而是把那些低质量、高概率出问题的样本提前过滤掉。这样剩下的代码再让人看，压力就小很多了。
2. CMP-D V2.0-Lite 的双轨思路：静态特征识别 + 动态行为验证
2.1 第一轨：静态轨（static_review.py）
CMP-D V2.0-Lite 的“双轨”，第一轨是纯静态检查，也就是不编译、不运行，只对 CUDA 源码做特征识别。
这一轨的核心思路，是把过去人工 review 时总结出来的“垃圾特征”固化成规则。具体实现上，我混用了三种手段：
- 基于正则表达式的快速扫描。适合匹配 API 幻觉、危险函数、明显拼写错误。
- 基于简单词法分析的上下文判断。比如判断
threadIdx.x是否出现在全局内存索引表达式中，是否跟blockIdx一起出现等。 - 基于 clang-format / tree-sitter 的结构解析。对 C++ 模板、宏定义较多的代码，正则容易误伤，需要用结构级别的 AST 信息做二次校验。
静态轨的好处是轻量、快、不依赖 GPU。一次扫描通常几十毫秒到几百毫秒，非常适合在生成端做实时过滤。
2.2 第二轨：动态轨（dynamic_check.py）
第二轨是动态校验，核心目标是“让代码自己跑一遍，用事实说话”。
动态轨包括两部分操作：
- 第一步，编译。调用本机已安装的 CUDA 工具链，对生成的
.cu文件执行nvcc编译。这一步能直接拦住编译层面的问题。 - 第二步，微型运行验证。把 kernel 加载到小尺寸数据上，和 CPU 参考实现做数值对比；如果有条件，再自动跑一组不同 shape 的规格化测试，检查边界行为。
动态轨的费用比静态轨高很多，所以我把它放在静态轨之后。只有静态检查没有硬伤的代码，才会进入动态轨。这样能避免大量垃圾代码白白占用编译时间。
2.3 两轨怎么汇合
两条轨道的输出最终汇成一个三态结论：
- REJECT：明确不合格，直接打回给生成侧，附带原因标签。
- REVIEW：存在疑点，但暂时没有一票否决的证据，需要人工看一眼。
- PASS：双轨都没有发现问题，可以进入下一步集成。
这个设计有一个很关键的好处：它不是一个“单点”拦截器，而是一个分级过滤器。只要静态轨看到明显的错误模式，就不会浪费 GPU 时间去动态验证；只有那些看起来很正经、需要运行时证据的代码，才值得花编译时间来做确认。这样组合下来，单条算子的平均检查时间能压到秒级。
3. 静态轨里真正有效果的 8 条检查规则
3.1 规则一：核函数是否存在有效的全局内存访问
LLM 特别喜欢生成空壳 kernel。它可能把数学公式写对了，计算结果也存到某个寄存器里，但最后忘了写回全局内存。这条规则的做法是：先定位__global__函数体，再检查里面是否存在[]下标写入，或者atomicAdd、memcpy等真实生效的内存修改动作。如果没有，直接 REJECT。
这里有个容易误伤的地方：比如 kernel 里调用了别的辅助函数做内存写入，主函数体里没有直接写全局内存。所以静态轨先做标记，不直接判死，而是等动态轨做一次数值一致性的强制校验。
3.2 规则二：索引空间与数据尺寸是否匹配
这是重灾区。LLM 生成的 kernel 经常出现 grid 维度、block 维度和数据 size 对不上的问题。典型情况是：
- 总线程数算多了，导致越界 read；
- 总线程数算少了，导致部分数据没被处理；
gridDim和blockDim的乘除关系写反。
我对这个规则的做法是：先正则提取 kernel launch 配置里的<<<...>>>参数，再和 kernel 内部引用的 buffer 长度声明做比对。如果发现launch的线程总数明显小于或大于数据长度，就标记为“索引空间不匹配”。
这里需要提醒一句：CUDA 的索引是三维的，很多 LLM 会生成二维 grid 去访问一维数组，导致blockIdx.y参与计算时出现不可控偏差。静态轨虽然做不到精确证明，但可以通过检查 “blockIdx.y 是否被使用但 launch 配置只在 x 维度上有值” 来拦掉一批。
3.3 规则三：API 调用黑名单
LLM 对 CUDA API 的幻觉问题很严重。常见的 API 黑名单包括：
- 在
__global__函数里调用printf做调试输出，这个虽然能编译，但部署时就是垃圾； malloc出现在 device code 里，除非是特殊情况，否则就是性能隐患；- 把
cudaMemcpy的cudaMemcpyHostToDevice和cudaMemcpyDeviceToHost搞反，或者复制字节数用错； - 使用不存在的
cudaThreadSynchronize之类的错误拼写。
API 黑名单有很强的时效性，因为 CUDA Toolkit 的接口一直在演进。我目前的做法是维护一个 JSON 配置文件，把已知的合法 API 和非法高风险 API 都列进去。静态轨遍历源码里的所有 API 调用点，做白名单和黑名单的双重匹配。
3.4 规则四到规则八：宏、模板、同步和类型推断
- 规则四：危险的宏定义。比如
#define TILE_SIZE 32后面又出现#undef TILE_SIZE，或者宏名跟 kernel 内的局部变量同名，导致展开后变成无法阅读的怪物。 - 规则五：过度模板化。模板本身不是问题，但 LLM 经常会把内部循环也模板化，生成大量实例化代码，编译时间和二进制体积都会爆炸。遇到这类代码我会先打 REVIEW，让工程师判断是否有必要保留模板。
- 规则六：共享内存边界失控。动态共享内存和静态共享内存混用时，LLM 经常把大小算错。静态轨会检查
extern __shared__的声明和 kernel launch 配置里动态共享内存大小是否匹配。 - 规则七：同步缺失或同步错位。在多阶段写读的 kernel 里，
__syncthreads()如果缺失，动态轨大概率能抓到错误结果；但静态轨可以先检查一个非常明显的模式：先写 shared memory，后在同线程内读其它线程写入的数据，这中间必须有__syncthreads()。 - 规则八：类型宽度不匹配。比如
float*指针却被当double*来解引用，或者int和unsigned int比较时出现隐式转换问题。
这些规则看起来零散，但它们都是我在过去半年里一条条攒出来的，对应的都是 LLM 真实翻车案例。静轨的关键不是规则多，而是规则长短结合、相互配合。能用正则的尽量轻量，正则表达不清的再用 AST 兜一层。
4. 动态轨怎么不把 GPU 当试验田
4.1 编译拦截：能省一次编译就省一次
动态轨的第一步当然是编译。但这里有几个很实际的工程细节：
- 首先，不要直接把整个工程拉起来编译。CMP-D V2.0-Lite 只会对生成的
.cu文件做单文件编译，并且只加-c参数生成 object，不做完整链接。这样能把编译时间压缩到最低。 - 其次，要严格控制编译参数。我在内部统一采用
-std=c++17 -arch=sm_80 -O2 -lineinfo。这里-lineinfo是为了后续做性能 profiling 时能看到行号映射，方便追问题。 - 第三，做一个编译缓存。相同的源码 hash 不做二次编译，直接复用上一次的编译结果。别小看这一步，它能拦下一大半重复生成的无效代码。
需要指出的是，不同 CUDA 版本对同一个.cu文件的解析结果有差异。之前遇到过同一段代码在 CUDA 11.8 能过、在 CUDA 12.3 报错的情况。所以我在脚本里写死了当前使用的 nvcc 路径，防止系统里出现多版本 CUDA 时，动态轨和实际编译环境不一致。
4.2 数值等价性校验：用一段尽量短的 host 参考实现
能编译通过，不代表算子是对的。动态轨的第二个关键动作，是加载一个轻量的 CPU 参考实现，对比 GPU 输出。
我在脚本里提前放了常见算子的参考 kernel，例如 ReLU、Sigmoid、LayerNorm、MatMul、自定义 tanh 融合算子。每次校验时，CMP-D 会生成一个随机输入的小张量，分别跑 CPU 参考实现和 GPU 生成 kernel，然后对比输出张量的误差。
这个对比的细节比较多：
- 对于 float 类型，误差阈值不能卡太死。我一般用 atol + rtol 的组合，例如
abs(a-b) <= 1e-4 * abs(b) + 1e-5。如果是 half 类型，阈值要放宽到1e-2级别。 - 如果 kernel 输出形状和参考实现形状不一致，直接判 REJECT。
- 如果输出有 NaN 或 Inf，直接判 REJECT，不用再跑更多测试。
动态轨的这一步能抓到相当多静态检查无法发现的“逻辑写对但索引写错”问题，尤其是对全局内存越界读写的算子来说，小尺寸测试通常就能暴露问题。
4.3 性能衰减检测：比基线慢两倍就直接打回
数值正确只是底线，性能才是工程上能否接受的关键。CMP-D V2.0-Lite 在数值校验通过后，还会做一个简单的性能测试。
方法也比较朴素：用同一份数据，在同一块 GPU 上，分别跑手写的基线 kernel 和 LLM 生成的 kernel，各跑 50 次取中位数。如果生成版本的中位耗时超过基线版本的 2 倍，就判为 REJECT。
这里我有意没有去追求复杂 profiling，原因有二：
- 一是它不需要额外的 profiling 工具，只要在 host 测时函数里包一层 CUDA event 就能完成；
- 二是在单算子场景下，耗时中位数已经能反映大多数问题。真正需要 ncu 或者 nsight 去定位的复杂问题，发生在代码已经成功进入工程之后，而不是前置拦截阶段。
实测下来，性能衰减规则大约能额外拦住 5%~10% 的代码。这个比例不算高，但价值很高，因为这一部分代码最容易骗过所有人，包括经验丰富的工程师。
5. 完整跑通一次检查：从 LLM 输出到 REJECT 的记录
5.1 一个典型坏例具体长什么样
下面这段代码，就是某个 LLM 在任务描述“实现一个 tanh 融合的 add 算子”时输出的 kernel。它看起来完整、头文件齐全、有注释，甚至还有 launch 配置。但拿 CMP-D 跑一遍，立刻原形毕露。
#include <cuda_runtime.h> #include <stdio.h> __global__ void tanh_add(float* a, float* b, float* out, int n) { int idx = threadIdx.x + blockIdx.x * blockDim.x; if (idx < n) { float x = a[idx] + b[idx]; out[idx] = tanhf(x); } } void launch_tanh_add(float* a, float* b, float* out, int n) { int threads = 256; int blocks = (n + threads - 1) / threads; tanh_add<<<blocks, threads>>>(a, b, out, n); cudaDeviceSynchronize(); }这段代码真正跑起来其实是没问题的，索引逻辑、边界判断、tanhf 调用都对。但请注意：这就是它的问题所在——太“常规”了，常规到掩盖了一个关键错误：任务要求的是“融合 add 算子”，也就是把add和tanh合并进同一个 kernel，而这个实现虽然确实融合了加法到 tanh 之前，但它的访存模式是每个线程只处理一个元素，完全没有向量化，也没有考虑实际硬件上的带宽瓶颈。如果 baseline 是手写 float4 向量化版本，这个版本大概率会慢三倍以上。
在这种“逻辑对但性能糟”的场景里，静态轨给的是 REVIEW 而不是 REJECT，但动态轨的 50 次耗时中位数会直接把它打回 REJECT。
5.2 运行 CMP-D 之后的输出结果
我在命令行里执行检查后，CMP-D V2.0-Lite 会输出类似下面的信息：
[static_review] + kernel tanh_add: global memory write detected. + index mapping matched launch config. + API blacklist: none. + shared memory check: none. verdict: PASS (enter dynamic track) [dynamic_check] + compile: OK (sm_80, 14.6ms) + numerical test: PASS atol=1e-4, rtol=1e-5 + perf test: baseline median=0.213ms, generated median=0.712ms (3.34x) verdict: REJECT reason: perf_regression_ratio_exceeded result: REJECT这里最直观的信号就是那个3.34x。如果我没有性能衰减检测，这个算子很可能会被人肉 review 放行，因为它的正确性是没问题的。但部署到线上之后，这种 3 倍以上的性能差距，会在整个模型推理链路里被放大。
5.3 另一个“勉强通过”的 REVIEW 样例
并非所有通过检查的代码都值得直接进工程。还有一类代码属于“能跑、但结构可疑”，CMP-D 会标成 REVIEW。比如下面这种写法：
__global__ void vec_copy(float* in, float* out, int n) { int idx = blockIdx.x * blockDim.x + threadIdx.x; if (idx < n) { out[idx] = in[idx]; } }这个 kernel 是每个线程复制一个元素，没有做float4向量化。性能上可能只比基线慢 1.5 倍，没触发 2 倍阈值，所以动态轨给 PASS。但我很清楚它的扩展性不好：当数据规模变大、显存带宽更高时，非向量化的访存模式会导致巨大的性能浪费。因此 CMP-D 会打上 REVIEW 标记，提醒我看到这类代码后，尽量再让 LLM 做一次“向量化优化”迭代。
REVIEW 存在的意义，就是不让脚本把所有决策都替人做了。它负责把风险暴露出来，而最终拍板还是靠人。
6. 这套系统在实际工作流里的落地姿势
6.1 接在 LLM 生成后端后面，而不是接在 CI 里
CMP-D V2.0-Lite 最初是我放在 CI 里的，想着每次提交直接自动跑。但后来我发现这个接入位置有一个致命问题：CI 的反馈链路太长了。等代码提交、流水线编译、测试、再返回结果，一个好一点的循环也要十分钟以上，效率上划不来。
更好的做法是把检查器直接挂在 LLM 生成后端后面，也就是让“生成”和“校验”成为同一套流水线里的两个连续步骤。每次模型输出一段代码，CMP-D 立即同步检查；如果 REJECT，就把错误原因和具体行号作为下一轮生成的辅助提示送回给模型，让它自己修改。这样一轮生成加检查的时间能控制在 20 秒左右，比走一遍 CI 快了一个数量级，也更适合高频迭代的开发节奏。
这个做法有一个额外的好处：错误信息能变成结构化数据。很多 LLM 的底层链路并不希望把一段自然语言丢回去，而 CMP-D 输出的reason: perf_regression_ratio_exceeded这样的标签，刚好可以直接拼到 prompt 里，实现“机器可读”的反馈。
6.2 用检查结果反哺 Prompt 或微调数据
被 CMP-D 标记为 REJECT 的代码样本，其实是非常宝贵的数据资产。我每周都会把这些坏例导出来，按原因分类，然后挑一部分整理成“反面对比”数据。这里说的不是拿去微调一个大模型，而是在 Prompt 层面构建 few-shot 坏例提示。具体做法是：在给 LLM 的上下文里加入“以下是之前生成的错误 CUDA kernel，你不应重犯这些错误”的示例。
做了这个动作之后，静态轨的拦截率又降了一点。为什么是“降”？因为模型看到坏例后主动避开了部分错误，生成的整体代码质量提高了，REJECT 率自然下降。但 CMP-D 作为最后一道安全网的定位不会变，因为大模型的输出依然有随机性，我们不能指望它永远不犯错。
6.3 误杀率与放行标准的取舍
拦截系统最怕的就是“误杀”。如果静态轨规则写得过严，大量正确代码被 REJECT，开发者很快就会失去耐心，直接绕过检查器。所以我在设计规则时特别注意两点：
- 第一，不要把“风格问题”升级为“正确性问题”。比如代码没有按公司规范加注释，这不该触发 REJECT。
- 第二，给规则分三档。致命错误（编译失败、数值错误、越界）直接 REJECT；可疑行为（无向量化、模板过度、共享内存边界存疑）标 REVIEW；风格提示（变量命名不清晰、缺少头文件注释）只记录，不参与结果判定。
我统计过一个月的数据，在接近 3000 条生成代码里，被 CMP-D 直接 REJECT 的约 2100 条，打上 REVIEW 的约 500 条，真正通过进入工程的约 400 条。这个通过率有一定波动，但拦掉七成左右垃圾代码的判断是成立的。
7. 常见问题与避坑速查
7.1 动态编译慢怎么办
动态轨最大的痛点是编译时间。常见做法有三个：
- 添加缓存，按源码 hash 缓存 nvcc 的 object 文件，重复生成不重复编译。
- 只编译
.cu文件，不要链接 final executable，能省掉大量链接时间。 - 对同一批代码做批量编译，把 10 个文件合到一条 nvcc 命令里，减少编译器的初始化开销。
如果实在急，还可以把-O2临时降成-O1，但要注意可能测不出部分性能问题，所以我一般只在功能验证阶段这么做。
7.2 正则误伤模板代码怎么办
如果你的代码大量使用模板，纯正则规则会非常容易误伤。比如匹配“显式调用cudaMalloc”的正则，可能会把模板函数里的cudaMalloc(ptr, size)匹配到，但这个模板其实只是做一层封装，并非真的在 device code 里分配显存。这种误伤会让人很崩溃。
我的解决方案是给静态轨加一份“豁免清单”。当某个 API 调用点处于模板函数内部，且该模板函数名命中白名单（比如safe_malloc、memory_pool_alloc），静态轨就不把它当作非法模式。更保险的方案是用 tree-sitter 的 C++ grammar 先拿到 AST 节点，再判断这个 API 调用所在的函数上下文。虽然复杂一点，但对模板代码的误伤会大幅降低。
7.3 到底要不要用 AST 替代正则
很多人一上来就希望用 clang AST 全量替代正则，这个想法我能理解，但实践下来并不划算。原因是：LLM 生成的 CUDA 代码往往包含宏、CUDA 特殊关键字和 C++ 模板混合，一个完整的 clang AST 解析需要引入额外的依赖，而且解析速度比正则慢一个数量级。
我的建议是分层处理：先用正则做快速淘汰，杀掉那些明确错误且成本很低的样本；只有正则拿不准、或者类型结构异常复杂的样本，才走 AST 通道。这个顺序让 95% 的代码都能走最快的那条路径，只有 5% 的疑难杂症才动用重型工具。
7.4 没有 NVIDIA GPU 时怎么办
CMP-D V2.0-Lite 的动态轨依赖 nvcc 和一块 NVIDIA GPU，所以不是所有机器都能直接跑。没有 GPU 的开发者有两个变通方案：
- 只跑静态轨。它能拦截掉 API 幻觉、索引空间不匹配、危险函数调用等大约三分之一的问题，依然有价值。
- 把编译动作放到远程 GPU 服务器上执行。CMP-D 支持把源文件通过网络传输到指定服务器，让远端执行 nvcc 和 kernel 测试，再把结果传回本地。
后续我还计划把这套检查逻辑改造成更通用的“算子生成质量闸门”，目标是让同一套规则也能支持 Ascend C 这类自定义算子开发环境。不过 Ascend C 的语法树和 runtime API 跟 CUDA 差异挺大，静态规则可能需要重写一批，动态轨的编译和运行机制也要跟着适配。这块我还在摸索，等跑通后会再写一篇。
7.5 一点个人体会
这套双轨脚本并没有用任何玄乎的算法，核心就是把“审查经验”变成了“自动化规则体”。我觉得最有价值的不是拦截了多少垃圾代码，而是它逼着我把自己对大模型生成代码的每一次骂街，都转化成了显式、可量化的检查项。每次加规则的时候，我都会问一句：这个错误下一次能不能被规则自动抓住？如果能，就值得固化；如果不能，那就说明我对这个错误的理解还不够深。
如果你也在用 LLM 写 CUDA 或者其他高性能算子，我建议先别急着收集各类提示词模板，而是花一个周末，把自己最近踩过的坑整理成一张规则表，再用脚本跑起来。哪怕一开始只有五条规则，也一定能省下后面无数个小时的返工时间。
