Pull Request已成功合入, 合并人@ascend-robot
(感谢 gsl 的贡献)变更摘要
本 PR 修复了 2 个及以上 r-node 广播归约在编译期出现的 broadcast/rank 失败问题(涉及 test_cauchy、test_mark_unbacked_slice、test_buffer_use_after_remove 等用例)。核心思路是在 torch_npu/_inductor/triton_experimental/codegen/triton.py 中为 promoted r-tree 重写后的 masked load 指针索引与 store 偏移量补齐与 promoted 秩/块形状对齐的重写逻辑,并在 torch_npu/_inductor/triton_experimental/config.py 中新增开关 promoted_rtree_shape_fix 控制该行为。
主要改动
-
新增配置项
promoted_rtree_shape_fix: 在config.py中新增布尔配置(默认True),用于开启/关闭 promoted r-tree 形状修复,即对 masked load 指针(bar 形索引 vs 块形掩码)和 store 偏移(promotion 前秩 vs rank-N 值)按 promoted 块形状重新对齐,规避 triton-ascend 拒绝隐式 ptr-vs-mask 广播的问题。 -
构建 per-slot 块 token 并传入重写流程: 在
_npu_rewrite_promoted_rtree_body中为每个自由 r-node 生成对应的块 token(动态节点用{nm}_blk,静态节点用real_block_{nm},其余槽位为"1"),构成slot_block_tokens,与real_sizes一起传递给_npu_apply_promoted_rtree_lines。 -
扩展
_npu_apply_promoted_rtree_lines签名与 store 形状重写: 为该函数新增real_sizes、slot_block_tokens参数;在_rewrite_reduction_store_shape中,除保留原有的 value 侧 broadcast 槽位折叠(r-slots 置 1)外,新增 offset 侧修复:当开启promoted_rtree_shape_fix时,将 store 偏移tl.full([1,..], 0, dt).broadcast_to(XBLOCK, 1)的 full 形状秩前缀重写为[1]*real_ndim,并把 broadcast 目标改为按槽位取值(r-slots 为"1",保留槽位为对应real_sizes),以匹配 rank-N 的 value 形状。 -
新增
_fix_masked_load_ptr修复 load 侧指针秩: 当 masked load 的 ptr 索引仅覆盖部分自由 r-node(bar 形)而掩码已被重写为完整块形状时,对指针索引显式执行tl.broadcast_to(..., slot_block_tokens)后再与基址相加,使指针与掩码形状一致;该修复作为新的一步处理接在_rewrite_reduction_store_shape之后执行。


