GPU 性能优化做到后期很多人会撞上同一个尴尬CUDA 代码里的 kernel 已经写得很精简甚至连循环展开都手动做过但最终性能仍然不理想。真正决定执行效率的并不是 C 或 CUDA 源码而是经过 ptxas 编译后真正运行在 NVIDIA GPU 上的 SASS 机器码。SASS2MLIR 正是从这个最底层切入的方案先把 SASS 指令提升成 MLIR 中间表示再用 MLIR 的优化规则重新整理和生成指令。它的公开实验结论在 NVIDIA GPU 上能带来约 20% 到 100% 的性能提升虽然这个数字依赖应用场景但足够说明把机器码再当作可优化 IR 处理是一条值得掌握的优化思路。本文从 NVIDIA 编译链路起步讲清楚 SASS2MLIR 的机制、构建方法、运行流程、基准测试和落地注意事项。1. 为什么面向 SASS 做优化先看懂 NVIDIA GPU 的编译链路1.1 CUDA 到 SASS 的完整编译路径学习 SASS2MLIR 之前需要先理清一个 CUDA 内核从源码到硬件执行的完整链路。一个.cu文件经过nvcc编译时并不会直接被翻译成 GPU 最终执行的机器码而是先经过多层表示CUDA C / Triton | v NVVM IR / LLVM IR | v PTX虚拟指令集与具体 GPU 架构解耦 | v SASS对应具体 GPU 微架构的机器码 | v NVIDIA GPU 硬件执行在工程实践中生成这几层表示通常使用下面的命令。假设有一个名为saxpy.cu的简单内核# 生成 PTX nvcc -archsm_80 -ptx -o saxpy.ptx saxpy.cu # 生成 cubin内部已经包含 SASS nvcc -archsm_80 -cubin -o saxpy.cubin saxpy.cu # 从 cubin 中提取 SASS 汇编 cuobjdump -sass saxpy.cubin saxpy.sassPTX 是 NVIDIA 定义的虚拟指令集它的存在是为了让代码可以在不同代际 GPU 上迁移。同一个 PTX 可以被 ptxas 编译成 Ampere、Hopper、Ada 等不同架构的 SASS。SASS 则是真正交给 GPU 执行的指令每条指令都对应具体的寄存器、内存操作和调度约束。理解这条链路的意义在于PTX 并不是最终性能来源SASS 才是。很多时候两个看起来不同的 CUDA 写法可能编译出几乎相同的 SASS而另外一些看似等价的写法由于寄存器分配和指令顺序不同性能差异可能达到百分之几十。1.2 传统编译优化为什么会在 SASS 层失效ptxas在做 SASS 生成时确实会做寄存器分配、指令调度、循环展开、内存访问合并等优化。但它有一个天然限制编译单元是单个 kernel。ptxas 看不到两个 kernel 之间的数据复用关系也看不到同一个 kernel 在不同调用上下文中的调用模式。举一个典型场景先写一个 kernel 对矩阵做分块转置再写一个 kernel 做矩阵乘。如果由人工优化可以把转置结果直接留在共享内存中复用减少一次全局内存写回和读取。但在传统的 CUDA 编译链路里这两个 kernel 是分开编译、分开运行的ptxas 没有任何机会做这种跨 kernel 优化。另一个问题在于一旦 SASS 生成并打包进 cubin常规编译器不会再对它做进一步优化。GPU 驱动虽然可以在加载 cubin 时做部分 JIT 处理但绝大多数时候 Vcache 和 warp 调度器拿到的就是 ptxas 产出的那条指令序列。应用层信息、数据布局信息、tensor core 使用意图都在编译过程中丢失了。1.3 SASS2MLIR 想切入的优化空间SASS2MLIR 的核心思路是把“已经生成的 SASS”再次提升到一个可优化的中间表示MLIR。它把每条 SASS 指令、寄存器、控制流和数据流重新建模成 MLIR 操作然后运行一系列 MLIR 优化 passes最后生成新的 SASS。这个思路背后有一个关键判断机器码虽然在硬件上已经可执行但它并不等于最优代码。只要能把机器码还原成一个保留了语义、可分析、可变换的 IR就存在继续优化空间。MLIR 正好提供了这样一个生态它有丰富的 dialect 机制、完善的 pattern rewrite 框架、pass 管理器和目标代码生成工具链。性能提升正是从这个环节开始的原来只能被 ptxas 处理一次的 SASS现在可以被反复分析、变换、验证。这样编译器就能基于真实执行的指令重新做指令融合、布局优化、同步优化和调度最终拿回一部分被早期编译决策放弃的性能。2. SASS2MLIR 的核心机制把机器码抬回可优化表示2.1 SASS 与 MLIR 的关系SASS 是 NVIDIA GPU 的底层指令集它的指令往往带有明确的目标架构特征。例如FFMA是浮点乘加LDG是加载全局内存STS是存储共享内存BAR.SYNC是同步屏障。这些指令本身是稳定的但一个完整的 kernel 会包含大量指令之间的依赖关系、寄存器生命周期和线程同步点。MLIR 是一种多级中间表示框架。它可以描述从高层循环结构到底层机器指令的完整过程。SASS2MLIR 通常会定义一个用于表示 SASS 指令的 dialect比如sassdialect每条 SASS 指令都被建模为一个 MLIR operation操作数对应寄存器或立即数结果对应新的寄存器。这样做的好处在于MLIR 自身就带有 SSA 形式、控制流图、数据流分析等基础设施不需要重新发明轮子。只要 SASS 指令能正确落到 dialect 中后续的依赖分析、死代码消除、指令融合就都可以复用 MLIR 生态里的 pass。2.2 SASS2MLIR 工作流程一次典型优化过程可以分为四个阶段输入收集从 cubin、fatbin 或纯文本 SASS 文件中读取未优化的指令序列。反汇编与建模把二进制 SASS 或汇编文本解析为内存中的指令对象再映射到 MLIR dialect operation。优化变换在 MLIR 上运行 canonicalize、loop unroll、layout optimize、barrier elimination、instruction fusion 等 pass。代码生成把优化后的 MLIR 重新翻译成 SASS并尝试封装回 cubin 或供运行时加载。下面是一个简化流程示意图saxpy.cubin | v cuobjdump / 反汇编工具 | v SASS 指令序列 | v SASS2MLIR (lift to MLIR) | v MLIR SASS dialect | v MLIR optimization passes | v Optimized MLIR | v SASS backend | v opt.sass / opt.cubin这套流程有一个明显优势每个阶段都可以独立验证。你可以在反汇编后检查指令是否完整在 MLIR 提升后检查控制流是否正确在优化后检查指令序列是否仍然等价在最终生成后跑一遍正确性测试。对编译器开发来说这种可组合性非常重要。2.3 MLIR 中如何表达 GPU 指令为了直观理解 SASS 提升到 MLIR 后的形态可以看一个简化示例。假设有一段 SASS 风格指令LDG.E R2, [R0.64] FFMA R4, R2, R6, R4 STG [R8.64], R4 BAR.SYNC 0这些指令在 MLIR 中可能被建模成如下形式%val sass.ldg %ptr : memreff32 %acc sass.ffma %val, %weight, %acc : f32 sass.stg %ptr, %acc : memreff32 sass.bar.sync 0这里的%ptr、%val、%acc都是 SSA value编译器可以很方便地追踪依赖关系。例如sass.ffma依赖%val和%acc如果后续发现%acc在某个分支中没有被使用就可以做死代码消除。这一点在纯汇编文本中很难自动完成因为寄存器不断被复用别名分析非常复杂。需要说明的是不同版本和不同实现的 SASS2MLIR 可能选择不同的 dialect 名称和操作名称。上面这段只是用于说明思路实际工具输出应当以对应版本的--help和示例输出为准。真正落地时你大概率需要翻一遍sass dialect的 operation 定义理解每条指令映射到 MLIR 后的 operand 排列和类型约束。2.4 性能提升的关键点SASS2MLIR 的性能提升来自多个方向的综合效果。根据目前公开实验和编译器优化的常见做法主要可以归纳为以下几点优化方向核心思想典型收益来源指令融合把多条简单指令合成为一条等效指令减少取指和执行开销FFMA、FMULFADD、reduction 融合数据布局优化调整共享内存、全局内存访问顺序减少 bank conflict 和 cache miss分块、数组转置、swizzle 模式同步和 barrier 消除在保证语义的前提下减少不必要的 warp 同步BAR.SYNC冗余消除寄存器生命周期优化缩短寄存器占用窗口降低寄存器峰值压力降低 spill提高占用率循环和头尾处理优化把循环展开、predication 和边界分支处理得更高效大循环场景收益明显要特别说明的是这些优化并不总是同时生效。计算密集型 kernel 更容易从指令融合和 tensor core 映射中获益访存密集 kernel 则更容易从布局优化和访问合并中获益。如果一个 kernel 已经接近硬件理论极限那么 SASS2MLIR 能提供的空间通常较小甚至没有明显变化。3. 环境准备与构建搭一个能跑通 SASS2MLIR 的环境3.1 硬件和驱动要求SASS2MLIR 的输入和输出都依赖 NVIDIA GPU 生态因此硬件环境首先要有一条可用的 NVIDIA GPU。这里建议优先选用官方支持较好的服务器 GPU例如 NVIDIA A100、V100、H100 或 RTX 30/40 系列。架构越新SASS 指令集越复杂对工具的兼容性要求也越高。驱动层面的要求包括检查项建议要求验证方式GPU 是否可见至少能运行 CUDA 程序nvidia-smi出现 GPU 列表驱动版本与 CUDA Toolkit 版本匹配nvidia-smi输出 CUDA VersionCUDA Toolkit建议 11.8 以上nvcc --version编译器工具链C17 以上编译器g --version构建工具CMake 3.20 以上、Ninjacmake --version在 WSL 或容器环境中需要额外确认 GPU 是否透传成功。如果遇到类似Failed to initialize NVML: GPU access blocked by the operating system的报错说明宿主机的 GPU 权限没有正确传递给虚拟机或容器。这时要先检查 WSL 的 GPU 支持配置或确认容器是否安装了正确的容器工具包再回到 SASS2MLIR 构建流程。3.2 依赖清单构建 SASS2MLIR 通常需要以下几类依赖依赖用途注意事项LLVM / MLIR提供 MLIR 框架和 pass 管理需包含 NVPTX 后端CUDA Toolkit生成 PTX、SASS运行 cuobjdump版本要和工具匹配Python 3部分构建脚本或测试脚本使用推荐 3.10 以上Git拉取源码和子模块确保网络可访问Ninja / CMake编译构建建议用 Ninja 加速LLVM/MLIR 是最大的外部依赖。SASS2MLIR 很可能需要和特定版本的 LLVM 对齐因为 MLIR API 在不同版本之间有较大变化。不要默认使用任意最新版本建议先查看项目 README 或CMakeLists.txt中要求的 LLVM 版本。3.3 编译 LLVM/MLIR如果项目要求使用私有修改版 LLVM通常官方仓库会提供说明。如果只需要系统级 LLVM/MLIR可以按下面的典型流程编译git clone --depth 1 -b llvmorg-18.1.8 https://github.com/llvm/llvm-project.git cd llvm-project cmake -G Ninja -S llvm -B build \ -DLLVM_ENABLE_PROJECTSmlir \ -DLLVM_TARGETS_TO_BUILDNVPTX;X86 \ -DLLVM_ENABLE_ASSERTIONSON \ -DCMAKE_BUILD_TYPERelease cmake --build build --target mlir-opt这个步骤会生成mlir-opt、mlir-translate等核心工具。编译 MLIR 本身可能需要较长时间建议使用-j指定并行度或者在内存充足的机器上执行。3.4 构建 SASS2MLIR 本体拿到 SASS2MLIR 源码后通常构建流程与普通 CMake 工程类似。由于不同时间点的仓库结构可能不同这里给出一个通用流程具体以官方 README 为准git clone sass2mlir-repo cd sass2mlir cmake -G Ninja -S . -B build \ -DMLIR_DIRllvm-project/build/lib/cmake/mlir \ -DLLVM_DIRllvm-project/build/lib/cmake/llvm \ -DCMAKE_BUILD_TYPERelease cmake --build build需要说明的是如果构建时找不到MLIR_DIR通常说明传入的路径不正确或者 target 没有包含mlir项目。另外SASS2MLIR 的项目结构里可能包含python/、include/、lib/等多个目录构建完成后的二进制可能分布于build/bin或类似位置可以先检查一下输出目录。3.5 验证安装构建完成后先运行版本命令确认可执行文件能正常启动sass2mlir --version如果工具提供了自测用例继续运行测试ctest --test-dir build在进入真实优化前还可以准备一个最小的 SASS 输入文件尝试完成一次“SASS - MLIR - SASS”的往返。即便第一次生成的 SASS 不一定更快只要链路能走通就说明环境依赖已经对齐。4. 用 SASS2MLIR 跑通一次优化流程4.1 准备一个 CUDA 内核为了观察优化效果先准备一个计算型 kernel。下面以一个简单的 SAXPY 为例便于核验正确性// saxpy.cu __global__ void saxpy(float* y, const float* x, float a, int n) { int i blockIdx.x * blockDim.x threadIdx.x; if (i n) { y[i] a * x[i] y[i]; } }这个 kernel 足够简单但已经包含全局内存加载、乘加运算、条件分支和全局内存写回。后面可以用它熟悉完整流程再替换成真实项目中的 GEMM 或卷积 kernel。4.2 生成 PTX 和 SASS使用 nvcc 生成 cubinnvcc -archsm_80 -cubin -o saxpy.cubin saxpy.cu再从 cubin 中导出 SASS 汇编cuobjdump -sass saxpy.cubin saxpy.sass检查saxpy.sass内容确认其中包含多条LDG、FFMA、STG、BRA等指令。这里要注意一点如果nvcc生成的 cubin 同时包含了多个架构的 SASScuobjdump可能会打印多个函数版本建议在实际优化前确认选择的目标sm_架构避免混淆。4.3 SASS 转 MLIR将 SASS 文件或 cubin 输入给 SASS2MLIR得到一个 MLIR 文件。命令可能类似sass2mlir saxpy.sass -o saxpy.mlir如果你的版本支持直接从 cubin 读取也可以sass2mlir saxpy.cubin -o saxpy.mlir打开saxpy.mlir后你会看到类似 SSA 形式的 operation。如果这一步出现 “unknown opcode” 或 “unsupported instruction”说明工具当前不支持该架构的某条指令或者 SASS 输入中包含了工具未建模的扩展指令。这类问题在后面的第 6 节会专门讨论。4.4 应用优化 passes 与生成代码优化阶段通常需要独立的驱动工具类似mlir-opt的使用方式。假设工具提供sass2mlir-opt可以检查可用 passsass2mlir-opt --help | grep sass然后选择常见优化组合sass2mlir-opt saxpy.mlir \ -canonicalize \ -sass-fuse-ffma \ -sass-layout-optimize \ -o saxpy_opt.mlir这里的 pass 名称是示例性写法不同版本的 pass 名称可能完全不同。执行前应当先通过--help查看 tool 提供的实际 pass 列表再组合你需要的优化。canonicalize属于 MLIR 通用 pass负责规整产物sass-fuse-ffma、sass-layout-optimize则是面向 SASS 的领域优化。4.5 生成 SASS 并替换优化后的 MLIR 需要翻译回 SASS。这一步通常由后端工具完成sass2mlir-backend saxpy_opt.mlir -o saxpy_opt.sass如果工具支持打包 cubin还可以直接生成新的 cubinsass2mlir-backend saxpy_opt.mlir --cubin -o saxpy_opt.cubin得到新的 SASS 后需要验证两件事第一新 SASS 是否仍然可以正确运行第二新 SASS 是否真的更快。对于第一点可以把新的 SASS 通过工具链嵌入 cubin再用自定义的 CUDA loader 加载执行。如果当前工具不支持直接替换 fatbin也可以先用cuobjdump对比新旧指令序列确认没有明显语义错误再编写一个小的加载测试。5. 性能收益怎么量Benchmark 方法和结果解读5.1 基准测试环境配置谈到性能提升必须要有可信的测量方法。如果环境设置不一致20% 到 100%的结论很容易被复现为负优化。建议在基准测试前做固定配置关闭 GPU 动态超频固定 SM 频率和显存频率。使用同一个 NVIDIA 驱动版本和 CUDA Toolkit 版本。测试代码使用相同编译参数只切换 SASS 来源。多次运行取中位数并记录最小值和最大值。避免在 GPU 被其他任务占用的节点上做对比。在 Linux 环境下可以通过nvidia-smi -lgc锁定 GPU 时钟但不同型号的 GPU 支持范围不同。也可以直接通过nvidia-smi确认当前运行状态。5.2 指标选择基准测试指标不能只看时间。以下指标适合结合分析指标含义优化评估建议Kernel 执行时间内核在 GPU 上的耗时最直接但受噪声影响大GFLOPS / TFLOPS每秒浮点运算次数适合计算密集 kernel内存吞吐每秒有效读写字节数适合访存密集 kernelSM 占用率warp 占用的硬件资源比例反映并行度是否充分寄存器溢出量local memory 访问次数溢出通常导致明显回退指令数执行的总指令数量指令融合是否能带来收益推荐同时使用ncu或nsys这类性能分析工具观察 stall 原因和 warp state 分布。只看 kernel 时间可能会掩盖真正原因。5.3 如何判断提升是真实的判断性能提升真实存在至少要做以下几个检查第一正确性检查。优化后的 SASS 计算出来的数值必须与原始版本一致。如果目标 kernel 有浮点误差就要明确误差阈值不能让优化后的结果偏差过大。第二多次重复。性能测试至少执行 20 次以上观察方差。如果中位数和最小值都稳定提升结论才更可信。第三对照分析。除了执行时间还要看寄存器数、指令数、spill 次数、cache 命中率是否有变化。比如指令数下降了 15%但访存时间不变最终时间可能没有明显下降。第四多规模验证。不要只测试一个 block size 和输入 size。一个 kernel 在 1024x1024 矩阵上有提升不代表在 8192x8192 上也一样有提升。5.4 为什么有些场景可能没有提升甚至退化SASS2MLIR 并不是所有 kernel 的银弹。以下几种情况需要特别留意场景可能原因建议处理kernel 已接近理论峰值没有可压缩的冗余接受现状不要过度优化访存带宽成为瓶颈指令融合不改变访存模式优先优化布局和数据复用优化后寄存器溢出新 pass 增大了寄存器生命周期调整并行度或 pass 组合工具版本过旧不支持当前 GPU 架构某些指令更新工具或降级目标架构小 kernel 频繁启动启动开销掩盖了指令优化收益配合 CUDA Graph 减小启动开销标题中的~20% to 100%属于一些优化场景下的测量结果不代表所有内核都能达到。实际项目中更常见的是“某些 kernel 提升 20%某些没有变化个别下降 5%”。带着这个预期去评估才不会被单一 benchmark 误导。6. 常见问题与排查路径6.1 驱动和 CUDA 版本不匹配现象nvcc可以编译但运行nvidia-smi显示的 CUDA Version 低于编译要求或者明明安装了 GPU 驱动运行时仍然报failed to initialize NVML。可能原因驱动版本太老或者容器/WSL 没有正确透传 GPU。检查方式nvidia-smi nvcc --version解决方式升级 NVIDIA 驱动到与 CUDA Toolkit 匹配的版本。在容器中安装匹配的nvidia-container-toolkit。WSL 环境下确认 Windows 侧驱动支持 WSL GPU paravirtualization。6.2 SASS 反汇编失败或指令不识别现象工具运行到一半提示unknown instruction、unsupported opcode或者生成的 MLIR 中明显缺失指令。可能原因GPU 架构太新工具没有为该架构添加指令支持或者 SASS 文件中混入了不支持的扩展指令。检查方式cuobjdump -sass --dump-sass saxpy.cubin查看失败点对应的原始 SASS 指令再到工具源码中搜索该 opcode 是否被处理。解决方式切换目标架构例如统一使用sm_80或sm_90。合入工具的最新分支看是否已支持该架构。如果指令属于少量特殊扩展可以在工具 config 中关闭对应指令的建模把未知指令当作 opaque operation 透传。6.3 转出 MLIR 后语义变化现象优化后的 SASS 执行结果与原始 kernel 不一致或者数值误差超过预期。可能原因控制和数据流提升存在 bugbarrier 被错误消除浮点操作被错误融合或重排。检查方式对比原始 SASS 和优化后 SASS 的指令数量、同步指令数量。使用相同输入数据分别运行原始 cubin 和优化后 cubin。在 MLIR 层面做 pattern 前 dump 和 pattern 后 dump 对比。解决方式减少优化 pass逐个开启定位是哪个 pass 导致语义变化。确认浮点合约是否允许 FMA 融合。如果应用层对精度敏感可能需要禁用这类融合。在没有把握前不要直接替换生产环境的 cubin。6.4 寄存器压力过高导致 spill现象优化后 SASS 中出现了大量STL和LDL指令性能反而下降。可能原因某些优化 pass 延长了寄存器生命周期导致寄存器峰值升高最终触发局部内存溢出。检查方式cuobjdump -res-usage saxpy.cubin观察 register count 是否超过目标 GPU 的 64K 寄存器预算。解决方式减少线程数或 block size降低每个线程的寄存器压力。调整优化 pass去掉与当前 kernel 不匹配的展开类优化。在生成 SASS 时限制最大寄存器数例如加入--maxrregcount语义的配置。6.5 性能提升不稳定现象同一份优化结果在生产节点上有时候快 30%有时候只快 5%。可能原因节点被其他任务抢占GPU 时钟未锁定输入数据 shape 变化导致内存布局不同CPU 端启动和同步开销影响整体时间。解决方式使用ncu只测量 kernel 时间而不是包含 CPU 启动的总体时间。在专用 GPU 节点上测试并固定时钟。建立回归测试把性能基线纳入 CI。7. 实际项目中的落地建议与最佳实践7.1 适合 SASS2MLIR 的代码特征不是所有 CUDA 代码都值得进入 SASS2MLIR 处理链路。结合这类工具的优化特征下面几类代码更容易获得收益计算密集 kernel例如 GEMM、卷积、张量收缩。多 kernel 连续调用且数据强烈复用的场景适合利用 MLIR 做融合分析。在共享内存和寄存器布局上有明显优化空间的 kernel。运行时间长、调用频率高的热 kernel值得投入优化成本。如果一个 kernel 只运行一次或者运行时间只有几十微秒那么 SASS2MLIR 的编译、验证和发布成本很可能无法回收。7.2 落地流程建议建议按下面的顺序接入使用nsys/ncu对现有应用做 profiling选出时间占比最高的 top kernel。针对 top kernel 单独剥离一个 benchmark 工程固定输入和运行参数。跑通 SASS2MLIR 的 SASS 转 MLIR 流程先生成基线优化结果。做正确性对比和性能对比保留多组输入规模。如果收益足够再把这些优化过的 cubin 接入生产加载路径。加入回滚机制通过环境变量或配置控制是否使用优化版本。7.3 与既有优化手段的组合SASS2MLIR 可以和其他 GPU 优化方法组合。比如 CUDA Graph 能减少 kernel launch 开销SASS2MLIR 能减少 kernel 内部指令数两者互不冲突。NVTX 标记能帮助 profiling 定位到具体阶段SASS2MLIR 又能从指令层提高那个阶段的效率。如果项目本身使用 Triton 或 torch.compile生成的高层 kernel 最终也会落到 SASS。对这类路径SASS2MLIR 更像是一个“最后一道优化关卡”它不替代高层编译器而是修改编译器已经生成的底料。7.4 回归测试和风险控制任何针对机器码的自动优化都伴随风险。建议在回归测试中做四层检查检查层级覆盖内容推荐工具语义等价优化前后输出一致自建 CUDA 对比程序数值误差浮点误差在阈值内自定义 RMSE/MaxErr 检查性能基线指定输入下不慢于原版本benchmark 脚本硬件兼容目标 GPU 架构可正常加载cuobjdump / 加载测试编码上建议把 SASS2MLIR 的 passes 组合固化为配置文件使用与源码相同的版本管理流程。不要让业务线随意修改 pass 序列否则排查性能回退时很难定位根因。7.5 可以继续学习的方向如果看完这篇文章后想继续深入可以从以下几个方面入手学习 MLIR 的基础 dialect 和 pass 机制用mlir-opt编写自己的 pattern。研究 NVIDIA 官方 SASS 指令手册和cuobjdump用法提高读懂 SASS 的能力。用ncu分析真实 kernel 的 stall 原因把 profiling 结论与 SASS 指令序列对应起来。关注 SASS2MLIR 上游仓库的更新特别是新增架构支持和 pass 变化。尝试在一个小型计算 kernel 上跑通完整流程并记录优化前后指令数、寄存器数、执行时间的变化。GPU 优化的核心能力说到底是对“源码到机器码”这条链路的理解。SASS2MLIR 的价值不只在于它可能带来的性能提升更在于它提供了一套把最终代码重新打开、分析和优化的工程方法论。拿到一个新 kernel 时先做 profiling再决定是否引入 SASS2MLIR 这类工具最后用回归测试守住正确性和性能底线这才是生产环境中更稳妥的落地方案。