已合并
fix(inductor): triton_experimental UT fixes + promoted rtree shape alignment #45171
fix(inductor): triton_experimental UT fixes + promoted rtree shape alignment #45171
已合并
gsl创建于 17 天前
gsl成员
17 天前

【合入来源】

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

【修改方案】

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)时触发变体。

  • load:索引张量显式 tl.broadcast_to 到逐槽位块形状(非 r 槽位取 real_sizes[s] 真实保持轴块宽——存活 x 轴得 XBLOCK 而非 1);
  • store:偏移 tl.full rank 前缀改写为 [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 拒绝:

ValueError: Cannot broadcast ... [3,3,1] vs [1,3,1]
→ 全部候选 config 编译失败 → NoTritonConfigsError → test_cauchy(_dynamic_shapes) FAIL

1.2 store 侧:偏移 rank 滞后(large_broadcast_reduction 类图)

tl.store(out_ptr0 + tl.full([1, 1], 0, tl.int32).broadcast_to(XBLOCK, 1), tmp4, None)
  • value tmp4:promoted 重写经 tl.sum + resize 升至 rank-3 [1,1,1]
  • offset:tl.full([1,1], 0, dt).broadcast_to(XBLOCK, 1) 是上游在 flat 单轴时代发射的
    标准偏移形态,rank-2。重写对循环体按 token 匹配改造,该表达式不含任何被匹配的名字
    (长在循环外、由 tl.full 与数字常量构成),因此未被升级。
ValueError: rank mismatch: [1, 1, 1], [1, 1]

【资料变更】

不涉及

【接口变更】

不涉及

【功能验证】

1f2fca88-b95b-47bd-b59e-2f65c1284018.png

1.png

【CheckList】

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

likedislike
Pull Request已成功合入, 合并人@ascend-robot
(感谢 gsl 的贡献)
Ggsl成员
17 天前 创建了 pull request,commit d8a9f204
atomgit-bot
atomgit-bot
17 天前 评论:

变更摘要

本 PR 修复了 2 个及以上 r-node 广播归约在编译期出现的 broadcast/rank 失败问题(涉及 test_cauchytest_mark_unbacked_slicetest_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_sizesslot_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 之后执行。

likedislike
不准确?
atomgit-bot
atomgit-bot
17 天前 评论:

代码审查

✅ 未发现问题

likedislike
不准确?
ascend-robotascend-robot成员
17 天前 添加了label:ascend-cla/yes
此处折叠了146条消息 查看更多
AtlasAccount
AtlasAccount成员
7 天前 评论:
流水线 PR-pipeline_pytorch#65251 [ commitID:6af532d7 ] 已完成
>>>代码风格自动修复执行成功(无修复内容)
阶段 任务名 状态 详情
编译构建 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 >>>
此流水线已支持下列评论快捷指令,仅PR创建者和白名单成员[wujinyuan1, huangjingwei, liangsongwei, yashi999, culechan, Dring, wuyouqi1, L1919_snow, qq_52711437, WhiteNight12, nomiz, xiu_21, ffmh, wanglijun55, hss-shuai, husichao, smallsilly, lanshaozuishuai, jimmyisme1, lzy0920232, alpha-junh, Sunshine_Youngster, wei_zhuoyi, zhangyihuiben, zyw-hw, zzzkeke, rmch, yangch0324, LucciC, AACAES, renyujin, wjlflyer, senzhen-town, pengjingyou, qsc97, limuan, yule100, xiaoqi-zhou, kuhn7, chenxingying, hanye02, zichun_ye, anyrenwei, kkjocker, wangzili121, Lu_G, yvjc, puddingfjz, HandsoemLemon, bigprestigee1, huawuyi, zhenyu10, dairenjie, du-jin-hang, zou-jieyu, adelaideliu, TrHan, wanlinan, Windwindzzz, pengqihw, kisnwang, yuheng_wang, honghao_wang, jizewei, zhangguoguang, sunyu-xuan, chenrayray, hbhu_bin, liujunzhu, c_34, LiNuoh, maoyuanpeng1, zzhongmin, zhaoyu65, bellatan, jiabaolin, zhuofanshen, wencaiwen, lu_zhuge, caoshuyang, molly12, lyx324521, LQ1206, gitcode-chenjiao, cai-weiwei1989, CHDong, ogqin, yuanlipingGit, xuqinglin1, lqz2, zouwei1, chaoluoa, paradox325, jackzhang1116, yaoyao, akh, yujiacheng, dengjie0116, Hubert11111, Shine_Ws, wslhj555, longqiand, OYtao666, JiaqingQiang, luyyyy, Kingbelial, zhanghaiyu0101, wenxp1018, yanliu-luoluo, ksun_sekiro, liyong328, wgzheng, tangky, vivi_is_coding, aoiaoisola, weixin_44494597, wangmengmengwang65667, hid57809721, qq_35468730, comeonup, C547032, gcw_m5OQChA4, yao_yao_ling_xian, cnnbwcy, szqfes_12, cora_19, cann_lilin, can, shawnylee233, fanglanyue0916, hhz0, LiNuohang, taohuoquan, Jesse, WSs_321, SCh_zx]评论有效
  • compile、compile_inductor、compile_torchair : 运行流水线
  • retry : 重试流水线所有失败子任务
  • retry <任务名> : 仅重试指定失败子任务
  • stop : 停止流水线
likedislike
ascend-robotascend-robot成员
7 天前 合入了pull request
ascend-robot
ascend-robot成员
7 天前 评论:

Pull Request 已合并或已关闭。

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

likedislike
ascend-robot
ascend-robot成员
7 天前 评论:

Pull Request 已合并或已关闭。

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

likedislike
AtlasAccount
AtlasAccount成员
7 天前 评论:
流水线 pytorch_gitcode_PR_multiVersion#14828 [ commitID:6af532d7 ] 已完成
likedislike