| 阶段 | 任务名 | 状态 | 详情 |
|---|---|---|---|
| 编译构建 | Build_X86 | ✅ COMPLETED | >>> |
| Build_ARM | ✅ COMPLETED | >>> | |
| Build_X86_torchair | ⚪ IGNORED | >>> | |
| Build_ARM_torchair | ⚪ IGNORED | >>> | |
| patch_test | ⚪ IGNORED | >>> | |
| Build_X86_213 | ✅ COMPLETED | >>> | |
| Build_ARM_213 | ✅ COMPLETED | >>> | |
| 恶意代码检查 | Antipoison | ✅ COMPLETED | >>> |
| 编码安全与规范检查 | codecheck_pre-commit | ✅ COMPLETED | >>> |
| check_error | ✅ COMPLETED | >>> | |
| lintrunner | ✅ COMPLETED | >>> | |
| 开源片段检查 | SCA | ✅ COMPLETED | >>> |
| 开发者测试 | UT_ARM_A3_Part_01 | ⚪ IGNORED | >>> |
| UT_ARM_A3_Part_02 | ⚪ IGNORED | >>> | |
| UT_ARM_A2_Part_01 | ✅ COMPLETED | >>> | |
| UT_ARM_A2_Part_02 | ✅ COMPLETED | >>> | |
| UT_ARM_A2_Part_03 | ✅ COMPLETED | >>> | |
| UT_inductor_Part_01 | ⚪ IGNORED | >>> | |
| UT_inductor_Part_02 | ⚪ IGNORED | >>> | |
| UT_inductor_Part_03 | ⚪ IGNORED | >>> | |
| UT_inductor_Part_04 | ⚪ IGNORED | >>> | |
| UT_DIST_ARM_Part_01 | ⚪ IGNORED | >>> | |
| UT_DIST_ARM_Part_02 | ⚪ IGNORED | >>> | |
| UT_DIST_ARM_Part_03 | ⚪ IGNORED | >>> | |
| UT_DIST_ARM_Part_04 | ⚪ IGNORED | >>> | |
| UT_ARM_A2_Select_Part_01 | ✅ COMPLETED | >>> | |
| UT_ARM_A2_Select_Part_02 | ✅ COMPLETED | >>> | |
| UT_ARM_A2_Part_213 | ✅ COMPLETED | >>> | |
| UT_inductor_Part_213 | ⚪ IGNORED | >>> | |
| UT_DIST_ARM_Part_213 | ⚪ IGNORED | >>> | |
| UT_ARM_A2_Select_Part_213 | ✅ COMPLETED | >>> | |
| 流水线 | PR-pipeline_pytorch | ✅ COMPLETED | >>> |
- compile、compile_inductor、compile_torchair : 运行流水线
- retry : 重试流水线所有失败子任务
- retry <任务名> : 仅重试指定失败子任务
- stop : 停止流水线


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


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




【合入来源】
【修改方案】
promoted rtree 形状对齐修复
promoted rtree flat-loop 重写把归约 mask 升级为全槽位块形状,但 (a) 被 mask 的 load 指针在索引只覆盖部分自由 r-node 时仍为条形;(b) store 偏移仍停留在提升前 rank——两者均依赖上游 Triton宽容、triton-ascend 拒绝的隐式 ptr-vs-mask 广播(NoTritonConfigsError: Cannot broadcast ...)。所有 2+ r-node 广播归约在编译期命中**;图带存活 x 轴(如 layernorm backward)时触发变体。
tl.broadcast_to到逐槽位块形状(非 r 槽位取real_sizes[s]真实保持轴块宽——存活 x 轴得 XBLOCK 而非 1);tl.fullrank 前缀改写为[1]*real_ndim,broadcast_to 目标逐槽位对齐(r 槽 1、保持槽真实块宽),与 value 侧塌缩镜像。1. 修复前的故障现场
1.1 load 侧:ptr 与 mask 形状不齐
重写后 kernel 的 load 段(3×3 等价形式):
r0_1 = tl.arange(0, 3)[:, None, None] # [3,1,1] 行索引 r0_0 = tl.arange(0, 3)[None, :, None] # [1,3,1] 列索引 tmp1 = tl.load(y + (r0_0), (r0_1mask & r0_0mask), other=0.0)(r0_1mask & r0_0mask)逐格相与,广播为[3,3,1];r0_0只携带列轴,停留在[1,3,1]。即
tl.load(ptr[1,3,1], mask[3,3,1]):掩码为 3×3 每格发放守卫位,而地址表只有第一行。上游 Triton 以隐式广播容纳该写法;triton-ascend 拒绝:
1.2 store 侧:偏移 rank 滞后(large_broadcast_reduction 类图)
tl.store(out_ptr0 + tl.full([1, 1], 0, tl.int32).broadcast_to(XBLOCK, 1), tmp4, None)tmp4:promoted 重写经tl.sum+ resize 升至 rank-3[1,1,1];tl.full([1,1], 0, dt).broadcast_to(XBLOCK, 1)是上游在 flat 单轴时代发射的标准偏移形态,rank-2。重写对循环体按 token 匹配改造,该表达式不含任何被匹配的名字
(长在循环外、由
tl.full与数字常量构成),因此未被升级。【资料变更】
【接口变更】
【功能验证】
【CheckList】