已合并
import triton backend for mm and triton cv fusion to master branch #44574
import triton backend for mm and triton cv fusion to master branch #44574
已合并
shi-yufeng99创建于 21 天前
shi-yufeng99
21 天前

【合入来源】

如有社区issue,请关联issue链接
请勿携带内部流程信息(需求链接、问题单、内部issue等)

【修改方案】

#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。

【资料变更】

请确认是否涉及资料变更。如涉及,需要在PR中体现,并简要说明修改内容。如不涉及,需填写“不涉及”

【接口变更】

请确认是否涉及跨代码仓或者客户面可见的接口变更。如涉及,需要详细说明接口以及对应的变更内容,同时需要在资料中体现。如不涉及,需填写“不涉及”

【功能验证】

说明测试场景,测试方法。如果本次测试方式与常规单元测试不同,请详细说明您的测试步骤
新增/变更内容是否已新增/适配UT测试用例看护,并补充测试自验证截图

【CheckList】

PR提交人对以下CheckList自检项进行全量自检,自检通过或不涉及,均修改 [ ] 为 [x]

likedislike
Pull Request已成功合入, 合并人@ascend-robot
(感谢 shi-yufeng99 的贡献)
Sshi-yufeng99
21 天前 创建了 pull request,commit 65bb303e
Sshi-yufeng99
21 天前 关联了issue:Triton cv fusion
atomgit-bot
atomgit-bot
21 天前 评论:

变更摘要

本 PR 主要将 NPU Triton 后端的 mm/bmm 模板及 CV(Compute/Vector)epilogue 融合能力合入 master。核心改动包括:新增 NPU Triton 矩阵乘模板(npu_triton_mm_templatenpu_persistent_mm_templatenpu_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_configsadd_npu_triton_mm_choices,并在 tuned_mm 中通过 use_triton_template(layout) 注入候选,使 mm 支持 Triton 内核及 {{store_output}} epilogue 融合。
  • 新增 NPU persistent MM 模板:新增 npu_persistent_mm_templatenpu_persistent_mm_grid_get_npu_persistent_mm_configsadd_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
likedislike
atomgit-bot
atomgit-bot
21 天前 评论:

代码审查

我已完成对全部 8 个变更文件的审查。以下是按优先级汇总的审查结论。

发现汇总

  • P1(2 个)

    1. torch_npu/_inductor/kernel/mm.py:574addmm_epilogue 未导入,addmm Triton 融合路径必然 NameError(被 except 吞掉,功能静默失效)。
    2. torch_npu/_inductor/kernel/bmm.py:50npu_bmm_grid 缺少 min 参数,SymbolicGridFn 无条件注入 min 后必抛 TypeError,bmm Triton 融合失效。
  • P2(2 个)
    3. torch_npu/_inductor/kernel/mm.py:301even_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-3Path/Optional/load_template 未使用。
    6. torch_npu/_inductor/kernel/bmm.py:1-15functools/Path/fallback_handler/ExternKernelChoice/TritonTemplate(及 inductor_bmm)未使用。
    7. torch_npu/_inductor/__init__.py:97-100aggresive_autotune 未使用、log 冗余。
    8. torch_npu/_inductor/select_algorithm.py:57gen_common_triton_imports 未使用。

各文件审查确认

  1. torch_npu/_inductor/__init__.py — 已审查(P3 死导入)。
  2. torch_npu/_inductor/codegen/ir.py — 已审查(patch_fixed_indexer 为有意的全局补丁,文档详尽,no issues)。
  3. torch_npu/_inductor/codegen/scheduling.py — 已审查(codegen_kernel_benchmark 去掉 grid 参数与 NPU 子类签名对齐,no issues)。
  4. torch_npu/_inductor/kernel/bmm.py — 已审查(P1 grid 签名、P2 even_k、P3 死导入)。
  5. torch_npu/_inductor/kernel/mm.py — 已审查(P1 addmm_epilogue、P2 even_k、P3 死导入)。
  6. torch_npu/_inductor/runtime/triton_heuristics.py — 已审查(新增 template 与上游镜像,签名与 cached_autotune 匹配,no issues)。
  7. torch_npu/_inductor/scheduler.py — 已审查(非进程池分支补上 PyCodeCache.load,是修复,no issues)。
  8. 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

⛔ 需要修改

likedislike
此处折叠了142条消息 查看更多
dezheng889成员
17 天前 评论:

/lgtm
/approve

likedislike
ascend-robotascend-robot成员
17 天前 添加了label:approvedlgtm
ascend-robotascend-robot成员
17 天前 合入了pull request
ascend-robot
ascend-robot成员
17 天前 评论:

Pull Request 已合并或已关闭。

If you want to solve this problem, you can click here to do it in the FAQs.

likedislike
ascend-robot
ascend-robot成员
17 天前 评论:
流水线 pytorch_gitcode_PR_multiVersion#14137 [ commitID:4344938c ] 已完成
likedislike