已合并
fix(triton_experimental): rebuild stale-rank store broadcast after r-tree promotion #44955
fix(triton_experimental): rebuild stale-rank store broadcast after r-tree promotion #44955
已合并
huyuchao创建于 16 天前
huyuchao成员
16 天前

【合入来源】

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

【合入来源】

【修改方案】

问题现象:npu_backend=triton_experimental 时,输出为单个标量的全量归约融合 kernel(CausalLM 的 log_softmax + nll_loss(ignore_index) 路径,典型如 triton_unk_fused_clone_nll_loss_forward_slice_view_*)编译失败:

ValueError('Cannot broadcast, rank mismatch: [1, 1, 1], [1, 1]')
tl.store(in_out_ptr0 + (tl.full([1, 1], 0, tl.int32).broadcast_to(XBLOCK, 1)), tmp15, None)

HuggingFace runner 的 ElectraForCausalLM / RobertaForCausalLM / XGLMForCausalLM 三模型均因同一 kernel 模板失败。

根因分析(上游 torch/_inductor 与 triton_experimental 的交互):

  1. 上游 store codegen(torch/_inductor/codegen/triton.py 的 integer-index store 分支)会手动追加 .broadcast_to(<value.shape>),其中 value.shape 取自发射时刻 CSE 记录的形状——此刻归约值与索引秩一致(均为 2),代码正确;
  2. triton_experimental 随后做 r 树提升(promoted r-tree)改写,把归约值表达式整形为 [None, None, :](秩 3),但 store 行中手动追加的 broadcast_to 参数与基座 tl.full 仍停留在提升前的秩 2;
  3. Triton 要求 tl.store 两侧秩相等,[1,1,1] vs [XBLOCK,1] 直接编译失败。

修复内容:在既有的 store 回改函数 _rewrite_reduction_store_shape(torch_npu/_inductor/triton_experimental/codegen/triton.py)内补全分支——当 broadcast_to 参数数小于 real_ndim(陈旧秩)时重建参数:r 槽强制置 1(常量索引 ⟹ 全量归约 ⟹ 值在 r 轴为 singleton,类型上保证)、旧的非 1 参数按序映射到其余槽位、基座 tl.full 的 shape 列表补齐到 real_ndim。写入地址、数值与冗余写次数均不变,仅修正索引表达式的形状元数据,语义严格等价;参数计数不匹配时保持原行不改(防御性 bail)。

【资料变更】

不涉及

【接口变更】

不涉及

【功能验证】

验证环境:Ascend NPU(aarch64)+ torch 2.13.0 + torch_npu(triton_experimental 后端)。

  1. 新增 UT:test/_inductor/test_triton_experimental_rtree_store.py,覆盖 CausalLM shift + log_softmax + nll_loss(ignore_index) 触发模式(双级 r 树提升 + 标量整数索引 store)。修复前该 UT 失败,报错与线上三模型完全一致(rank mismatch: [1, 1, 1], [1, 1]);修复后通过(含与 eager 数值一致性断言)。
  2. 既有冒烟回归:test/_inductor/test_triton_experimental_enable.py 13/13 通过,无回归。
  3. 端到端验证(HuggingFace benchmark runner,triton_experimental 后端):
    • ElectraForCausalLM:修复前编译失败 → 修复后 1.155x 加速、pass_accuracy;
    • RobertaForCausalLM:修复前编译失败 → 修复后 1.318x、pass_accuracy;
    • XGLMForCausalLM:修复前编译失败 → 修复后 1.502x、pass_accuracy;
    • 原正常模型(resnet18 / fastNLP_Bert 等)无劣化。

【CheckList】

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

变更摘要

该 PR 修复 npu_backend=triton_experimental 下全量归约(输出为单个标量)kernel 的编译失败问题。根因是上游 store codegen 追加的 .broadcast_to(<value.shape>) 参数取自 r 树提升(promoted r-tree)之前的 CSE 形状,提升后归约值秩升高(如 [None, None, :] 秩 3),而 store 行的 broadcast 参数与基座 tl.full 仍停留在提升前秩 2,导致 tl.store 两侧秩不匹配(Cannot broadcast, rank mismatch: [1, 1, 1], [1, 1]),典型触发路径为 CausalLM 的 log_softmax + nll_loss(ignore_index)(ElectraForCausalLM / RobertaForCausalLM / XGLMForCausalLM 均受影响)。修复在既有 store 回改函数 _rewrite_reduction_store_shapetorch_npu/_inductor/triton_experimental/codegen/triton.py)内补全分支:当 broadcast 参数个数小于 real_ndim 时重建参数,将 r 槽位置为 1、旧的非 1 参数按序映射到其余槽位,并把基座 tl.full 的 shape 列表补齐到 real_ndim;写入地址、数值与冗余写次数不变,仅修正索引表达式形状元数据,语义严格等价。同时新增回归 UT 覆盖该触发模式。

