把“稀疏化”从论文里搬到 GPU 上为什么会卡住这么多人模型训练完剪枝之后理论计算量确实降下来了可是在 GPU 上一跑延迟反而没降多少甚至某些算子还变慢了。问题往往不是稀疏化本身而是没有一个匹配当前稀疏模式的高效内核。不同稀疏模式需要完全不同的访存策略、索引计算和并行划分方式手写 CUDA Kernel 不仅门槛高而且面对 2:4 结构化稀疏、Block 稀疏、随机稀疏等不同形态几乎是一套模式就要重写一套代码。SparseDitto 这个项目解决的就是这个矛盾用 LLM 驱动的 Agentic System把“从稀疏模式描述到 CUDA 内核实现”的过程自动化。它的核心思路不是让大模型直接硬生成一段代码而是构建一个会搜索、会生成、会编译、会根据错误信息反复修改的智能体闭环。读完这篇文章你会理解为什么稀疏内核难写SparseDitto 这类系统解决了哪些工程难点以及如果你想在自己的项目里尝试接入应该从哪个切入点开始。1. 这篇文章真正要解决的问题在 LLM 训练和推理的落地场景里越来越多人开始关注稀疏化——把权重矩阵里的零值跳过从而减少计算量和显存占用。但这里有个容易被低估的事实GPU 是为稠密连续计算设计的它擅长的是整齐划一的矩阵乘法而不是“遇到零就跳过”。如果你只是简单地把稀疏矩阵里的小值置零然后继续用稠密 Kernel 计算性能不会有任何提升因为 GPU 的线程仍然会被调度仍然会做多余的乘法甚至因为增加了判断分支而变慢。真正的收益来自专用 Kernel比如结构化稀疏中常见的 2:4 模式需要按固定 pattern 压缩数据、重排索引Block 稀疏则需要把非零块打包成紧凑的 tile再调整加载逻辑。这带来一个现实问题稀疏模式千变万化性能友好的 Kernel 必须跟着模式定制。传统方式下这是高性能计算工程师的工作需要懂 GPU 架构、懂内存访问模式、懂汇编级优化。对大部分算法工程师和 LLM 应用开发者来说这个门槛太高了。SparseDitto 的价值在于它把“Kernel 定制”这件事的入口从“手写 CUDA”变成了“描述稀疏模式 让 LLM Agent 去实现”。注意它并不是简单地把需求扔给大模型生成一次代码就结束而是围绕稀疏内核开发流程构建了一个能够迭代验证的智能体系统。这篇文章会带你拆解它的设计思路并给你一套可以借鉴的工程落地框架。2. 稀疏模式与 GPU Kernel 的基础关系要理解 SparseDitto 的价值得先明确一个基础问题为什么稀疏模式不同Kernel 就不同2.1 稀疏模式的常见分类稀疏模式本质上描述的是“非零元素在矩阵中的位置分布规律”。实际项目中常见的类型包括稀疏模式特征典型场景非结构化稀疏Unstructured非零元素位置完全随机无规律剪枝后的模型权重结构化稀疏Structured按行、按列或按固定块删除硬件友好的模型压缩2:4 结构化稀疏每 4 个连续元素中保留 2 个英伟达 Ampere 及以上架构的稀疏加速Block 稀疏Block Sparse非零元素按固定大小的块分布大模型 MoE、长序列注意力模式化稀疏Pattern-based非零位置遵循某种重复 pattern定制硬件或推理引擎同一个稀疏率下两种不同模式的非零元素分布完全不同Kernel 的数据加载方式、索引计算方式、线程块划分方式都不同。2.2 Kernel 针对稀疏做的核心优化一个稀疏 Kernel 的核心工作不只是“跳过零”而是围绕四点优化压缩存储把稀疏矩阵转换成 CSR、CSC、BSR 或自定义紧凑格式减少显存占用。索引映射在 Kernel 内部把逻辑坐标映射到压缩后的物理坐标避免每次计算都重新遍历查找表。访存合并让相邻线程访问相邻内存地址最大化内存带宽利用率。负载均衡稀疏模式下非零元素分布常常不均匀如果某些线程处理大量非零块而另一些线程在空转整体性能就会被拖垮。一个为 2:4 稀疏设计的 Kernel默认假设每 4 个元素恰好有 2 个非零数据可以预编码成固定格式而一个为 Block 稀疏设计的 Kernel则更关心如何把非零块打包成 tile然后对 tile 做矩阵乘。用错了模式轻则没有收益重则结果出错。这就是为什么“用一种 Kernel 适配所有稀疏模式”在工程上几乎行不通。SparseDitto 正是抓住了这个点它要按稀疏模式定制而不是提供一个万能 Kernel。3. 传统手工优化内核的三大痛点把时间线拉回到没有 SparseDitto 这类系统的时代。如果算法工程师想要一个支持某种稀疏模式的算子实际要经历什么3.1 痛点一手写 CUDA Kernel 的门槛太高写一个能跑、结果正确的 CUDA Kernel 已经不容易写一个性能能逼近 cuBLAS 或 cuSPARSE 的稀疏 Kernel 更是难上加难。你需要理解GPU 的线程层级Thread、Warp、Block、Grid共享内存的 bank conflict寄存器压力和 occupancy 的关系内存事务的合并规则不同架构比如 Ampere、Hopper、Ada的计算特性差异。这些知识通常分布在高性能计算课程、硬件白皮书和几年实战经验里不是算法工程师日常的技能栈。3.2 痛点二模式的组合数量爆炸稀疏模式本身就有多种维度结构非结构化、2:4、Block、数据类型FP16、BF16、FP8、硬件架构A100、H100、L40S、消费级显卡、算子类型GEMM、GEMV、注意力计算。把这些维度组合起来可能的场景数量是几十甚至上百种。传统做法是每个组合都单独写一个 Kernel然后做性能调优。这个工作量在人力上基本不可持续很多团队最终只在最常用的几种模式上做了优化其他场景退回到稠密计算。3.3 痛点三迭代验证环节非常耗时即使写出一版 Kernel也不意味着结束。你要对比正确性、测试不同 block size、不同 unroll 参数、不同访存顺序。这种调优循环本质上是重复性劳动却消耗了大量工程师时间。SparseDitto 这类系统能出现正是因为这三个痛点的叠加LLM 可以读文档、写代码、理解报错信息Agent 可以自动执行“生成-编译-运行-反馈-修改”的循环而这个循环正好对应手工优化 Kernel 的完整流程。4. SparseDitto 的核心思路LLM 智能体如何定制 KernelSparseDitto 的基本思路可以概括为一句话把“稀疏模式”作为输入让 LLM 智能体自动完成 Kernel 的搜索、生成、编译和验证最终输出一个针对该模式定制的高性能 GPU Kernel。4.1 它不只是“用 AI 写 CUDA”很多人看到这个标题第一反应是“这不就是用大模型生成 CUDA 代码吗”。如果是这样那意义确实有限——现在很多大模型都能写 CUDA但写出来的代码经常有 bug性能也不稳定。SparseDitto 更关键的设计是Agentic System智能体系统。它不是单次生成而是形成一个闭环理解任务解析输入的稀疏模式描述和算子需求搜索参考从已有的内核模板、开源算子库、相关论文代码中寻找可借鉴的方案生成候选代码基于搜索和推理结果生成 Kernel 源码编译验证自动编译捕获编译错误和运行错误反馈迭代把错误信息、profile 结果反馈给 LLM让它修改代码性能评估对比基线性能决定是否继续调优。这个闭环的本质是把“有经验的 Kernel 工程师”的日常工作流程自动化。工程师写 Kernel 时会查资料、参考已有实现、编译报错后看错误信息、修改、再 profileSparseDitto 让 LLM Agent 复刻了这个过程。4.2 为什么需要“Agentic”而不是“一次生成”代码生成任务里单次生成的失败率非常高。尤其在 CUDA Kernel 这种对细节极度敏感的领域可能只是 block 尺寸设置不合理或者索引计算差了一个偏移就导致结果错误或性能暴跌。Agentic 系统最大的优势是可以利用反馈信号持续改进。编译器的报错信息是明确的反馈运行结果和参考结果的差异也是反馈性能 Profile 的数据更是反馈。LLM 具备理解这些反馈并修改代码的能力这才是它能在 Kernel 定制场景发挥作用的关键。4.3 SparseDitto 在技术栈中的定位从技术栈角度看SparseDitto 位于三者的交汇点底层CUDA / GPU Kernel 编程技术中层稀疏矩阵计算与存储格式上层LLM Agent 的规划、工具调用与自我修正能力。它不是一个通用代码生成工具而是聚焦在“稀疏模式适配”这一垂直场景。正因为场景垂直任务目标明确Agent 的搜索空间可控反馈信号清晰所以比通用的“AI 编程助手”更容易产生实际价值。5. 系统工作流拆解从模式描述到可用内核下面我们来拆解一个符合 SparseDitto 设计思路的典型工作流。这里需要说明不同版本的项目实现细节可能有差异以下流程是基于智能体系统和稀疏内核开发的通用逻辑整理的示意框架具体接口以项目官方仓库为准。5.1 输入稀疏模式描述用户不是直接写 Kernel而是描述自己的需求。一个最小描述可能包括稀疏模式类型如 block_sparse矩阵维度与块大小数据类型目标 GPU 架构算子类型如 matmul。5.2 内核生成与搜索Agent 拿到描述后先做两件事从本地或远程代码库中检索相似的 Kernel 实现或模板结合 LLM 自身的训练知识生成候选 Kernel。这里有个容易踩坑的地方不能让 Agent 一上来就写一个超大文件而是应该引导它先生成最小可用版本再逐步优化。否则问题会集中在“代码太长、错误太多、难以定位”上。5.3 编译与执行验证生成的 Kernel 不能直接信任。SparseDitto 类系统会做这几步验证编译检查在模拟数据上运行与稠密参考结果对比正确性记录运行时间和性能指标。如果任何一步失败错误信息会作为下一轮生成的输入。5.4 迭代优化Agent 根据错误信息和性能数据修改 Kernel重复编译、运行、评估。这个过程可能循环很多轮直到正确性通过且性能达到目标。5.5 输出可用内核与配套信息最终输出包括CUDA Kernel 源码调用示例编译和运行命令性能测试报告使用该 Kernel 的注意事项。从用户视角看整个流程从“我需要一个支持 2:4 稀疏的矩阵乘 Kernel”到“拿到了一个经过验证的 Kernel”中间的人工介入被降到了最低。6. 环境准备与接入思路虽然 SparseDitto 项目本身的具体安装方式需要参考仓库文档但从这类系统的通用依赖来看环境准备通常包括以下部分。如果你要在一个 Kernel 定制项目中尝试类似思路需要准备一台可用 GPU 的 Linux 服务器CUDA 工具链nvcc 编译器Python 3.8 以上环境LLM API 或本地模型推理服务支持 Agent 工具调用的 LLM 框架或直接用 OpenAI Function Calling 风格接口。下面的配置示例是示意性质的目的是展示“如何把一个稀疏模式描述传给智能体系统”不是官方 API。{ task: sparse_matmul_kernel, sparsity_pattern: { type: block_sparse, block_size: [32, 32], sparsity_ratio: 0.8 }, matrix_shape: [4096, 4096], dtype: fp16, gpu_arch: sm_90, baseline: dense_cublas }这段描述告诉系统客户需要一个稀疏矩阵乘 Kernel稀疏模式是 32x32 的 Block 稀疏稀疏率约 80%矩阵维度 4096x4096数据类型 FP16目标架构是 Hoppersm_90要和稠密 cuBLAS 做性能对比。在设计自己的 Agent 时输入描述越结构化Agent 的理解成本就越低。这里真正值得留意的是稀疏模式描述字段。如果缺失这个字段Agent 只能靠猜生成的 Kernel 很可能与你的数据分布完全不匹配。7. 完整示例一个最小 Agentic Kernel 生成循环这一段给出一个简化但结构完整的 Python 示意代码展示 Agent 如何控制“生成 - 编译 - 反馈 - 再生成”的循环。# 文件路径agent_loop_demo.py # 说明这是一个简化版 Agentic 内核生成流程示意不是 SparseDitto 官方实现 import subprocess import json from typing import Optional def generate_kernel_with_llm(prompt: str, error_feedback: Optional[str] None) - str: 调用 LLM 生成或修改 CUDA Kernel 源码。 实际项目中使用你选择的 LLM API 或本地模型。 messages [ {role: system, content: 你是一名 CUDA Kernel 优化专家擅长稀疏矩阵计算。}, {role: user, content: prompt}, ] if error_feedback: messages.append({role: user, content: f上一版代码存在以下问题请修复\n{error_feedback}}) # 这里替换为真实 LLM 调用 # response llm_chat(messages) # return response[kernel_code] return __global__ void placeholder_kernel() {} def compile_kernel(source_file: str) - tuple[bool, str]: 使用 nvcc 编译 CUDA 源码。返回是否成功及输出信息。 result subprocess.run( [nvcc, -archnative, -o, test_kernel, source_file], capture_outputTrue, textTrue, ) return result.returncode 0, result.stderr def run_kernel_test(binary: str) - tuple[bool, str]: 运行测试程序检查结果正确性与性能。 result subprocess.run([f./{binary}], capture_outputTrue, textTrue) return result.returncode 0, result.stdout def main(): task_desc { operator: sparse_matmul, pattern: block_sparse_32x32, shape: [4096, 4096], dtype: fp16, } prompt f请为以下任务编写一个 CUDA Kernel\n{json.dumps(task_desc, indent2)} error_feedback None max_iterations 5 for i in range(max_iterations): print(f 迭代 {i 1} ) kernel_code generate_kernel_with_llm(prompt, error_feedback) with open(generated_kernel.cu, w, encodingutf-8) as f: f.write(kernel_code) ok, compile_output compile_kernel(generated_kernel.cu) if not ok: error_feedback f编译错误\n{compile_output} print(error_feedback) continue ok, run_output run_kernel_test(test_kernel) if not ok: error_feedback f运行错误\n{run_output} print(error_feedback) continue print(编译和运行成功进入性能评估阶段。) print(运行输出) print(run_output) break else: print(达到最大迭代次数需要人工介入。) if __name__ __main__: main()这段代码展示了一个最朴素的 Agent 循环。几个关键设计点错误反馈是循环的核心每一轮生成的代码都会被编译编译错误会被反馈给下一轮生成形成自我修正。最大迭代次数保护避免 Agent 陷入无限循环。真实场景中可以设置超时时间和最大迭代次数。模块化设计LLM 调用、编译、运行、反馈解析各自独立方便替换和调试。真实系统中每一轮迭代结束后还需要把性能 Profile 数据也加进反馈里。比如 kernel 运行时间、占用率、访存吞吐量这些都可以作为下一轮优化方向的参考信息。7.1 更贴近实际的反馈信息格式为了让 LLM 能有效利用反馈信息反馈内容不应该只是一段字符串保持结构化会更容易帮助模型定位问题。{ stage: compile, status: failed, error_type: compile_error, error_message: generated_kernel.cu:42: error: invalid argument type, suggestion: 检查第 42 行索引计算中 float4 的类型转换 }结构化反馈的另一个好处是可以让 Agent 基于错误类型选择不同的修复策略编译错误直接修改语法和类型运行错误则先检查越界和同步问题性能不达标则优化访存和并行划分。8. 运行验证与效果评估SparseDitto 这类系统最终输出的不只是代码还应该包括一套验证记录。用户在拿到 Kernel 后需要自己也会做基本确认。8.1 验证正确性正确性验证不能只看“跑通了”而是要对比结果。最直接的方式是用同样的输入数据分别运行稀疏 Kernel 和稠密参考实现再比较输出矩阵之间的误差。# 运行稀疏 Kernel 测试程序它会输出与稠密参考结果的最大误差 ./test_kernel --check-error # 预期输出类似 # max_abs_error 0.000732 # Test PASSED稀疏 Kernel 因为浮点累加顺序不同误差不必为 0但要在可接受范围内。FP16 计算场景下相对误差在 1e-2 量级通常是可接受的FP32 场景则应当更低。8.2 验证性能收益性能验证的基准不是“比稠密版本快”就算成功应该关注与稀疏率匹配的“理论加速比”之间有多大差距与 cuSPARSE 等成熟库的同等功能 Kernel 相比如何在不同矩阵规模下是否保持稳定。一个简单的性能测试方法是让 Kernel 执行多次取平均时间再与基线对比。# 运行性能测试 ./test_kernel --benchmark --iterations 100 # 输出示例 # Sparse Kernel: 0.342 ms # Dense Baseline: 1.875 ms # Speedup: 5.48x注意这里不要期望每次都能达到理论加速比。访存开销、索引计算、block 调度等因素都会影响最终效果。如果加速比远低于预期最值得看的地方是访存是否连续、是否有大量 bank conflict、以及 warp 内部是否出现分支发散。8.3 失败时的第一步排查顺序如果运行失败或者性能不达标建议按这个顺序排查先看编译日志有没有类型不匹配、未定义符号再看运行日志有没有 CUDA error、非法内存访问检查数据规模与 block 配置线程数是否覆盖了矩阵所有元素检查索引映射稀疏索引计算是否正确最后才看性能数据不要一上来就调优先保证正确。9. 常见问题与排查思路结合稀疏 Kernel 开发中常见的实际问题这里列出几个高频场景。问题现象可能原因排查方式解决方案编译报错invalid argument typeLLM 生成的索引变量类型不一致查看报错行附近的类型定义在 Prompt 中明确要求统一使用int或size_t并附上类型检查规则运行时报illegal memory access稀疏索引计算越界使用compute-sanitizer检查越界访问增加边界条件判断校验压缩后的索引范围Kernel 能跑但结果全为 0稀疏数据加载逻辑错误未正确映射非零块打印部分中间索引值对比 CSR 偏移表检查 row offset 和 block index 的计算公式性能提升远低于理论值访存不连续或 bank conflict 严重用 Nsight Compute 查看内存吞吐量调整数据布局使用共享内存做 tile 重排Agent 反复修改后仍编译失败反馈信息不足或 Prompt 约束不够检查反馈内容是否包含完整错误信息把编译错误、源码位置和当前 Kernel 代码一并传给 LLM不同矩阵规模下性能波动很大Kernel 是针对固定 block 大小优化的测试多组 shape观察性能曲线设计自适应 block 参数或在描述中增加 shape 范围这些问题的核心规律是稀疏 Kernel 的 bug 大多不在算法逻辑而在索引计算和内存访问。Agent 系统迭代时需要把 sanitizer 和 profiler 的输出纳入反馈否则 LLM 很难发现这类隐藏在运行期的问题。10. 最佳实践与工程建议如果你准备在自己的项目里尝试“LLM Agent 定制 GPU Kernel”这条路下面是几条有实际价值的建议。10.1 把任务描述做结构化不要用自然语言长段落描述任务而是像上文那样使用 JSON 等结构化格式。结构化描述能显著降低 LLM 的误解率尤其在稀疏模式、数据类型、GPU 架构等关键参数上。字段必须包括算子类型、稀疏模式、矩阵维度、数据类型、目标架构、基线实现、容错范围。10.2 设计清晰的反馈回路Agent 系统的上限取决于反馈回路的质量。如果只反馈“运行失败”LLM 很难定位问题。好的反馈应该包含编译错误原文、报错位置、运行日志、期望结果与实际结果的差异、性能数据。把这些信息组织成结构化文本再传给下一轮生成。10.3 设置人工介入阈值不要把 Agent 完全放在无人值守模式。合理的设计是迭代超过 N 次仍失败或者性能达不到预设目标就暂停并通知人工。否则 Agent 可能在错误方向上浪费大量 GPU 时间和 LLM 调用成本。10.4 保护 GPU 资源与成本跑 Agent 循环会消耗大量 GPU 时间尤其是多次编译、运行、profile 的循环。建议在小型数据上先验证正确性再上真实数据限制单次任务的最大迭代次数设置单轮编译和运行的超时时间对 LLM 调用增加预算控制。10.5 关注安全边界与稳定性在自动化生成代码的流程中生成的代码虽然在隔离环境编译执行接入生产前仍然要做代码审查。特别是不要直接在生产环境跑 Agent 生成的未经审查的 Kernel对 Kernel 的调用参数做严格校验保留基线实现的回退路径一旦新 Kernel 出问题可以快速切换。10.6 维护一份验证基线库每次 Agent 成功生成一个 Kernel都应该把对应的任务描述、生成代码、验证结果、性能数据保存下来。这相当于为一个团队积累“稀疏 Kernel 经验库”后续 Agent 可以从这些历史案例中检索相似实现少走弯路。11. 总结与后续学习方向SparseDitto 真正想解决的问题不是“让 AI 写几行 CUDA”而是“让 GPU 内核定制从高门槛的手工劳动变成可复用的自动化流程”。它的核心资产是那个 LLM Agentic System 的闭环描述稀疏模式、搜索参考、生成内核、编译验证、错误反馈、迭代优化。这套思路对任何与高性能计算相关的领域都有借鉴价值。如果你想继续深入建议从这几个方向入手学习稀疏矩阵的存储格式和基础 Kernel 实现理解 CSR、BSR、2:4 稀疏的基本优化思路熟悉 CUDA 的性能调试工具比如 Nsight Compute 和 compute-sanitizer它们是 Agent 反馈回路里最有效的“眼睛”尝试自己构建一个最小闭环的 Kernel 生成 Agent从一个简单的稀疏向量加法或稀疏矩阵乘开始不需要一开始就挑战大规模复杂算子关注多智能体协作方向比如用一个 Agent 负责生成代码、另一个 Agent 负责审查性能效果通常比单个 Agent 反复迭代更好。最后提醒一句这类工具解决的是“Kernel 定制”的效率问题但 GPU 底层原理和稀疏计算基础仍然无法绕过。工具的自动化程度越高越需要使用者在关键时刻能判断生成结果到底对不对、值不值得用。把这篇文章收藏起来在你下次被稀疏算子性能卡住的时候再翻出来对着做一遍应该能省下不少时间。