Pull Request已成功合入, 合并人@ascend-robot
(感谢 shi-yufeng99 的贡献)变更摘要
本 PR 主要将 NPU Triton 后端的 mm/bmm 模板及 CV(Compute/Vector)epilogue 融合能力合入 master。核心改动包括:新增 NPU Triton 矩阵乘模板(npu_triton_mm_template、npu_persistent_mm_template、npu_triton_bmm_template)及对应的 tile 配置生成与候选注入逻辑,支持通过 {{store_output}} 融合下游逐元素算子;新增 addmm 的 bias epilogue 融合路径;同时补充了 patch_fixed_indexer 修复 size-1 维度的索引生成问题,并新增 Triton 模板启发式 template、调整 codegen_kernel_benchmark 调用签名及非进程池分支的模块加载逻辑。
主要改动
- 新增 NPU Triton MM 模板:在
kernel/mm.py中新增npu_triton_mm_template、_MM_TEMPLATE、_get_npu_mm_configs与add_npu_triton_mm_choices,并在tuned_mm中通过use_triton_template(layout)注入候选,使mm支持 Triton 内核及{{store_output}}epilogue 融合。 - 新增 NPU persistent MM 模板:新增
npu_persistent_mm_template、npu_persistent_mm_grid、_get_npu_persistent_mm_configs与add_npu_persistent_mm_choices,采用每个 program 处理多个 tile 的对角分核策略,并通过GROUP_M约束保证取模运算为编译期常量。 - 新增 NPU Triton addmm 与 bmm 模板:
add_npu_triton_addmm_choices复用npu_triton_mm_template并结合addmm_epilogue将 bias 融合进 epilogue;npu_triton_bmm_template通过tl.program_id(1)处理 batch 维,并在tuned_bmm中接入候选选择。 - 修复
_fixed_indexer的 size-1 维度跳过问题:新增patch_fixed_indexer(),覆盖torch._inductor.ir._fixed_indexer使其始终保留idx * stride项,避免mm输出存在 size-1 维时store_output退化为非连续路径并触发tt.store的 MLIR 编译错误。 - 补充 Triton 模板运行支撑:在
runtime/triton_heuristics.py新增template启发式函数并通过HeuristicType.TEMPLATE调用cached_autotune;同时在scheduler.py的非进程池分支直接PyCodeCache.load,并将codegen_kernel_benchmark调用调整为只传num_gb。


代码审查
我已完成对全部 8 个变更文件的审查。以下是按优先级汇总的审查结论。
发现汇总
-
P1(2 个):
torch_npu/_inductor/kernel/mm.py:574—addmm_epilogue未导入,addmm Triton 融合路径必然NameError(被except吞掉,功能静默失效)。torch_npu/_inductor/kernel/bmm.py:50—npu_bmm_grid缺少min参数,SymbolicGridFn无条件注入min后必抛TypeError,bmm Triton 融合失效。
-
P2(2 个):
3.torch_npu/_inductor/kernel/mm.py:301—even_k = (k % block_k == 0)对符号 K 不可靠(动态 shape 下EVEN_K失真),persistent 配置中同类问题一并指出。
4.torch_npu/_inductor/kernel/bmm.py:173— 同上的 bmm 版本。 -
P3(4 个,均为死导入/冗余导入):
5.torch_npu/_inductor/kernel/mm.py:2-3—Path/Optional/load_template未使用。
6.torch_npu/_inductor/kernel/bmm.py:1-15—functools/Path/fallback_handler/ExternKernelChoice/TritonTemplate(及inductor_bmm)未使用。
7.torch_npu/_inductor/__init__.py:97-100—aggresive_autotune未使用、log冗余。
8.torch_npu/_inductor/select_algorithm.py:57—gen_common_triton_imports未使用。
各文件审查确认
torch_npu/_inductor/__init__.py— 已审查(P3 死导入)。torch_npu/_inductor/codegen/ir.py— 已审查(patch_fixed_indexer为有意的全局补丁,文档详尽,no issues)。torch_npu/_inductor/codegen/scheduling.py— 已审查(codegen_kernel_benchmark去掉grid参数与 NPU 子类签名对齐,no issues)。torch_npu/_inductor/kernel/bmm.py— 已审查(P1 grid 签名、P2 even_k、P3 死导入)。torch_npu/_inductor/kernel/mm.py— 已审查(P1 addmm_epilogue、P2 even_k、P3 死导入)。torch_npu/_inductor/runtime/triton_heuristics.py— 已审查(新增template与上游镜像,签名与cached_autotune匹配,no issues)。torch_npu/_inductor/scheduler.py— 已审查(非进程池分支补上PyCodeCache.load,是修复,no issues)。torch_npu/_inductor/select_algorithm.py— 已审查(P3 死导入)。
整体风险判断
本 PR 引入的 mm/bmm/addmm Triton CV 融合功能存在两个高概率功能性缺陷(addmm 的 addmm_epilogue NameError 与 bmm 的 grid 函数缺参 TypeError),两者都会导致对应融合路径完全不可用且以静默方式失效;此外动态 shape 下的 EVEN_K 整除判断存在正确性隐患。建议合入前至少修复两个 P1 问题并确认动态 shape 场景的行为。
| 类型 | 数量 |
|---|---|
| 🔴 阻塞 | 1 |
| 🟡 建议 | 0 |
⛔ 需要修改