主要改动

  • 修复 _rewrite_reduction_store_shape 的陈旧秩分支torch_npu/_inductor/triton_experimental/codegen/triton.py):当 broadcast_to 参数个数小于 real_ndim 时,通过 ast.walk(store) 定位该调用并重建参数——r 槽强制置 1(常量索引 ⟹ 全量归约 ⟹ r 轴 singleton)、旧的非 1 参数按序填充空闲槽位,从而消除 store 两侧的秩不匹配。
  • 补齐基座 tl.full 的 shape 列表(同文件):当 tl.full 首参为 ast.List 且元素数小于 real_ndim 时,在头部插入常量 1 将列表补齐到 real_ndim,保证 broadcast 输入秩与目标秩一致。
  • 防御性 bail 逻辑(同文件):若未找到 broadcast_to 调用、参数个数已达 real_ndim,或旧非 1 参数数量与空闲槽位数不匹配,均保持原行不改,避免误改写。
  • 新增回归测试test/_inductor/test_triton_experimental_rtree_store.py):新增 TestPromotedRtreeScalarStore 测试类,通过 _causal_lm_loss(shift + log_softmax + nll_loss(ignore_index))复现双级 r 树提升加标量整数索引 store 的场景,断言编译成功、生成包含 triton_experimental 标记的 wrapper 代码,并与 eager 结果做数值一致性校验(assert_close)。
likedislike
atomgit-bot
atomgit-bot
16 天前 评论:

代码审查

✅ 未发现问题

likedislike
ascend-robotascend-robot成员
16 天前 添加了label:ascend-cla/yes
此处折叠了63条消息 查看更多
AtlasAccountAtlasAccount成员
5 天前 添加了label:ci-pipeline-passed
AtlasAccount
AtlasAccount成员
5 天前 评论:
流水线 PR-pipeline_pytorch#64908 [ commitID:21bbfd9d ] 已完成
>>>代码风格自动修复执行成功(无修复内容)
阶段 任务名 状态 详情
编译构建 Build_X86 >>>
Build_ARM >>>
Build_X86_torchair 🛑 >>>
Build_ARM_torchair 🛑 >>>
patch_test 🛑 >>>
Build_X86_213 >>>
Build_ARM_213 >>>
恶意代码检查 Antipoison >>>
编码安全与规范检查 codecheck_pre-commit >>>
check_error >>>
lintrunner >>>
开源片段检查 SCA >>>
开发者测试 UT_ARM_A3_Part_01 🛑 >>>
UT_ARM_A3_Part_02 🛑 >>>
UT_ARM_A2_Part_01 >>>
UT_ARM_A2_Part_02 >>>
UT_ARM_A2_Part_03 >>>
UT_inductor_Part_01 🛑 >>>
UT_inductor_Part_02 🛑 >>>
UT_inductor_Part_03 🛑 >>>
UT_inductor_Part_04 🛑 >>>
UT_DIST_ARM_Part_01 🛑 >>>
UT_DIST_ARM_Part_02 🛑 >>>
UT_DIST_ARM_Part_03 🛑 >>>
UT_DIST_ARM_Part_04 🛑 >>>
UT_ARM_A2_Select_Part_01 >>>
UT_ARM_A2_Select_Part_02 >>>
UT_ARM_A2_Part_213 >>>
UT_inductor_Part_213 🛑 >>>
UT_DIST_ARM_Part_213 🛑 >>>
UT_ARM_A2_Select_Part_213 >>>
流水线 PR-pipeline_pytorch >>>
此流水线已支持下列评论快捷指令,仅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成员
5 天前 关闭了关联的issue
ascend-robotascend-robot成员
5 天前 合入了pull request
AtlasAccount
AtlasAccount成员
5 天前 评论:
流水线 pytorch_gitcode_PR_multiVersion#14740 [ commitID:21bbfd9d ] 已完成
likedislike