已关闭
[Bug]: triton experimental inductor后端,不兼容triton ascend 3.6 #4253
AllenGuan创建于  8月21日关闭于  8月25日
AllenGuan成员
8月21日 创建

在提交新问题之前,请确保您已经在社区中搜索过相关问题,并使用了社区中提供的资源/工具后,仍未找到满意的解决方式。

⚠️ 安全信息提醒:请仔细检查提供的文本内容,确保其不包含敏感数据信息,包括但不限于:

  • API 令牌或密钥
  • 密码或身份验证凭证
  • 私有网址或接口地址
  • 个人或机密数据
  • ...

在分享配置信息或代码示例时,请将敏感信息脱敏处理,或使用 <TOKEN> 等占位符替代原有内容。

环境信息

- 操作系统 openEuler 24.03 LTS SP3 (aarch64)
- 昇腾硬件信息 910B2
- CANN软件版本 9.1.0-beta.1
- 安装的对应软件版本 2.13(torch 2.13.0 + torch_npu 2.13.0)
- triton-ascend 3.6.0(whl 内 vendored triton core 3.5.0,import triton 后 triton.__version__ == '3.5.0')

🐛 问题描述

torch_npu 的 triton_experimental 后端(torch.compile(model, options={"npu_backend": "triton_experimental"}))在 triton-ascend 3.6.0 上无法完成任何 kernel 的编译/启动,共观察到三类症状(同一环境、不同触发条件)。在 triton-ascend 3.2.x(vendored core 3.2.0)上同样的代码工作正常;差异指向 vendored triton core 3.5.0 的 API 变化:launch hooks 从 CompiledKernel 类属性迁移到 triton.knobs.runtime.*,以及 AttrsDescriptor 类被整体移除(attrs 改用 plain dict 表达)。

症状一:任意 torch.compile 图在 launcher 创建阶段报 AttributeError(最常见的入口症状)

复现步骤:

import torch
import torch_npu  # noqa: F401

def f(a, b):
    return a + b

compiled = torch.compile(f, options={"npu_backend": "triton_experimental"})
a = torch.randn(1024, device="npu")
b = torch.randn(1024, device="npu")
print(compiled(a, b))

期望行为:编译成功,输出与 eager 一致的 a + b 结果。

实际行为:抛出异常,任何经过 triton kernel 的图均失败:

torch._inductor.exc.InductorError: AttributeError: type object 'CompiledKernel' has no attribute 'launch_enter_hook'

决定性栈帧:torch_npu/_inductor/triton_experimental/npu_triton_heuristics.py 的 NPUTritonCompileResult.make_launcher(构建 launcher scope 处直接访问 binary.__class__.launch_enter_hook / binary.__class__.launch_exit_hook)。triton core 3.5.0 中这两个属性已迁移到 triton.knobs.runtime.launch_enter_hook / launch_exit_hook,CompiledKernel 类上不再存在。

症状二:绕过症状一后,autotuner 报 "All tiling ... are not runnable"(launcher 参数表未剔除 constexpr)

若将 hooks 读取改为兼容写法,同一个最小复现片段继续失败,报错换为:

RuntimeError: All tiling for [triton_unk_fused_add_0] are not runnable.

该异常由 autotuner 在所有候选 tiling 的 PreRun 全部失败后抛出;per-tiling 的底层异常被 debug 日志吞掉,开启 debug 日志后可见根因为:

launcher() missing 1 required positional argument: 'XBLOCK'

定位:npu_triton_heuristics.py make_launcher 中按 triton_version_uses_attrs_dict()(core 3.5.0 上为 True)走 attrs_dict 分支构造 def_args/call_args 时,仅剔除了 num_warps/num_stages,未剔除 kernel 的 constexpr 形参(XBLOCK/YBLOCK/R0_BLOCK 等,已由选定的 Config 烘入编译产物,caller 不会传入),导致 launcher 的 Python 签名要求 XBLOCK 而调用方不提供。

症状三:动态 shape 的 OUTER 归约图在 codegen 阶段报 ImportError(rsplit combine kernel 路径)

复现步骤:

import torch
import torch_npu  # noqa: F401

class Model(torch.nn.Module):
    def forward(self, mm):
        view = mm.reshape(mm.shape[0], -1, 128)
        return view.to(torch.float32).sum(0)

model = torch.compile(Model().npu(), dynamic=None)
x = torch.randn(16384, 2048, device="npu", dtype=torch.bfloat16)
torch._dynamo.mark_dynamic(x, 0)
with torch.no_grad():
    print(model(x))

期望行为:编译成功,输出正确归约结果。

实际行为:整个 torch.compile 失败:

ImportError: cannot import name 'AttrsDescriptor' from 'triton.compiler.compiler' (.../site-packages/triton/compiler/compiler.py)

决定性栈帧:torch_npu/_inductor/triton_experimental/codegen/triton.py 的 _npu_emit_rsplit_combine(rsplit-outer 归约的跨核 partial+combine 拆分路径,对满足条件的单 kernel OUTER 归约默认启用),函数内无条件执行 from triton.compiler.compiler import AttrsDescriptor 并调用 AttrsDescriptor.from_dict(...) 构造 combine kernel 的 triton_meta["configs"]。triton core 3.5.0 已移除 AttrsDescriptor(3.2.x 中位于 triton.backends.compiler,并由 triton.compiler.compiler re-export),因此任何触发该路径的图(动态 shape 的 view+sum、rms_norm backward 等 OUTER 归约)都会失败。除 import 失败外,3.5.0 上 configs 亦要求 plain dict 形态的 attrs 表达,与 3.2.x 的 AttrsDescriptor 对象形态不同。

补充信息

  • 症状一/二可用同一个最小片段观察(症状一修复后症状二显形);症状三需要满足 rsplit 触发条件的归约图(例:动态首维 + reshape 后沿新轴 sum(0))。
  • 排查症状二时需注意 autotuner 会吞掉 per-tiling 异常,仅对外抛 "All tiling ... are not runnable",建议对 torch_npu._inductor.triton_experimental.npu_triton_heuristics logger 开启 DEBUG 级别查看底层报错。

欢迎加入社区,感谢您对社区的贡献 🎉!

likedislike
AAllenGuan成员
8月21日 关联了看板:FrameworkPTAdapter 版本issue看板
TorchNPU-BotTorchNPU-Bot成员
8月21日 添加了label:triage-review
TorchNPU-Bot
TorchNPU-Bot成员
8月21日 评论:

issue待分派,添加triage-review标签

likedislike
ascend-robotascend-robot成员
8月21日 添加了label:bug
AAllenGuan成员
8月21日 修改了issue 的描述
ascend-robotascend-robot成员
8月25日 关闭了 issue
ascend-robotascend-robot成员
8月25日 添加了label:resolved
chenrayraychenrayray成员
13 天前 关联了里程碑:v26.2.0
ascend-robotascend-robot成员
13 天前 issue状态由 TODO 改变为 DONE