/lgtm
/approve


Pull Request 已合并或已关闭。
If you want to solve this problem, you can click here to do it in the FAQs.




【合入来源】
【修改方案】
#3016
本 PR 主要将 NPU Triton 后端的 mm/bmm 模板及 CV(Compute/Vector)epilogue 融合能力合入 master。核心改动包括:新增 NPU Triton 矩阵乘模板(npu_triton_mm_template、npu_persistent_mm_template、npu_triton_bmm_template)及对应的 tile 配置生成与候选注入逻辑,支持通过 {{store_output}} 融合下游逐元素算子;新增 addmm 的 bias epilogue 融合路径;同时补充了 patch_fixed_indexer 修复 size-1 维度的索引生成问题,并新增 Triton 模板启发式 template、调整 codegen_kernel_benchmark 调用签名及非进程池分支的模块加载逻辑。
主要改动
新增 NPU Triton MM 模板:在 kernel/mm.py 中新增 npu_triton_mm_template、_MM_TEMPLATE、_get_npu_mm_configs 与 add_npu_triton_mm_choices,并在 tuned_mm 中通过 use_triton_template(layout) 注入候选,使 mm 支持 Triton 内核及 {{store_output}} epilogue 融合。
新增 NPU persistent MM 模板:新增 npu_persistent_mm_template、npu_persistent_mm_grid、_get_npu_persistent_mm_configs 与 add_npu_persistent_mm_choices,采用每个 program 处理多个 tile 的对角分核策略,并通过 GROUP_M 约束保证取模运算为编译期常量。
新增 NPU Triton addmm 与 bmm 模板:add_npu_triton_addmm_choices 复用 npu_triton_mm_template 并结合 addmm_epilogue 将 bias 融合进 epilogue;npu_triton_bmm_template 通过 tl.program_id(1) 处理 batch 维,并在 tuned_bmm 中接入候选选择。
修复 _fixed_indexer 的 size-1 维度跳过问题:新增 patch_fixed_indexer(),覆盖 torch._inductor.ir._fixed_indexer 使其始终保留 idx * stride 项,避免 mm 输出存在 size-1 维时 store_output 退化为非连续路径并触发 tt.store 的 MLIR 编译错误。
补充 Triton 模板运行支撑:在 runtime/triton_heuristics.py 新增 template 启发式函数并通过 HeuristicType.TEMPLATE 调用 cached_autotune;同时在 scheduler.py 的非进程池分支直接 PyCodeCache.load,并将 codegen_kernel_benchmark 调用调整为只传 num_gb。
【资料变更】
【接口变更】
【功能验证】
【CheckList】